如何将 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 与 NVPTX 后端和 LLVM 的 AMDGPU 后端一起使用。
但如果已经有其他方法，我想使用它。
（我希望我还没有找到的 OpenCL 前端在 clang 中存在）

我的程序的场景是：

1. 有opencl内核源文件(.cl)
2. 将文件编译为 LLVM IR
3. IR 级流程到 IR
4. 编译（使用 llc）IR 到二进制
   • 每个 gpu 目标（nvptx, amdgcn..）
5. 使用二进制文件，使用 clCreateProgramWithBinary() 运行主机（.c 或带有 lib OpenCL 的 .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 内核。您可以通过传递 -emit-llvm 标志告诉 Clang 发出 LLVM-IR（添加 -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.

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

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

（我希望我还没有找到的OpenCL前端在clang中存在）

clang 中有一个 OpenCL 前端 - 您正在使用它，否则您无法使用 clang 编译一行 OpenCL。前端是识别 OpenCL 语言的 Clang。 LLVM 中没有 OpenCL 任何类型的后端，这不是 LLVM 的工作；提供适当的库是各种 OpenCL 实现的工作。 Clang+LLVM 只是识别语言并将其编译为位码和机器二进制文件，仅此而已。

在结果 opencl 二进制文件的程序集（ptx）中，它直接进入函数调用，而不是将其转换为内联程序集。

您可以尝试链接到其他库而不是 libclc，如果找到的话。也许 NVidia 的 CUDA 在某处有一些位码库，然后又是许可问题……顺便说一句，你 100% 确定你需要 LLVM IR 吗？使用 OpenCL 运行时或使用 SPIR-V 获取 OpenCL 二进制文件可能会让您更快地生成二进制文件，并且使用起来肯定不会那么痛苦。即使您设法获得了一个不错的 LLVM IR，您也需要一些实际接受它的运行时（我可能错了，但我怀疑专有的 AMD/NVIDIA OpenCL 只会接受随机 LLVM IR 作为输入）。

Answer 3

Clang 不提供标准的 CL 声明头文件（例如，C 的 stdio.h），这就是为什么你会得到“未定义类型浮点”之类的。

如果您得到一个这样的标头，则可以使用“clang -include cl.h -x cl [your filename here]”将其标记为隐式包含

一个这样的声明头可以从参考 OpenCL 编译器实现中检索到

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

顺便说一句，考虑使用这个生成 SPIR（尽管是 1.0）的编译器，它可以作为输入输入到 OpenCL 驱动程序中。