1

产生错误的内核代码：

__kernel void testDynamic(__global int *data)
{
    int id=get_global_id(0);
    atomic_add(&data[1],2);
}

__kernel void test(__global int * data)
{
    int id=get_global_id(0);
    atomic_add(&data[0],2);
    if (id == 0) {
        queue_t q = get_default_queue();
        ndrange_t ndrange = ndrange_1D(1,1);
        void (^my_block_A)(void) = ^{testDynamic(data);};
        enqueue_kernel(q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
                       ndrange,
                       my_block_A);
    }

}

我测试了下面的代码以确保 OpenCL 2.0 编译器正常工作。

__kernel void test2(__global int *data)
{
    int id=get_global_id(0);
    data[id]=work_group_scan_inclusive_add(id);
}

扫描功能提供 0,1,3,6 作为输出，因此 OpenCL 2.0 缩减功能正在工作。

动态并行是 OpenCL 2.0 的扩展吗？如果我删除enqueue_kernel命令，结果等于预期值（省略子内核）。

设备：AMD RX550，驱动：17.6.2

是否有需要在主机端运行的特殊命令才能在get_default_queue队列上运行子内核？目前，命令队列是使用 OpenCL 1.2 方式创建的，如下所示：

commandQueue = cl::CommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &err);

是否get_default_queue()必须是调用父内核的相同命令队列？问这个是因为我使用相同的命令队列将数据上传到 GPU，然后在一次同步中下载结果。

4

1 回答 1

1

将解决方案从一个问题移到另一个答案：

编辑：下面的 API 命令是解决方案：

commandQueue = cl::CommandQueue(context, device,
  CL_QUEUE_ON_DEVICE|
  CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | 
  CL_QUEUE_ON_DEVICE_DEFAULT, &err);

在创建这个队列（每个设备只有 1 个）之后，没有将它用于其他任何事情，并且父内核也在任何其他主机队列中排队，因此看起来 get_default_queue() 不必是父调用队列。

文档说，如果指定了 CL_QUEUE_ON_DEVICE，则会抛出 CL_INVALID_QUEUE_PROPERTIES，但对于我的机器，动态并行可以使用它并且不会抛出该错误（作为上层 commandQueue 构造函数参数）。

于 2017-07-02T04:26:04.890 回答