c - OpenCl共享虚拟内存测试奇怪的结果

Question

我正在尝试测试新的 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 正常。但是，至少要以正确的方式打印数组。

有什么想法如何获得如此奇怪的输出吗？

Answer 1

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

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