cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

hariseldon
Journeyman III

Running OpenCL SPIR programs on GPU

Hi all,

as reported here: Re: Re: Standard Portable Intermediate 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

0 Likes
9 Replies
gopal
Staff

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!

0 Likes
hariseldon
Journeyman III

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

0 Likes

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

0 Likes

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.

0 Likes

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,

0 Likes

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

0 Likes

The SPIR binary seems OK. I can load it in my program. Can you post your code loading the SPIR binary and building the program?

You may also try this sample code for loading SPIR.

Thank you smaliu,

your version runs fine on my machine.

In my program the problem was the missing null character at the end of the input .bc file.

Cheers,

Hari

0 Likes
samliu
Adept I

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?

0 Likes