From: Luo Xionghu <[email protected]> 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 <[email protected]> --- 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 _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
