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 放在指定的位置，但它不喜欢第一个参数：

错误：将“void (^)(const cl_ndrange *, cl_int, cl_int, cl_int *, size_t)”传递给不兼容类型“cl_kernel”（又名“struct _cl_kernel *”）的参数

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

Answer 1

通过放弃 Apple 的扩展并使用标准的 OpenCL 1.2 调用来解决这个问题。这意味着用clCreateBuffer 替换gcl_malloc，用clEnqueueNDRangeKernel 替换dispatch_sync，最重要的是，在局部变量的最后一个参数中使用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;
}

Answer 2

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