GBE: Use actual bti information to determine a pointer's addressspace.

Due to the private constant buffer support, it introduces private address
space mixed with constant address space some time. And more generic, one
constant address space may be located in private address space in LLVM IR
layer. Such as the following code:
  __kernel ...
  {
    const int2 foo[] = {{0, 1}, {2, 3}};
    int2 data = foo[get_global_id(0) % 2];
  }

The foo is in private address space but we finally will use __constant bti
to access it in Gen backend. The the above code will cause a assertion fail
in gen insturcion selection stage, because it generate a vector loading
instruction on a __constant buffer.

So we should use the actual BTI data to determine one pointer's address space
rather than get it from the LLVM IR layer.

Signed-off-by: Zhigang Gong <zhigang.gong@intel.com>
Reviewed-by: "Yang, Rong R" <rong.r.yang@intel.com>

5 files changed, 23 insertions, 8 deletions

diff --git a/backend/src/backend/gen_insn_selection.cpp b/backend/src/backend/gen_insn_selection.cpp
index becb1c99..68ca7610 100644
--- a/backend/src/backend/gen_insn_selection.cpp
+++ b/backend/src/backend/gen_insn_selection.cpp
@@ -3412,7 +3412,8 @@ namespace gbe
       GBE_ASSERT(insn.getAddressSpace() == MEM_GLOBAL ||
                  insn.getAddressSpace() == MEM_CONSTANT ||
                  insn.getAddressSpace() == MEM_PRIVATE ||
-                 insn.getAddressSpace() == MEM_LOCAL);
+                 insn.getAddressSpace() == MEM_LOCAL ||
+                 insn.getAddressSpace() == MEM_MIXED);
       //GBE_ASSERT(sel.isScalarReg(insn.getValue(0)) == false);
       const Type type = insn.getValueType();
       const uint32_t elemSize = getByteScatterGatherSize(type);
diff --git a/backend/src/backend/program.h b/backend/src/backend/program.h
index 4065a17e..554fb16d 100644
--- a/backend/src/backend/program.h
+++ b/backend/src/backend/program.h
@@ -70,6 +70,7 @@ enum gbe_get_arg_info_value {
 #define BTI_MAX_WRITE_IMAGE_ARGS  8
 #define BTI_WORKAROUND_IMAGE_OFFSET 128
 #define BTI_MAX_ID 253
+#define BTI_LOCAL 0xfe
 
 /*! Constant buffer values (ie values to setup in the constant buffer) */
 enum gbe_curbe_type {
diff --git a/backend/src/ir/instruction.cpp b/backend/src/ir/instruction.cpp
index f5580cbc..698e33ee 100644
--- a/backend/src/ir/instruction.cpp
+++ b/backend/src/ir/instruction.cpp
@@ -1262,6 +1262,7 @@ namespace ir {
       case MEM_LOCAL: return out << "local";
       case MEM_CONSTANT: return out << "constant";
       case MEM_PRIVATE: return out << "private";
+      case MEM_MIXED: return out << "mixed";
       case MEM_INVALID: return out << "invalid";
     };
     return out;
diff --git a/backend/src/ir/instruction.hpp b/backend/src/ir/instruction.hpp
index ddcefcee..f9fdfd4a 100644
--- a/backend/src/ir/instruction.hpp
+++ b/backend/src/ir/instruction.hpp
@@ -58,6 +58,7 @@ namespace ir {
     MEM_LOCAL,      //!< Local memory (thread group memory)
     MEM_CONSTANT,   //!< Immutable global memory
     MEM_PRIVATE,    //!< Per thread private memory
+    MEM_MIXED,      //!< mixed address space pointer.
     MEM_INVALID
   };
 
diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index 30f5cb33..9f4ed48c 100644
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -290,6 +290,19 @@ namespace gbe
     return ir::MEM_GLOBAL;
   }
 
+  static INLINE ir::AddressSpace btiToGen(const ir::BTI &bti) {
+    if (bti.count > 1)
+      return ir::MEM_MIXED;
+    uint8_t singleBti = bti.bti[0];
+    switch (singleBti) {
+      case BTI_CONSTANT: return ir::MEM_CONSTANT;
+      case BTI_PRIVATE: return  ir::MEM_PRIVATE;
+      case BTI_LOCAL: return ir::MEM_LOCAL;
+      default: return ir::MEM_GLOBAL;
+    }
+    return ir::MEM_GLOBAL;
+  }
+
   static Constant *extractConstantElem(Constant *CPV, uint32_t index) {
     ConstantVector *CV = dyn_cast<ConstantVector>(CPV);
     GBE_ASSERT(CV != NULL);
@@ -1445,7 +1458,7 @@ namespace gbe
                 incBtiBase();
               break;
               case ir::MEM_LOCAL:
-                ctx.input(argName, ir::FunctionArgument::LOCAL_POINTER, reg,  llvmInfo, ptrSize, align, 0xfe);
+                ctx.input(argName, ir::FunctionArgument::LOCAL_POINTER, reg,  llvmInfo, ptrSize, align, BTI_LOCAL);
                 ctx.getFunction().setUseSLM(true);
               break;
               case ir::MEM_CONSTANT:
@@ -2822,12 +2835,11 @@ namespace gbe
     CallSite::arg_iterator AE = CS.arg_end();
     GBE_ASSERT(AI != AE);
 
-    unsigned int llvmSpace = (*AI)->getType()->getPointerAddressSpace();
-    const ir::AddressSpace addrSpace = addressSpaceLLVMToGen(llvmSpace);
     const ir::Register dst = this->getRegister(&I);
 
     ir::BTI bti;
     gatherBTI(&I, bti);
+    const ir::AddressSpace addrSpace = btiToGen(bti);
     vector<ir::Register> src;
     uint32_t srcNum = 0;
     while(AI != AE) {
@@ -3605,7 +3617,7 @@ namespace gbe
               new_bti = BTI_CONSTANT;
               break;
             case 3:
-              new_bti = 0xfe;
+              new_bti = BTI_LOCAL;
               break;
             default:
               GBE_ASSERT(0 && "address space not unhandled in gatherBTI()\n");
@@ -3699,15 +3711,14 @@ namespace gbe
   template <bool isLoad, typename T>
   INLINE void GenWriter::emitLoadOrStore(T &I)
   {
-    unsigned int llvmSpace = I.getPointerAddressSpace();
     Value *llvmPtr = I.getPointerOperand();
     Value *llvmValues = getLoadOrStoreValue(I);
     Type *llvmType = llvmValues->getType();
     const bool dwAligned = (I.getAlignment() % 4) == 0;
-    const ir::AddressSpace addrSpace = addressSpaceLLVMToGen(llvmSpace);
     const ir::Register ptr = this->getRegister(llvmPtr);
     ir::BTI binding;
     gatherBTI(&I, binding);
+    const ir::AddressSpace addrSpace = btiToGen(binding);
 
     Type *scalarType = llvmType;
     if (!isScalarType(llvmType)) {
@@ -3754,7 +3765,7 @@ namespace gbe
       const ir::RegisterFamily pointerFamily = ctx.getPointerFamily();
       const ir::RegisterFamily dataFamily = getFamily(type);
 
-      if(dataFamily == ir::FAMILY_DWORD && addrSpace != ir::MEM_CONSTANT) {
+      if(dataFamily == ir::FAMILY_DWORD && addrSpace != ir::MEM_CONSTANT && addrSpace != ir::MEM_MIXED) {
         // One message is enough here. Nothing special to do
         if (elemNum <= 4) {
           // Build the tuple data in the vector