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

问题描述 投票:0回答:3

这个问题与 LLVM/clang 有关。我已经知道如何使用 OpenCL API(

opencl-kernel-file.cl
clBuildProgram()
)编译
clGetProgramBuildInfo()

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

我尝试了各种方法从 OpenCL 内核生成 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

此方法工作正常,但由于 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版本。

谢谢指教!

clang llvm opencl llvm-ir
3个回答
5
投票

Clang 9(及更高版本)可以编译用 OpenCL C 语言编写的 OpenCL 内核。您可以通过传递

-emit-llvm
标志(添加
-S
以文本形式而不是字节码格式输出 IR)来告诉 Clang 发出 LLVM-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 包含

标准 OpenCL 标头
,例如

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

0
投票

(我希望我还没有找到的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 作为输入)。


0
投票

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

如果您得到一个这样的标头,则可以使用“clang -include cl.h -x cl [此处的文件名]”将其标记为隐式包含

可以从参考 OpenCL 编译器实现中检索这样一个声明头,网址为

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

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

© www.soinside.com 2019 - 2024. All rights reserved.