CL_INVALID_KERNEL_NAME 带有在 Intel 上用 clang 创建的 SPIRV 程序

Question

我正在尝试为我正在研究的库实现我的 OpenCL 计算着色器的离线编译，以准备使用 SPIRV 进行开发。最好的方法似乎是使用 clang 和 LLVM，我一直在遵循 Khronos 博客的指南：

https://www.khronos.org/blog/whats-new-in-clang-release-14-for-opencl-developers

https://www.khronos.org/blog/offline-compilation-of-opencl-kernels-into-spir-v-using-open-source-tooling

使用以下命令（我也尝试过使用 spir64 目标）：

> clang -cl-std=CL1.2 -c -target spir -O0 -emit-llvm -o test_cl_spv.bc test_cl_spv.cl
> llvm-spirv test_cl_spv.bc -o test_cl_spv.spv
> spirv-opt test_cl_spv.spv -o test_cl_spv.spv
> spirv-link test_cl_spv.spv -o test_cl_spv.spv

我已经在我的 Ubuntu AWS 实例上下载了相关的存储库（clang/llvm v17、SPIRV-Tools、LLVM-SPIRV 翻译器），并构建了项目（在构建选项中指定了 LLVM 的所有各种插件/项目）。

我遇到的问题是，即使我可以使用 clang 和 LLVM-SPIRV Translator 构建一个 SPIRV 文件，它作为一个程序使用

clCreateProgramWithIL()

成功加载，它在尝试获取内核时返回 CL_INVALID_KERNEL_NAME (-46)在我的 Intel Iris Xe 显卡上使用

clCreateKernel()

。

我可以用 clang 将这个简单的计算内核编译成 LLVM 二进制文件，并且可以使用 LLVM-SPRIV Translator 无错误地输出 spirv 文件：

OpenCL源代码：

__kernel void work(__global int * ptr) {
    size_t id = get_global_id(0);

    ptr[id] = (int)id;
}

这里是 LLVM 转换为 SPIRV 文件之前的可读版本：

; ModuleID = 'test_cl_spv.cl'
source_filename = "test_cl_spv.cl"
target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
target triple = "spir"

; Function Attrs: convergent noinline norecurse nounwind optnone
define dso_local spir_kernel void @work(ptr addrspace(1) noundef align 4 %0) #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
  %2 = alloca ptr addrspace(1), align 4
  %3 = alloca i32, align 4
  store ptr addrspace(1) %0, ptr %2, align 4
  %4 = call spir_func i32 @_Z13get_global_idj(i32 noundef 0) #2
  store i32 %4, ptr %3, align 4
  %5 = load ptr addrspace(1), ptr %2, align 4
  %6 = load i32, ptr %3, align 4
  %7 = getelementptr inbounds i32, ptr addrspace(1) %5, i32 %6
  store i32 7, ptr addrspace(1) %7, align 4
  ret void
}

; Function Attrs: convergent nounwind willreturn memory(none)
declare dso_local spir_func i32 @_Z13get_global_idj(i32 noundef) #1

attributes #0 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
attributes #1 = { convergent nounwind willreturn memory(none) "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #2 = { convergent nounwind willreturn memory(none) }

!llvm.module.flags = !{!0, !1}
!opencl.ocl.version = !{!2}
!opencl.spir.version = !{!2}
!llvm.ident = !{!3}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 7, !"frame-pointer", i32 2}
!2 = !{i32 1, i32 2}
!3 = !{!"clang version 17.0.0 (https://github.com/llvm/llvm-project.git be83a4b257c8f0dfd74a659261a544483c5df9af)"}
!4 = !{i32 1}
!5 = !{!"none"}
!6 = !{!"int*"}
!7 = !{!""}

虽然不完全相同，但它似乎与 Godbolt 上的输出相似。当我用

clCreateProgramWithSource()

直接加载源文件时，它能够找到内核并成功执行计算着色器。

我做了一个小例子来说明这个问题：

void compute_test::Run(Platform pltform, Device device)
{
    std::string file_path = "C:/Users/jdrurka1/source/repos/Dynamics-io/Dynamics.io-Testbench/CPP_Bench/shaders/OpenCL/test_cl_spv.spv";

    cl_platform_id platform_id = (cl_platform_id)pltform.platform;

    // context properties list - must be terminated with 0
    cl_context_properties properties[3];
    properties[0] = CL_CONTEXT_PLATFORM;
    properties[1] = (cl_context_properties)platform_id;
    properties[2] = 0;

    cl_device_id deviceID = (cl_device_id)device.cl_device;

    cl_int err;

    cl_context context = clCreateContext(properties, 1, &deviceID, NULL, NULL, &err);
    cl_command_queue command_queue = clCreateCommandQueue(context, deviceID, 0, &err);


    //open file
    std::ifstream infile(file_path, std::ios::binary);
    char* buffer;
    //get length of file
    infile.seekg(0, infile.end);
    size_t length = infile.tellg();
    infile.seekg(0, infile.beg);
    if (length == 0)
        return;

    //read file
    printf("Reading binary file of length %i\n", (int)length);
    buffer = new char[length];
    infile.read(buffer, length);

    cl_program program = clCreateProgramWithIL(context, buffer, length, &err);
    printf("CreateProgramWithIL: res %i\n", err);

    std::string args;
    cl_int build_res = clBuildProgram(program, 1, &deviceID, args.c_str(), NULL, NULL);
    printf("BuildProgram: res %i\n", build_res);

    
    cl_kernel kernel = clCreateKernel(program, "work", &err);
    printf("Create kernel '%s': %i\n", "work", err);

    char errorStr[1000];
    size_t e_size = 0;
    int res = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_LOG, 1000, errorStr, &e_size);
    printf("LOG (%i): %s\n", e_size, errorStr);
}

产生这个控制台输出（程序打印平台/设备并在调用 Run() 之前选择一个设备。在这种情况下它选择英特尔设备）：

Platform NVIDIA CUDA:
        NVIDIA Corporation - NVIDIA T500
Platform Intel(R) OpenCL HD Graphics:
        Intel(R) Corporation - Intel(R) Iris(R) Xe Graphics
Reading binary file of length 900
CreateProgramWithIL: res 0
BuildProgram: res 0
Create kernel 'work': -46 (CL_INVALID_KERNEL_NAME)
LOG (1):

我不确定我在这里缺少什么让 OpenCL 加载内核。当我查看扩展时，spirv 扩展存在。我尝试在文件上运行 spirv-opt 和 spirv-link，结果相同。

任何指导表示赞赏。