如何将 opencl-kernel-file(.cl) 编译为 LLVM IR

Question

这个问题与LLVM/clang有关。
我已经知道如何使用 OpenCL API（ clBuildProgram() 和 clGetProgramBuildInfo() ）编译 opencl-kernel-file(.cl)

我的问题是：
如何使用 OpenCL 1.2 或更高版本将 opencl-kernel-file(.cl) 编译为 LLVM IR？
换句话说，如何在没有 libclc 的情况下将 opnecl-kernel-file(.cl) 编译为 LLVM IR？

我尝试了各种方法来获取 OpenCL-Kernel-File 的 LLVM-IR。

我首先遵循 clang 用户手册。( https://clang.llvm.org/docs/UsersManual.html#opencl-features）但它没有运行。

其次，我找到了使用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<br>

此方法工作正常，但由于 libclc 是在 OpenCL 1.1 规范之上构建的，因此它无法与 OpenCL 1.2 或更高版本的代码（例如使用 printf 的代码）一起使用。

该方法使用libclc，它以新函数的形式实现了OpenCL内置函数。您可以观察到，在结果 opencl 二进制文件的程序集（ptx）中，它直接进入函数调用，而不是将其转换为内联程序集。我担心这会影响 GPU 的行为和性能，例如执行时间。

所以现在我正在寻找一种使用libclc代替编译的方法。
作为最后的手段，我正在考虑将 libclc 与 LLVM 的 NVPTX 后端和 AMDGPU 后端一起使用。
但如果已经有另一种方法，我想使用它。
（我希望我还没有找到的OpenCL前端存在于clang中）

我的程序的场景是：

1. 有opencl内核源文件(.cl)
2. 将文件编译为 LLVM IR
3. IR 级流程至 IR
4. 将 IR 编译（使用 llc）为二进制
   • 每个 GPU 目标（nvptx、amdgcn ..）
5. 使用二进制文件，使用 clCreateProgramWithBinary() 运行主机（带有 lib OpenCL 的 .c 或 .cpp）

现在，当我将内核源文件编译为 LLVM IR 时，我必须包含 libclc 标头（上述命令第一个命令中的 -include 选项）以编译内置函数。在将 IR 编译为二进制文件之前，我必须链接 libclc 库

我的环境如下：

• GTX960
  - NVIDIA 的二进制文件以 nvptx 格式显示
  - 我的 GPU 使用 sm_52 nvptx。
  
• Ubuntu Linux 16.04 LTS
  
• LLVM/Clang 5.0.0
  - 如果有其他方法，我愿意更改 LLVM 版本。

谢谢指教！

Answer 1

Clang 9（及更高版本）可以编译用 OpenCL C 语言编写的 OpenCL 内核。您可以通过传递标志来告诉 Clang 发出 LLVM-IR -emit-llvm（添加-S以文本形式而不是字节码格式输出 IR），并使用例如指定 OpenCL 标准的版本-cl-std=CL2.0。Clang 目前最高支持 OpenCL 2.0。

默认情况下，Clang 不会添加标准 OpenCL 标头，因此如果您的内核使用任何 OpenCL 内置函数，您可能会看到如下错误：

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.

您可以通过将标志传递给 Clang 前端来告诉 Clang 包含标准 OpenCL 标头，例如-finclude-default-header

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