The patchset could not apply on master. Could you check and rebase it to latest master?
And I have some other comments as below: On Sat, Feb 28, 2015 at 01:53:04PM +0800, [email protected] wrote: > 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")){ Use strcmp is not good here. Please check the Class Triple, and use getArch() to get the arch type and check whether the arch type is Triple::spir. > + module->setTargetTriple("spir"); Do we need to set the module triple here? Is it already "spir" target? > + } > 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) { Just as we discussed, we need to keep consistent with the standard "opencl_spir.h". And use a separate patch to fix all the related prototypes should be better here. > 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 || Use FROM_LLVM_SPIR should be better then FROM_INTERMEDIATE for me. Thanks, Zhigang Gong. _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
