OpenCl共享虚拟内存测试结果奇怪答案

【问题标题】：OpenCl shared virtual memory test strange resultsOpenCl共享虚拟内存测试结果奇怪
【发布时间】：2016-03-19 18:36:18
【问题描述】：

我正在尝试测试新的 OpenCl 2.0 SVM 功能。我在具有以下 SVM 功能的机器上使用 AMD-APP-SDK：

CL_DEVICE_SVM_COARSE_GRAIN_BUFFER：1（支持）
CL_DEVICE_SVM_FINE_GRAIN_BUFFER：1（支持）
CL_DEVICE_SVM_FINE_GRAIN_SYSTEM：0（不支持）
CL_DEVICE_SVM_ATOMICS：0（不支持）

我正在使用一个非常简单的示例进行测试，其中我执行以下操作（当然是在设置 OpenCL 对象之后）：

在主机端：

分配一个 SVM 缓冲区。
映射、填充和取消映射该缓冲区。
设置内核。
将 SVM 缓冲区作为参数传递。
调用内核。

以下是相关主机代码：

int status;
cl_int cl_status;
int num_elements = 10;

status = SetupKernel("test_svm_kernel.cl", "test_svm_kernel");
CHECK_ERROR(status, "SetupKernel");

svm_input = clSVMAlloc(context, CL_MEM_READ_WRITE, num_elements*sizeof(int),
                      0);
CHECK_ALLOCATION(svm_input, "svm_input");


cl_status = clEnqueueSVMMap(queue, CL_TRUE,
                            CL_MAP_WRITE_INVALIDATE_REGION,
                            svm_input, num_elements*sizeof(int), 0, NULL,
                            NULL);
CHECK_OPENCL_ERROR(cl_status, "clEnqueueSVMMap");

for(int i=0 ; i<num_elements ; i++)
{
    ((int*)svm_input)[i] = i;
}


for(int i=0 ; i<num_elements ; i++)
{
    std::cout << i << ", " << ((int*)svm_input)[i] << std::endl;
}

cl_status = clEnqueueSVMUnmap(queue, svm_input, 0, NULL, NULL);
CHECK_OPENCL_ERROR(cl_status, "clEnqueueSVMUnmap");

cl_status = clGetKernelWorkGroupInfo(
    kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t),
    &kernel_wrkgrp_size, NULL);
CHECK_OPENCL_ERROR(cl_status, "clGetKernelWorkGroupInfo");

cl_status = clGetKernelWorkGroupInfo(
    kernel, device, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(cl_ulong),
    &compile_wrkgrp_size, NULL);
CHECK_OPENCL_ERROR(cl_status, "clGetKernelWorkGroupInfo");

cl_status = clGetKernelWorkGroupInfo(
    kernel, device, CL_KERNEL_COMPILE_WORK_GROUP_SIZE,
    sizeof(size_t)*3, &compile_wrkgrp_size, NULL);
CHECK_OPENCL_ERROR(cl_status, "clGetKernelWorkGroupInfo");

size_t local_threads = 1;//kernel_wrkgrp_size;
size_t globl_threads = num_elements;

cl_status = clSetKernelArgSVMPointer(kernel, 0, (void*)(svm_input));
CHECK_OPENCL_ERROR(cl_status, "clSetKernelArgSVMPointer");

cl_event ndr_event;
cl_status = clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
                                   &globl_threads, NULL,
                                   0, NULL, &ndr_event);
CHECK_OPENCL_ERROR(cl_status, "clEnqueueNDRangeKernel");

cl_status = clFlush(queue);
CHECK_OPENCL_ERROR(cl_status, "clFlush");

内核方面：内核非常简单。

__kernel void test_svm_kernel(__global void* input)
{
    __global int* intInput = (__global int*)(input);

    int idx = get_global_id(0);
    printf("intInput[%d] = %d\n", idx, intInput[idx]);
}

我在主机中得到的输出是：

Host: input[0] =  0
Host: input[1] =  1
Host: input[2] =  2
Host: input[3] =  3
Host: input[4] =  4
Host: input[5] =  5
Host: input[6] =  6
Host: input[7] =  7
Host: input[8] =  8
Host: input[9] =  9

这是自然的预期输出。

在内核上，我得到奇怪的输出（有时会发生变化）：

input[0] = 0
input[1] = 2
input[3] = 1
input[5] = 5
input[1] = 7
input[8] = 1
input[0] = 0
input[0] = 0
input[0] = 0
input[0] = 0

我不希望设备上的 printf 正常。但是，至少要以正确的方式打印数组。

任何想法如何获得如此奇怪的输出？

【问题讨论】：

为什么你要传递给内核一个空指针然后你转换它？此外，您应该确保主机和设备中的 int 和 int 匹配。因为 int 可以是 16/32/64，具体取决于系统。虽然 cl_int（你应该使用的）总是 32 位。
检查在使用 OpenCL 1.x 时是否得到正确的输出。还要确保您使用最新的驱动程序 - 我遇到了类似的奇怪问题，后来通过驱动程序更新修复。
一般来说，使用 printf 进行检查是有风险的，因为它是完全调试和实验的（它也破坏了 OpenCL 执行模型[完全并行不同步]）。检查输出是一个更好的主意。例如，也许编译器优化了所有内核，因为您没有在任何地方写入输入数据。
谢谢你们的cmets。事实证明 printf 不可靠，当我更改输入数组并将其映射回主机时，内容完全正确。

标签： c kernel opencl shared-memory gpgpu

【解决方案1】：

一般来说，使用 printf 进行检查是有风险的，因为它是完全调试和实验的（它也破坏了 OpenCL 执行模型[完全并行不同步]）。检查输出是一个更好的主意。例如，也许编译器优化了所有内核，因为您没有在任何地方写入输入数据。 – DarkZeros

事实证明 printf 不可靠，当我更改输入数组并将其映射回主机时，内容完全正确。 – 卡里姆·厄高维

【讨论】：