How to compile opencl-kernel-file(.cl) to LLVM IR

Question

This question is related to LLVM/clang. I already know how to compile opencl-kernel-file.cl using the OpenCL API (clBuildProgram() and clGetProgramBuildInfo()).

My question is: How to compile opencl-kernel-file(.cl) to LLVM IR with OpenCL 1.2 or higher?
In the other words, How to compile opencl-kernel-file(.cl) to LLVM IR without libclc?

I have tried various methods to generate LLVM-IR from an OpenCL kernel.

I first followed the clang user manual.(https://clang.llvm.org/docs/UsersManual.html#opencl-features) but it did not run.

Secondly, I found a way to use libclc:

clang++ -emit-llvm -c -target -nvptx64-nvidial-nvcl -Dcl_clang_storage_class_specifiers -include /usr/local/include/clc/clc.h -fpack-struct=64 -o "$@".bc "$@" <br>
llvm-link "$@".bc /usr/local/lib/clc/nvptx64--nvidiacl.bc -o "$@".linked.bc <br>
llc -mcpu=sm_52 -march=nvptx64 "$@".linked.bc -o "$@".nvptx.s

This method worked fine, but since libclc was built on top of the OpenCL 1.1 specification, it couldn't be used with OpenCL 1.2 or later code such as code using printf.

And this method uses libclc, which implements OpenCL built-in functions in the shape of new function. You can observe that in the assembly (ptx) of result OpenCL binary, it goes straight to the function call instead of converting it to an inline assembly. I am concerned that this will affect GPU performance.

So I am looking for a way to replace compilation using libclc. As a last resort, I'm considering using libclc with the NVPTX backend and AMDGPU backend of LLVM. But if there is already another way, I want to use it. (I expect that the OpenCL front-end I have not found yet exists in clang)

My program's scenarios are:

1. There is opencl kernel source file(.cl)
2. Compile the file to LLVM IR
3. IR-Level process to the IR
4. Compile(using llc) the IR to Binary
   • with each gpu targets(nvptx, amdgcn..)
5. Using the binary, Run host(.c or .cpp with lib OpenCL) with clCreateProgramWithBinary()

Now, When I compile kernel source file to LLVM IR, I have to include header of libclc(-include option in first one of above command) for compiling built-in functions. And I have to link libclc libraries before compile IR to binary

My environments are below:

• GTX960
  
  • NVIDIA's Binary appears in nvptx format
    
  • I'm using sm_52 nvptx for my gpu.
    
• Ubuntu Linux 16.04 LTS
  
• LLVM/Clang 5.0.0
  
  • If there is another way, I am willing to change the LLVM version.

Thanks in advice!

Answer 1

Clang 9 (and up) can compile OpenCL kernels written in the OpenCL C language. You can tell Clang to emit LLVM-IR by passing the -emit-llvm flag (add -S to output the IR in text rather than in bytecode format), and specify which version of the OpenCL standard using e.g. -cl-std=CL2.0. Clang currently supports up to OpenCL 2.0.

By default, Clang will not add the standard OpenCL headers, so if your kernel uses any of the OpenCL built-in functions you may see an error like the following:

clang-9 -c -x cl -emit-llvm -S -cl-std=CL2.0 my_kernel.cl -o my_kernel.ll
my_kernel.cl:17:12: error: implicit declaration of function 'get_global_id' is invalid in OpenCL
  int i = get_global_id(0);
          ^
1 error generated.

You can tell Clang to include the standard OpenCL headers by passing the -finclude-default-header flag to the Clang frontend, e.g.

clang-9 -c -x cl -emit-llvm -S -cl-std=CL2.0 -Xclang -finclude-default-header my_kernel.cl -o my_kernel.ll

Answer 2

(I expect that the OpenCL front-end I have not found yet exists in clang)

There is an OpenCL front-end in clang - and you're using it, otherwise you couldn't compile a single line of OpenCL with clang. Frontend is Clang recognizing the OpenCL language. There is no OpenCL backend of any kind in LLVM, it's not the job of LLVM; it's the job of various OpenCL implementations to provide proper libraries. Clang+LLVM just recognizes the language and compiles it to bitcode & machine binaries, that's all it does.

in the assembly(ptx) of result opencl binary, it goes straight to the function call instead of converting it to an inline assembly.

You could try linking to a different library instead of libclc, if you find one. Perhaps NVidia's CUDA has some bitcode libraries somewhere, then again licensing issues... BTW are you 100% sure you need LLVM IR ? getting OpenCL binaries using the OpenCL runtime, or using SPIR-V, might get you faster binaries & certainly be less painful to work with. Even if you manage to get a nice LLVM IR, you'll need some runtime which actually accepts it (i could be wrong, but i doubt proprietary AMD/NVIDIA OpenCL will just accept random LLVM IR as inputs).

Answer 3

Clang does not provide a standard CL declaration header file (for example, C's stdio.h), which is why you're getting "undefined type float" and whatnot.

If you get one such header, you can then mark it as implicit include using "clang -include cl.h -x cl [your filename here]"

One such declaration header can be retrieved from the reference OpenCL compiler implementation at

https://github.com/KhronosGroup/SPIR-Tools/blob/master/headers/opencl_spir.h

And by the way, consider using this compiler which generates SPIR (albeit 1.0) which can be fed into OpenCL drivers as input.