[Beignet] [PATCH 1/2] enable cl_khr_spir extension to build and run from SPIR binary.

[Zhigang Gong]

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, xionghu.luo at intel.com wrote:
> 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")){
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.