diff options
| -rw-r--r-- | backend/src/llvm/llvm_gen_backend.cpp | 45 | ||||
| -rw-r--r-- | kernels/compiler_program_global.cl | 11 |
2 files changed, 54 insertions, 2 deletions
diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp index dda8c9f4..40f421e7 100644 --- a/backend/src/llvm/llvm_gen_backend.cpp +++ b/backend/src/llvm/llvm_gen_backend.cpp @@ -1497,6 +1497,36 @@ namespace gbe Type::TypeID id = type->getTypeID(); GBE_ASSERT(c); + if (isa<ConstantExpr>(c)) { + const ConstantExpr *expr = dyn_cast<ConstantExpr>(c); + Value *pointer = expr->getOperand(0); + if (expr->getOpcode() == Instruction::GetElementPtr) { + uint32_t constantOffset = 0; + CompositeType* CompTy = cast<CompositeType>(pointer->getType()); + for(uint32_t op=1; op<expr->getNumOperands(); ++op) { + int32_t TypeIndex; + ConstantInt* ConstOP = dyn_cast<ConstantInt>(expr->getOperand(op)); + GBE_ASSERTM(ConstOP != NULL, "must be constant index"); + TypeIndex = ConstOP->getZExtValue(); + GBE_ASSERT(TypeIndex >= 0); + constantOffset += getGEPConstOffset(unit, CompTy, TypeIndex); + CompTy = dyn_cast<CompositeType>(CompTy->getTypeAtIndex(TypeIndex)); + } + + ir::Constant cc = unit.getConstantSet().getConstant(pointer->getName()); + unsigned int defOffset = cc.getOffset(); + relocs.push_back(ir::RelocEntry(offset, defOffset + constantOffset)); + + uint32_t size = getTypeByteSize(unit, type); + memset((char*)mem+offset, 0, size); + offset += size; + } else if (expr->isCast()) { + Constant *constPtr = cast<Constant>(pointer); + getConstantData(constPtr, mem, offset, relocs); + offset += getTypeByteSize(unit, type); + } + return; + } if (isa<GlobalVariable>(c)) { ir::Constant cc = unit.getConstantSet().getConstant(c->getName()); unsigned int defOffset = cc.getOffset(); @@ -1594,8 +1624,21 @@ namespace gbe offset += sizeof(double); break; } - default: + case Type::TypeID::HalfTyID: + { + const ConstantFP *cf = dyn_cast<ConstantFP>(c); + llvm::APFloat apf = cf->getValueAPF(); + llvm::APInt api = apf.bitcastToAPInt(); + uint64_t v64 = api.getZExtValue(); + uint16_t v16 = static_cast<uint16_t>(v64); + *(unsigned short *)((char*)mem+offset) = v16; + offset += sizeof(short); + break; + } + default: { + c->dump(); NOT_IMPLEMENTED; + } } } diff --git a/kernels/compiler_program_global.cl b/kernels/compiler_program_global.cl index 405c53f5..fbe030fe 100644 --- a/kernels/compiler_program_global.cl +++ b/kernels/compiler_program_global.cl @@ -31,6 +31,9 @@ global int a = 1; global int b = 2; global int * constant gArr[2]= {&a, &b}; +global int a_var[1] = {0}; +global int *p_var = a_var; + __kernel void compiler_program_global0(const global int *src, int dynamic) { size_t gid = get_global_id(0); /* global read/write */ @@ -60,9 +63,15 @@ __kernel void compiler_program_global1(global int *dst, int dynamic) { dst[12] = *p; dst[13] = s; dst[14] = l; - dst[15] = *gArr[dynamic]; + if (p_var == a_var) + dst[15] = *gArr[dynamic]; if (gid < 11) dst[gid] = ba[gid]; } +__kernel void nouse(int dynamic) { + c[0].s1 = &s2; + p_var = a+dynamic; +} + |