[Beignet] [PATCH 1/2] support const indexed global constant array

[Homer Hsing]

support reading global constant arrays by CONST index.

example:
  constant int o[3] = {71, 72, 73};

  kernel void f(global int *dst) {
    dst[get_global_id(0)] = o[2]; // const index: 2
  }

in llvm converting phase, calculate offset from const index,
then add the offset to array head

Signed-off-by: Homer Hsing <homer.xing at intel.com>
---
 backend/src/llvm/llvm_gen_backend.cpp | 22 ++++++++++++++++++++++
 1 file changed, 22 insertions(+)

diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index 40f2667..260e3c1 100644
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -816,6 +816,28 @@ namespace gbe
         uint16_t reg = unit.getConstantSet().getConstant(name).getReg();
         return ir::Register(reg);
       }
+      if (isa<ConstantExpr>(CPV)) {
+        ConstantExpr *CE = dyn_cast<ConstantExpr>(CPV);
+        GBE_ASSERT(CE->isGEPWithNoNotionalOverIndexing());
+        SmallVector<Value*, 4> ValueOperands;
+        for (auto i = CE->op_begin(), e = CE->op_end(); i != e; ++i)
+          ValueOperands.push_back(cast<Value>(i));
+        ArrayRef<Value*> ops(ValueOperands);
+        GetElementPtrInst *GEP = GetElementPtrInst::CreateInBounds(ops[0], ops.slice(1));
+        auto pointer = GEP->getOperand(0);
+        auto offset1 = dyn_cast<ConstantInt>(GEP->getOperand(1));
+        GBE_ASSERT(offset1->getZExtValue() == 0);
+        auto offset2 = dyn_cast<ConstantInt>(GEP->getOperand(2));
+        int type_size = pointer->getType()->getTypeID() == Type::TypeID::DoubleTyID ? sizeof(double) : sizeof(int);
+        int type_offset = offset2->getSExtValue() * type_size;
+        auto pointer_name = pointer->getName().str();
+        ir::Register pointer_reg = ir::Register(unit.getConstantSet().getConstant(pointer_name).getReg());
+        ir::Register offset_reg = ctx.reg(ir::RegisterFamily::FAMILY_DWORD);
+        ctx.LOADI(ir::Type::TYPE_S32, offset_reg, ctx.newIntegerImmediate(type_offset, ir::Type::TYPE_S32));
+        ir::Register reg = ctx.reg(ir::RegisterFamily::FAMILY_DWORD);
+        ctx.ADD(ir::Type::TYPE_S32, reg, pointer_reg, offset_reg);
+        return reg;
+      }
       const ir::ImmediateIndex immIndex = this->newImmediate(CPV, elemID);
       const ir::Immediate imm = ctx.getImmediate(immIndex);
       const ir::Register reg = ctx.reg(getFamily(imm.type));
-- 
1.8.1.2