[Beignet] [PATCH 1/2] enable cl_khr_spir extension to build and run from SPIR binary.
xionghu.luo at intel.com
xionghu.luo at intel.com
Fri Feb 27 21:53:04 PST 2015
From: Luo Xionghu <xionghu.luo at intel.com>
the SPIR are built by clang generating a standard llvm Module file,
beignet need insert one byte before the module repesents binary type
then parse the module to link.
enable cl_khr_spir extension output string;
enable the SPIR calling conversion of CallingConv::SPIR_KERNEL;
get_global_id shoud be OVERLOADABLE; fix some bugs in prinf parse
and backend.
Signed-off-by: Luo Xionghu <xionghu.luo at intel.com>
---
backend/src/backend/gen_program.cpp | 4 ++++
backend/src/libocl/include/ocl_workitem.h | 2 +-
backend/src/libocl/src/ocl_workitem.cl | 2 +-
backend/src/libocl/tmpl/ocl_defines.tmpl.h | 1 +
backend/src/llvm/llvm_gen_backend.cpp | 5 ++++-
backend/src/llvm/llvm_printf_parser.cpp | 3 ++-
backend/src/llvm/llvm_scalarize.cpp | 1 +
src/cl_api.c | 1 +
src/cl_extensions.c | 4 ++++
src/cl_program.c | 21 +++++++++++++++++++--
src/cl_program.h | 3 ++-
11 files changed, 40 insertions(+), 7 deletions(-)
diff --git a/backend/src/backend/gen_program.cpp b/backend/src/backend/gen_program.cpp
index 4cfb703..0917909 100644
--- a/backend/src/backend/gen_program.cpp
+++ b/backend/src/backend/gen_program.cpp
@@ -260,6 +260,10 @@ namespace gbe {
llvm::MemoryBuffer* memory_buffer = llvm::MemoryBuffer::getMemBuffer(llvm_bin_str, "llvm_bin_str");
acquireLLVMContextLock();
llvm::Module* module = llvm::ParseIR(memory_buffer, Err, c);
+ // if load 32 bit spir binary, the triple should be spir-unknown-unknown.
+ if(!strcmp(module->getTargetTriple().c_str(), "spir-unknown-unknown")){
+ module->setTargetTriple("spir");
+ }
releaseLLVMContextLock();
if(module == NULL){
GBE_ASSERT(0);
diff --git a/backend/src/libocl/include/ocl_workitem.h b/backend/src/libocl/include/ocl_workitem.h
index 7534ee8..e1a59df 100644
--- a/backend/src/libocl/include/ocl_workitem.h
+++ b/backend/src/libocl/include/ocl_workitem.h
@@ -22,7 +22,7 @@
uint get_work_dim(void);
uint get_global_size(uint dimindx);
-uint get_global_id(uint dimindx);
+OVERLOADABLE uint get_global_id(uint dimindx);
uint get_local_size(uint dimindx);
uint get_local_id(uint dimindx);
uint get_num_groups(uint dimindx);
diff --git a/backend/src/libocl/src/ocl_workitem.cl b/backend/src/libocl/src/ocl_workitem.cl
index f4629f8..f14443e 100644
--- a/backend/src/libocl/src/ocl_workitem.cl
+++ b/backend/src/libocl/src/ocl_workitem.cl
@@ -52,6 +52,6 @@ DECL_PUBLIC_WORK_ITEM_FN(get_global_offset, 0)
DECL_PUBLIC_WORK_ITEM_FN(get_num_groups, 1)
#undef DECL_PUBLIC_WORK_ITEM_FN
-uint get_global_id(uint dim) {
+OVERLOADABLE uint get_global_id(uint dim) {
return get_local_id(dim) + get_local_size(dim) * get_group_id(dim) + get_global_offset(dim);
}
diff --git a/backend/src/libocl/tmpl/ocl_defines.tmpl.h b/backend/src/libocl/tmpl/ocl_defines.tmpl.h
index 4e210be..fe999b2 100644
--- a/backend/src/libocl/tmpl/ocl_defines.tmpl.h
+++ b/backend/src/libocl/tmpl/ocl_defines.tmpl.h
@@ -34,5 +34,6 @@
#define cl_khr_byte_addressable_store
#define cl_khr_icd
#define cl_khr_gl_sharing
+#define cl_khr_spir
#endif /* end of __OCL_COMMON_DEF_H__ */
diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index 0f84215..c8b0207 100644
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -1566,7 +1566,9 @@ error:
llvmInfo.typeName = (cast<MDString>(typeNameNode->getOperand(1 + argID)))->getString();
llvmInfo.accessQual = (cast<MDString>(accessQualNode->getOperand(1 + argID)))->getString();
llvmInfo.typeQual = (cast<MDString>(typeQualNode->getOperand(1 + argID)))->getString();
- llvmInfo.argName = (cast<MDString>(argNameNode->getOperand(1 + argID)))->getString();
+ if(argNameNode){
+ llvmInfo.argName = (cast<MDString>(argNameNode->getOperand(1 + argID)))->getString();
+ }
// function arguments are uniform values.
this->newRegister(I, NULL, true);
@@ -2212,6 +2214,7 @@ error:
case CallingConv::PTX_Kernel:
#else
case CallingConv::C:
+ case CallingConv::SPIR_KERNEL:
#endif
break;
default: GBE_ASSERTM(false, "Unsupported calling convention");
diff --git a/backend/src/llvm/llvm_printf_parser.cpp b/backend/src/llvm/llvm_printf_parser.cpp
index 52da2e5..d6894af 100644
--- a/backend/src/llvm/llvm_printf_parser.cpp
+++ b/backend/src/llvm/llvm_printf_parser.cpp
@@ -564,6 +564,7 @@ error:
case CallingConv::PTX_Kernel:
#else
case CallingConv::C:
+ case CallingConv::SPIR_KERNEL:
#endif
break;
default:
@@ -594,7 +595,7 @@ error:
continue;
}
- if (call->getCalledFunction()->getIntrinsicID() != 0)
+ if (call->getCalledFunction() && call->getCalledFunction()->getIntrinsicID() != 0)
continue;
Value *Callee = call->getCalledValue();
diff --git a/backend/src/llvm/llvm_scalarize.cpp b/backend/src/llvm/llvm_scalarize.cpp
index 4df849f..6dd7b37 100644
--- a/backend/src/llvm/llvm_scalarize.cpp
+++ b/backend/src/llvm/llvm_scalarize.cpp
@@ -788,6 +788,7 @@ namespace gbe {
case CallingConv::PTX_Kernel:
#else
case CallingConv::C:
+ case CallingConv::SPIR_KERNEL:
#endif
break;
default: GBE_ASSERTM(false, "Unsupported calling convention");
diff --git a/src/cl_api.c b/src/cl_api.c
index 972c687..c715c0b 100644
--- a/src/cl_api.c
+++ b/src/cl_api.c
@@ -941,6 +941,7 @@ clBuildProgram(cl_program program,
/* TODO support create program from binary */
assert(program->source_type == FROM_LLVM ||
program->source_type == FROM_SOURCE ||
+ program->source_type == FROM_INTERMEDIATE ||
program->source_type == FROM_BINARY);
if((err = cl_program_build(program, options)) != CL_SUCCESS) {
goto error;
diff --git a/src/cl_extensions.c b/src/cl_extensions.c
index d07a525..cea2dd8 100644
--- a/src/cl_extensions.c
+++ b/src/cl_extensions.c
@@ -34,8 +34,12 @@ void check_opt1_extension(cl_extensions_t *extensions)
{
int id;
for(id = OPT1_EXT_START_ID; id <= OPT1_EXT_END_ID; id++)
+ {
if (id == EXT_ID(khr_icd))
extensions->extensions[id].base.ext_enabled = 1;
+ if (id == EXT_ID(khr_spir))
+ extensions->extensions[id].base.ext_enabled = 1;
+ }
}
void
diff --git a/src/cl_program.c b/src/cl_program.c
index c30f85e..00b620a 100644
--- a/src/cl_program.c
+++ b/src/cl_program.c
@@ -231,7 +231,21 @@ cl_program_create_from_binary(cl_context ctx,
program->binary_sz = lengths[0];
program->source_type = FROM_BINARY;
- if(isBitcode((unsigned char*)program->binary+1, (unsigned char*)program->binary+program->binary_sz)) {
+ if(isBitcode((unsigned char*)program->binary, (unsigned char*)program->binary+program->binary_sz)) {
+
+ char* typed_binary;
+ TRY_ALLOC(typed_binary, cl_calloc(lengths[0]+1, sizeof(char)));
+ memcpy(typed_binary+1, binaries[0], lengths[0]);
+ *typed_binary = 1;
+ program->opaque = compiler_program_new_from_llvm_binary(program->ctx->device->vendor_id, typed_binary, program->binary_sz+1);
+ cl_free(typed_binary);
+ if (UNLIKELY(program->opaque == NULL)) {
+ err = CL_INVALID_PROGRAM;
+ goto error;
+ }
+
+ program->source_type = FROM_INTERMEDIATE;
+ }else if(isBitcode((unsigned char*)program->binary+1, (unsigned char*)program->binary+program->binary_sz)) {
if(*program->binary == 1){
program->binary_type = CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT;
}else if(*program->binary == 2){
@@ -499,6 +513,9 @@ cl_program_build(cl_program p, const char *options)
memcpy(p->build_opts, options, strlen(options));
p->source_type = p->source ? FROM_SOURCE : p->binary ? FROM_BINARY : FROM_LLVM;
+ if (strstr(options, "-x spir")) {
+ p->source_type = FROM_INTERMEDIATE;
+ }
}
}
@@ -526,7 +543,7 @@ cl_program_build(cl_program p, const char *options)
/* Create all the kernels */
TRY (cl_program_load_gen_program, p);
- } else if (p->source_type == FROM_LLVM) {
+ } else if (p->source_type == FROM_LLVM || p->source_type == FROM_INTERMEDIATE) {
if (!CompilerSupported()) {
err = CL_COMPILER_NOT_AVAILABLE;
goto error;
diff --git a/src/cl_program.h b/src/cl_program.h
index 3ab7acd..98e9f11 100644
--- a/src/cl_program.h
+++ b/src/cl_program.h
@@ -33,7 +33,8 @@ struct _gbe_program;
enum {
FROM_SOURCE = 0,
FROM_LLVM = 1,
- FROM_BINARY = 2
+ FROM_BINARY = 2,
+ FROM_INTERMEDIATE = 3
};
/* This maps an OCL file containing some kernels */
--
1.9.1
More information about the Beignet
mailing list