9 Replies Latest reply on Aug 19, 2014 11:57 AM by hariseldon

    Running OpenCL SPIR programs on GPU

    hariseldon

      Hi all,

       

      as reported here: Re: Re: Standard Portable Intermedopencl_spir.hiate Representation (SPIR)

      cl_khr_spir happears as an available extension for GPUs with the new OpenCL driver.

      I am currently trying to generate spir programs using open-source LLVM.

       

      These are my steps:

      1) clang -x cl -O3 -target spir64 -include  opencl_spir.h program.cl -S -emit-llvm -fno-builtin -o - | llvm-as -o program.bc

      2) I create the program from binary loading program.bc

      3) I build the program using the option -x spir, as documented here: cl_khr_spir (although this does not affect the final error)

       

      Then I get the following build error:

       

      loadBitcode failed
      Frontend phase failed compilation.
      Error: Compilation from LLVMIR binary to IL text failed!
      

       

      Is it currently possible to run SPIR programs on GPUs ?

      What am I missing in my procedure ?

      Is there some documentation on how to run SPIR programs ?

       

      Btw, the same problem arises when compiling for the CPU (an Intel I7 in my case) using the AMD CPU runtime.

       

      Thank you in advance.

      Hari

        • Re: Running OpenCL SPIR programs on GPU
          gopal

          Hi Hari,

           

          AFAIK, there are 2 ways to generate SPIR:

           

          1.  Generating SPIR by offline compiler;

               a. Use clang options for generating SPIR:  -cl-std=cl1.2  -emit-llvm  -triple spir[32][64]-unknown-unknown

               b. use this binary to create program object using clCreateProgramWithBinary().  [This is to load SPIR Binary]

               c. build program object using clBuildProgram(), with -x spir as compile option.

           

          2.  Generating SPIR in host program,

               a. first load OpenCL source code by clCreateProgramWithSource()

               b. compile program with clCompileProgram() with a vendor specific option for generating SPIR. Use -x spir option to indicate that binary is in SPIR format.

               c. get the binaries with clGetProgramInfo with CL_PROGRAM_BINARIES

               d. use this binaries to create program object using clCreateProgramWithBinary().  [This is to load SPIR Binary]

               e. now build program object using clBuildProgram(). (Note: If you have multiple binaries, then first compile this program object using clCompileProgram() and later link all the compiled program objects using clLinkProgram().)

           

          You have used offline approach to generate SPIR, which seems to be work fine. The reason of failed compilation could be because of some mistake the way you have generated SPIR.

          1. Can you re-check your first step, whether you have used right options to generate SPIR?

          2. Replace target with triple flag, and see whether it works?.

           

          Thanks!

          • Re: Running OpenCL SPIR programs on GPU
            hariseldon

            Hi Gopal,

             

            thank you for your reply.

            I am actually interested in generating SPIR using the offline compiler.

            Concerning you suggestion in point 1.a:

            1) -cl-std=cl1.2 is not a valid command line option for clang.

            It's an OpenCL specific options for clBuildProgram: https://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clBuildProgram.html

             

            2) passing "-triple spir64-unknown-unknown" to clang instead of "-target spir64" results in a compilation failure.

            Clang does not recognize the triple.

             

            I personally think that the way I am generating SPIR code is correct: the LLVM-IR has the right metadata.

            The problem might be a bug in the AMD runtime or a missing flag in the clBuildProgram (other than -x spir).

             

            Cheers,

            Hari

              • Re: Running OpenCL SPIR programs on GPU
                gopal

                Hi Hari,

                 

                1. @-cl-std=cl1.2 is not a valid command line option for clang.

                As per my understanding, this option controls the version of OpenCL C that the compiler should accept. This is because, to use SPIR, you must have atleast OpenCL 1.2.


                As per the SPIR repository from this site : KhronosGroup/SPIR · GitHub,

                if you want to generate SPIR binary from valid OpenCL C file, use following command:

                 

                clang -cc1 -emit-llvm-bc -triple <triple> <OpenCL compile options> -cl-spir-compile-options "<OpenCL compile options>" -include <opencl_spir.h> -o <output> <input>


                As you can see above that you can specify the OpenCL options even with Clang compiler too.

                 

                2. @passing "-triple spir64-unknown-unknown" to clang instead of "-target spir64" results in a compilation failure.

                I am not sure why it does not recognize the triple option. As per the above command, it should not fail. Can you go through the above link and verify that you installed LLVM/Clang tool correctly?


                Thanks

                  • Re: Running OpenCL SPIR programs on GPU
                    hariseldon

                    Hi gopal,

                     

                    thank you for the link to the github account. I did not know about that.

                     

                    So I have been following strictly the instructions in their README file.

                    I compiled LLVM and clang from scratch, as suggested.

                    So LLVM and clang are version 3.2 now.

                     

                    I run the following command:

                    clang -cc1 -emit-llvm-bc -triple spir64-unknown-unknown -cl-std=CL1.2 -cl-spir-compile-options "-cl-std=CL1.2" -include opencl_spir.h program.cl -o program.bc

                    (Notice that CL1.2 must be upper case)

                    (I have tried both triples spir64-unknown-unknown and spir-unknown-unknown: same error)

                     

                    I create the program from binary using program.bc and I compile it using options "-x spir -cl-std=CL1.2".

                    I still get the usual error:

                     

                    loadBitcode failed

                    Frontend phase failed compilation.

                    Error: Compilation from LLVMIR binary to IL text failed!

                     

                    At this point I think there is a bug in the AMD runtime library.

                     

                    Cheers,

                    Hari.

                      • Re: Running OpenCL SPIR programs on GPU
                        dipak

                        Hi,

                         

                        Please try to follow Sam's suggestion and if still face the problem, please post the SPIR binary code to investigate in details.

                        Note: There was some issue to compile the SPIR code using earlier driver e.g. 13.35 and 14.10, but it should work fine with later drivers from 14.20 or 14.30.

                        So, please try with the latest driver.

                         

                        Regards,

                          • Re: Running OpenCL SPIR programs on GPU
                            hariseldon

                            Hi Dipak,

                             

                            At the moment I am using a driver of version 1526.3

                            The device (HD 7970) is 32 bits so I am using "-triple spir-unknown-unknown"

                            Here follows the spir code that causes the problem. It's an empty kernel generated using the command:

                             

                            clang -cc1 -emit-llvm-bc -triple spir-unknown-unknown -cl-std=CL1.2 -cl-spir-compile-options "-cl-std=CL1.2" -include opencl_spir.h mt.cl -S -emit-llvm -o mt.ll

                             

                            ##########################################################

                            target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024"

                            target triple = "spir-unknown-unknown"

                             

                            define cc76 void @mt() nounwind readnone {

                              ret void

                            }

                             

                            !opencl.kernels = !{!0}

                            !opencl.enable.FP_CONTRACT = !{}

                            !opencl.spir.version = !{!6}

                            !opencl.ocl.version = !{!6}

                            !opencl.used.extensions = !{!7}

                            !opencl.used.optional.core.features = !{!7}

                            !opencl.compiler.options = !{!8}

                             

                            !0 = metadata !{void ()* @mt, metadata !1, metadata !2, metadata !3, metadata !4, metadata !5}

                            !1 = metadata !{metadata !"kernel_arg_addr_space"}

                            !2 = metadata !{metadata !"kernel_arg_access_qual"}

                            !3 = metadata !{metadata !"kernel_arg_type"}

                            !4 = metadata !{metadata !"kernel_arg_type_qual"}

                            !5 = metadata !{metadata !"kernel_arg_base_type"}

                            !6 = metadata !{i32 1, i32 2}

                            !7 = metadata !{}

                            !8 = metadata !{metadata !"-cl-std=CL1.2"}

                            ##########################################################

                             

                            __kernel void mt() {

                            }

                             

                            I am giving in input to the compiler the mt.bc file generated by:

                            clang -cc1 -emit-llvm-bc -triple spir-unknown-unknown -cl-std=CL1.2 -cl-spir-compile-options "-cl-std=CL1.2" -include opencl_spir.h mt.cl -o mt.bc

                             

                            Can you find any problems in the mt.ll file ?

                             

                            Cheers,

                            Hari

                    • Re: Running OpenCL SPIR programs on GPU
                      samliu

                      One thing to notice: The target triple of SPIR should match the device address bits. Please check device address bits with clGetDeviceInfo/CL_DEVICE_ADDRESS_BITS in your host program and load SPIR with target triple corresponding to the device address bits.

                       

                      Did you try some simple kernel? Does it fail for all kernels, or only fails for some specific kernels?