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.g...@intel.com>
---
backend/src/backend/gen_insn_selection.cpp | 3 ++-
backend/src/backend/program.h | 1 +
backend/src/ir/instruction.cpp | 1 +
backend/src/ir/instruction.hpp | 1 +
backend/src/llvm/llvm_gen_backend.cpp | 25 ++++++++++++++++++-------
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 e025698..1d999fa 100644
--- a/backend/src/backend/gen_insn_selection.cpp
+++ b/backend/src/backend/gen_insn_selection.cpp
@@ -3246,7 +3246,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 4065a17..554fb16 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 583bab5..12bc1bf 100644
--- a/backend/src/ir/instruction.cpp
+++ b/backend/src/ir/instruction.cpp
@@ -1254,6 +1254,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 09b0148..f7024d4 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 ec79628..0487bcb 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);
@@ -1443,7 +1456,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:
@@ -2817,12 +2830,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) {
@@ -3646,7 +3658,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");
@@ -3740,15 +3752,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)) {
@@ -3795,7 +3806,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
--
1.9.1
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
