三个clCreateProgramWith函数

函数	是否需要 `clBuildProgram`？
`clCreateProgramWithSource`	是，必须调用。源CL C代码
`clCreateProgramWithBinary`	否，绝对不要调用。可能是二进制形式的SPIRV，或者其他二进制，和具体的编译器有关
`clCreateProgramWithIL`	是，必须调用，因为 IL (SPIR-V) 是中间语言，需要翻译成设备二进制。

clCreateProgramWithIL、clCreateProgramWithSource 和 clCreateProgramWithBinary 是 OpenCL 中用于创建程序对象的三种不同方法。以下是它们的区别和作用：

1. `clCreateProgramWithIL`¶

C++

cl_program clCreateProgramWithIL(
    cl_context context,
    const void *il,
    size_t length,
    cl_int *errcode_ret
);

参数说明¶

context：一个有效的 OpenCL 上下文，指定了程序将在哪个环境中执行。
il：指向包含中间语言代码的内存区域的指针。
length：中间语言代码的字节长度。
errcode_ret：指向一个整数的指针，用于返回函数执行的错误码。如果不需要错误码，可以传入 NULL。

返回值¶

如果函数执行成功，返回一个新创建的 cl_program 对象；如果失败，返回 NULL。

作用: 从中间语言 (Intermediate Language, IL) 创建程序对象。
输入: 接收中间语言的二进制数据（例如 SPIR-V）。
特点:
- 中间语言是一种高级的设备无关格式，可以跨平台使用。
- 使用中间语言允许更灵活地将程序移植到不同的设备上，因为中间语言可以被设备的编译器进一步优化。
- 通常用于更现代的 OpenCL 版本（如 OpenCL 2.1 及以上）。
典型场景:
- 从高层语言（如 OpenCL C 或 OpenCL C++ kernel）编译成中间语言后加载到设备。
- 整合工具链（如使用 LLVM 编译成 SPIR-V）。

示例:

c cl_program program = clCreateProgramWithIL(context, il_binary, il_size, &err);

这个用例假设我们有一个中间语言（IL），例如 SPIR-V 二进制数据。以下代码展示如何从 IL 创建 OpenCL 程序。

C

#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>

// 假设这是你的内核函数名
// 这个名字必须与生成 SPIR-V 二进制文件时使用的 __kernel 函数名完全一致
#define KERNEL_NAME "your_kernel_name" 

int main() {
    // 假设 context 和 device 已经创建
    cl_context context = ...;
    cl_device_id device = ...;
    cl_int err;

    // SPIR-V 二进制数据 (这只是示例，需替换为实际的 SPIR-V 数据)
    // 注意：这个二进制数据必须是从合法的 OpenCL C/C++ 代码编译而来的
    const unsigned char spirv_binary[] = {
        0x03, 0x02, 0x23, 0x07, // SPIR-V header (示例数据)
        /* 后续为实际的 SPIR-V 数据 */
    };
    size_t spirv_size = sizeof(spirv_binary);

    // 1. 从 SPIR-V IL 创建程序对象
    cl_program program = clCreateProgramWithIL(context, spirv_binary, spirv_size, &err);
    if (err != CL_SUCCESS) {
        printf("Failed to create program with IL: %d\n", err);
        return -1;
    }

    // 2. 构建/翻译程序 (将 SPIR-V 翻译为设备二进制)
    err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
    if (err != CL_SUCCESS) {
        printf("Failed to build program: %d\n", err);

        // 输出构建日志以排查问题
        char buffer[2048];
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, NULL);
        printf("Build log:\n%s\n", buffer);

        clReleaseProgram(program); // 清理
        clReleaseContext(context);
        return -1;
    }

    // 3. 从构建好的程序中创建内核对象 (这是缺失的关键步骤)
    cl_kernel kernel = clCreateKernel(program, KERNEL_NAME, &err);
    if (err != CL_SUCCESS) {
        printf("Failed to create kernel '%s': %d\n", KERNEL_NAME, err);
        // 常见错误是内核名字写错，或者 SPIR-V 模块中不包含此内核
        clReleaseProgram(program);
        clReleaseContext(context);
        return -1;
    }

    printf("Program created from IL, built, and kernel '%s' created successfully!\n", KERNEL_NAME);

    // 此时，你可以继续进行下一步操作，例如：
    // - 创建命令队列 (cl_command_queue)
    // - 创建内存对象 (cl_mem)
    // - 设置内核参数 (clSetKernelArg)
    // - 执行内核 (clEnqueueNDRangeKernel)

    // 4. 清理所有创建的 OpenCL 资源 (按相反顺序释放)
    clReleaseKernel(kernel);   // 释放内核对象
    clReleaseProgram(program); // 释放程序对象
    clReleaseContext(context); // 释放上下文

    return 0;
}

2. `clCreateProgramWithSource`¶

C++

cl_program clCreateProgramWithSource(
    cl_context context,
    cl_uint count,
    const char **strings,
    const size_t *lengths,
    cl_int *errcode_ret
);

参数说明¶

context：一个有效的 OpenCL 上下文，指定了程序将在哪个环境中执行。
count：strings 数组中的字符串数量。
strings：指向包含 OpenCL 程序源代码的字符串数组的指针。
lengths：指向一个数组的指针，该数组包含 strings 数组中每个字符串的长度。如果传入 NULL，则函数会自动计算字符串的长度。
errcode_ret：指向一个整数的指针，用于返回函数执行的错误码。如果不需要错误码，可以传入 NULL。

返回值¶

如果函数执行成功，返回一个新创建的 cl_program 对象；如果失败，返回 NULL。

作用: 从源代码字符串创建程序对象。
输入: 接收以字符串形式表示的 OpenCL C 内核源代码。
特点:
- 这是最常见的创建程序的方法，尤其是在开发和调试阶段。
- 源代码会在运行时由 OpenCL 平台的编译器编译成设备可执行的二进制代码。
- 编译过程可能会因目标设备的不同而有所不同。
典型场景:
- 开发阶段，直接加载内核的源代码。
- 动态生成内核代码。

示例:

c const char* source = "__kernel void my_kernel(...) { ... }"; cl_program program = clCreateProgramWithSource(context, 1, &source, NULL, &err);

这是最常见的用法，直接从 OpenCL 内核的源代码字符串创建程序对象。

C

#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>

int main() {
    // 假设 context 和 device 已经创建
    cl_context context = ...;
    cl_device_id device = ...;
    cl_int err;

    // OpenCL 内核源代码
    const char* source = 
        "__kernel void vector_add(__global const float* a, __global const float* b, __global float* c) {"
        "   int id = get_global_id(0);"
        "   c[id] = a[id] + b[id];"
        "}";

    // 1. 创建程序对象
    cl_program program = clCreateProgramWithSource(context, 1, &source, NULL, &err);
    if (err != CL_SUCCESS) {
        printf("Failed to create program with source: %d\n", err);
        return -1;
    }

    // 2. 构建程序
    err = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
    if (err != CL_SUCCESS) {
        printf("Failed to build program: %d\n", err);
        char buffer[2048];
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, NULL);
        printf("Build log:\n%s\n", buffer);
        clReleaseProgram(program); // 清理
        clReleaseContext(context);
        return -1;
    }

    // 3. 从构建好的程序中创建内核对象 (这是缺失的关键步骤)
    //    内核的名字 "vector_add" 必须与源代码中的 __kernel 函数名完全一致
    cl_kernel kernel = clCreateKernel(program, "vector_add", &err);
    if (err != CL_SUCCESS) {
        printf("Failed to create kernel 'vector_add': %d\n", err);
        // 常见错误是内核名字写错，或者构建失败
        clReleaseProgram(program);
        clReleaseContext(context);
        return -1;
    }

    printf("Program created, built, and kernel 'vector_add' created successfully!\n");

    // 此时，你可以继续进行下一步操作，例如：
    // - 创建命令队列 (cl_command_queue)
    // - 创建内存对象 (cl_mem) 用于 a, b, c 数组
    // - 设置内核参数 (clSetKernelArg)
    // - 执行内核 (clEnqueueNDRangeKernel)
    // - 读取结果 (clEnqueueReadBuffer)

    // 4. 清理所有创建的 OpenCL 资源 (按相反顺序释放)
    clReleaseKernel(kernel);   // 释放内核对象
    clReleaseProgram(program); // 释放程序对象
    clReleaseContext(context); // 释放上下文

    return 0;
}

3. `clCreateProgramWithBinary`¶

C++

cl_program clCreateProgramWithBinary(
    cl_context context,
    cl_uint num_devices,
    const cl_device_id *device_list,
    const size_t *lengths,
    const unsigned char **binaries,
    cl_int *binary_status,
    cl_int *errcode_ret
);

参数说明¶

context：一个有效的 OpenCL 上下文，指定了程序将在哪个环境中执行。
num_devices：device_list 数组中的设备数量。
device_list：指向包含设备 ID 的数组的指针，这些设备将使用该二进制程序。
lengths：指向一个数组的指针，该数组包含 binaries 数组中每个二进制代码的长度。
binaries：指向包含预编译二进制代码的数组的指针。
binary_status：指向一个整数数组的指针，用于返回每个设备的二进制代码加载状态。如果不需要状态信息，可以传入 NULL。
errcode_ret：指向一个整数的指针，用于返回函数执行的错误码。如果不需要错误码，可以传入 NULL。

返回值¶

如果函数执行成功，返回一个新创建的 cl_program 对象；如果失败，返回 NULL。

作用: 从已编译的二进制文件（设备特定的二进制）创建程序对象。
输入: 接收程序的设备特定二进制文件。
特点:
- 二进制文件是在某设备上提前编译的（通过 clGetProgramInfo 获取或离线工具生成）。
- 适用于运行时加载预编译的程序，避免运行时编译，从而减少启动时间。
- 设备二进制文件是特定于设备的，不能跨平台使用。
典型场景:
- 部署阶段，直接加载之前已编译的二进制文件。
- 在运行时减少编译时间。

示例:

c size_t binary_size = ...; unsigned char* binary = ...; cl_program program = clCreateProgramWithBinary(context, 1, &device, &binary_size, (const unsigned char**)&binary, NULL, &err);

这个用例展示如何从设备特定的二进制文件创建程序对象。二进制文件通常通过离线编译工具生成，或者通过 clGetProgramInfo 从已有的程序提取。

C

#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>

// 假设这是你的内核函数名
#define KERNEL_NAME "your_kernel_name" 

int main() {
    // 假设 context 和 device 已经创建
    cl_context context = ...;
    cl_device_id device = ...;
    cl_int err;

    // 从文件读取二进制数据
    FILE* file = fopen("kernel.bin", "rb");
    if (!file) {
        printf("Failed to open binary file.\n");
        return -1;
    }

    fseek(file, 0, SEEK_END);
    size_t binary_size = ftell(file);
    rewind(file);

    unsigned char* binary = (unsigned char*)malloc(binary_size);
    fread(binary, 1, binary_size, file);
    fclose(file);

    // 创建程序对象
    cl_program program = clCreateProgramWithBinary(context, 1, &device, &binary_size, 
                                                   (const unsigned char**)&binary, NULL, &err);
    free(binary); // 二进制数据已被驱动加载，本地内存可以释放

    if (err != CL_SUCCESS) {
        printf("Failed to create program with binary: %d\n", err);
        return -1;
    }

    // --- clBuildProgram 调用已被删除 ---

    // 直接从已加载的二进制程序创建内核
    cl_kernel kernel = clCreateKernel(program, KERNEL_NAME, &err);
    if (err != CL_SUCCESS) {
        printf("Failed to create kernel '%s': %d\n", KERNEL_NAME, err);
        // 注意：这里不能获取构建日志，因为没有执行构建操作
        clReleaseProgram(program);
        clReleaseContext(context);
        return -1;
    }

    printf("Program created from binary and kernel '%s' created successfully!\n", KERNEL_NAME);

    // 清理资源
    clReleaseKernel(kernel);   // 释放内核对象
    clReleaseProgram(program); // 释放程序对象
    clReleaseContext(context); // 释放上下文

    return 0;
}

总结对比¶

clCreateProgramWithIL 示例:
假设使用 SPIR-V 数据作为输入。
必须确保设备支持中间语言（如 SPIR-V），通常需要 OpenCL 2.1 或更高版本。
clCreateProgramWithSource 示例:
最常见的用法，从源码字符串直接创建和编译。
适合开发阶段和动态代码生成。
clCreateProgramWithBinary 示例:
需要准备设备特定的二进制文件，文件可以通过离线工具生成或通过 clGetProgramInfo 提取。
适合部署阶段，避免运行时编译。

方法	输入类型	优点	缺点	适用场景
`clCreateProgramWithIL`	中间语言（如 SPIR-V）	可移植性和优化性强，支持现代 OpenCL 特性	需要支持 IL 的设备和工具链	高级优化、跨平台开发
`clCreateProgramWithSource`	源代码字符串	开发和调试方便，动态生成代码灵活	运行时编译耗时，性能模型依赖设备	开发和调试，动态代码生成
`clCreateProgramWithBinary`	设备特定二进制	启动快，无需运行时编译	设备依赖性强，无法跨平台	部署阶段，固定设备的应用场景

实际开发中的使用建议¶

开发阶段：
使用 clCreateProgramWithSource，直接加载源代码调试。
优化和跨平台部署：
使用 clCreateProgramWithIL，生成中间语言以提高移植性和效率。
固定设备的高效部署：
使用 clCreateProgramWithBinary，提前编译好设备二进制文件，减少运行时开销。

通过合理选择这些方法，可以在开发效率、跨平台兼容性和运行性能之间找到平衡。

三个clCreateProgramWith函数

1. clCreateProgramWithIL¶

参数说明¶

返回值¶

2. clCreateProgramWithSource¶

参数说明¶

返回值¶

3. clCreateProgramWithBinary¶

参数说明¶

返回值¶

总结对比¶

实际开发中的使用建议¶

1. `clCreateProgramWithIL`¶

2. `clCreateProgramWithSource`¶

3. `clCreateProgramWithBinary`¶