OpenCL 本地内存和 Xcode

Question

我正在尝试在 Mac 上学习 OpenCL，这似乎与我正在阅读的 OpenCL 书籍在实现上有所不同。我希望能够在 GPU 上动态分配本地内存。我正在阅读的是我需要使用 clSetKernelArg 函数，但这在 Xcode 6.4 中不起作用。这是目前的代码（不要介意这是一个毫无意义的程序，只是想学习共享内存的语法）。在Xcode中，内核被编写为类似于CUDA的独立.cl文件，因此这是一个单独的文件。

add.cl:

kernel void add(int a, int b, global int* c, local int* d)
{
    d[0] = a;
    d[1] = b;
    *c = d[0] + d[1];
}

main.c:

#include <stdio.h>
#include <OpenCL/opencl.h>
#include "add.cl.h"

int main(int argc, const char * argv[]) {

    int a = 3;
    int b = 5;
    int c;
    int* cptr = &c;

    dispatch_queue_t queue = gcl_create_dispatch_queue(CL_DEVICE_TYPE_GPU, NULL);

    void* dev_c = gcl_malloc(sizeof(cl_int), NULL, CL_MEM_WRITE_ONLY);

    // attempt to create local memory buffer
    void* dev_d = gcl_malloc(2*sizeof(cl_int), NULL, CL_MEM_READ_WRITE); 
    // clSetKernelArg(add_kernel, 3, 2*sizeof(cl_int), NULL);

    dispatch_sync(queue, ^{

        cl_ndrange range = { 1, {0, 0, 0}, {1, 0, 0}, {1, 0, 0} };

        // This gives a warning: 
        // Warning: Incompatible pointer to integer conversion passing 'cl_int *' 
        //     (aka 'int *') to parameter of type 'size_t' (aka 'unsigned long')
        add_kernel(&range, a, b, (cl_int*)dev_c, (cl_int*)dev_d);

        gcl_memcpy((void*)cptr, dev_c, sizeof(cl_int));

    });

    printf("%d + %d = %d\n", a, b, c);

    gcl_free(dev_c);    
    dispatch_release(queue);
    return 0;
}

我试过将 clSetKernelArg 放在指示的位置，但它不喜欢第一个参数：

Error: Passing 'void (^)(const cl_ndrange *, cl_int, cl_int, cl_int *, size_t)' to parameter of incompatible type 'cl_kernel' (aka 'struct _cl_kernel *')

我看了又看，但找不到在 Xcode 环境中说明这一点的任何示例。你能给我指出正确的方向吗？

Answer 1

在常规 OpenCL 中，对于声明为本地指针的内核参数，您不会分配主机缓冲区并将其传入（就像您在 dev_d 中所做的那样）。相反，您执行一个 clSetKernelArg，其大小为所需本地存储的大小，但为 NULL 指针（如下所示：clSetKernelArg(kernel, 2, sizeof(cl_int) * local_work_size[0], NULL)）。如果您坚持特定于平台，则必须将其转换为 Xcode 方式。

Answer 2

通过放弃 Apple 的扩展并使用标准的 OpenCL 1.2 调用设法解决了这个问题。这意味着用 clCreateBuffer 替换 gcl_malloc，用 clEnqueueNDRangeKernel 替换 dispatch_sync，最重要的是，在 local 的最后一个参数中使用 clSetKernelArg 和 NULL变量。很有魅力。

这是新版本：

char kernel_add[1024] =
"kernel void add(int a, int b, global int* c, local int* d) \
{\
    d[0] = a;\
    d[1] = b;\
    *c = d[0] + d[1];\
}";

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <OpenCL/opencl.h>

int main(int argc, const char * argv[]) {

    int a = 3;
    int b = 5;
    int c;

    cl_device_id device_id;
    int err = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);

    cl_context context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);    
    cl_command_queue queue = clCreateCommandQueue(context, device_id, 0, &err);

    const char* srccode = kernel;
    cl_program program = clCreateProgramWithSource(context, 1, &srccode, NULL, &err);

    err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    cl_kernel kernel = clCreateKernel(program, "kernel_add", &err);

    cl_mem dev_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int), NULL, NULL);

    err = clSetKernelArg(kernel, 0, sizeof(int), &a);
    err |= clSetKernelArg(kernel, 1, sizeof(int), &b);
    err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &dev_c);
    err |= clSetKernelArg(kernel, 3, sizeof(int), NULL);

    size_t one = 1;
    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &one, NULL, 0, NULL, NULL);
    clFinish(queue);

    err = clEnqueueReadBuffer(queue, dev_c, true, 0, sizeof(int), &c, 0, NULL, NULL);

    clReleaseMemObject(dev_c);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);

    return 0;
}

OpenCL 本地内存和 Xcode

OpenCL Local memory and Xcode

opencl

xcode6