OpenCL kernel 函数中包含多个源文件

写 OpenCL 的 kernel 函数时，是利用原来在 CPU 上已经验证过的代码，往往是由多个源文件组成的，那么怎么实现包含多个源文件？

上网搜索，https://www.cnblogs.com/willhua/p/13362837.html 写了方法，我这里写个简单的例子：

// main.c
    cl_int errNum;
    cl_program program;
    size_t program_length;

    char *const source = ReadKernelSourceFile ("vecAdd.cl", &program_length);    // 读取 kernel 函数源代码，此处无需被包含源文件

    program = clCreateProgramWithSource(context, 1, (const char **)&source, NULL, NULL);
    if (program == NULL)
    {
        printf("Failed to create CL program from source.\n" );
        return NULL;
    }

    errNum = clBuildProgram(program, 0, NULL, "-I ../opencl_test/", NULL, NULL);  // 这句是关键
    if (errNum != CL_SUCCESS)
    {
        char buildLog[16384];
        clGetProgramBuildInfo(program, device,
                              CL_PROGRAM_BUILD_LOG,
                              sizeof(buildLog),
                              buildLog, NULL);
        printf("Error in kernel:%s \n", buildLog);
        clReleaseProgram(program);
        return NULL;
    }

关键就在 errNum = clBuildProgram(program, 0, NULL, "-I ../opencl_test/", NULL, NULL) 的 "-I ../opencl_test/"，这是被包含文件存放的目录。

kernel 为：

// vecAdd.cl

#include <test.cl>  // 包含文件

__kernel void vector_add(__global const float *a,
                         __global const float *b,
                         __global float *result)
{
    int gid = get_global_id(0);
//    result[gid] = a[gid] * b[gid];
    result[gid] = test(a[gid] , b[gid]);
}

包含的 test.cl 为：

// test.cl

#include <test1.cl>  // 包含文件

float test(float a, float b)
{
    return test1(a,b);
}

再包含的 test1.cl 为：

// test1.cl

float test1(float a, float b)
{
    return a*b;
}

这样的话，kernel 函数就可以组织为多个源文件，便于理解和调试。