为什么 clEnqueueMapBuffer 返回随机数据？答案

【问题标题】：Why is clEnqueueMapBuffer returning random data?为什么 clEnqueueMapBuffer 返回随机数据？
【发布时间】：2014-05-02 22:51:39
【问题描述】：

我正在尝试通过编写一个内核来了解 OpenCL 中的 clEnqueueMapBuffer，该内核在输入缓冲区中查找值的平方，但使用 clEnqueueMapBuffer 在输出缓冲区中一次只返回两个项目。据我了解，此函数在主机内存中返回一个指针，该指针指向设备中的缓冲内存。然后 clEnqueueUnmapMemObject 必须取消映射此缓冲区以允许内核继续其计算。现在，当我调用 clEnqueueMapBuffer 时，它正在返回随机数据。

这是我的内核

__kernel void testStream(
    __global int *input_vector,
    __global int *output_vector,
    __global int *mem_flag) // informs the host when the workload is finished
{
    mem_flag[0] = 1;
}

我的来源

#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#include <CL/opencl.h>

#include "utils.h"

int main(void)
{
    unsigned int n = 24;
    int BUFF_SIZE = 2;
    // Input and output vectors
    int num_bytes = sizeof(int) * n;
    int *output_buffer = (int *) malloc(num_bytes);
    int output_buffer_offset = 0;
    int *mapped_data = NULL;
    // use mapped_flag for determining if the job on the device is finished
    int *mapped_flag = NULL;
    int *host_in = (int *) malloc(num_bytes);
    int *host_out = (int *) malloc(num_bytes);
    // Declare cl variables
    cl_mem device_in;
    cl_mem device_out;
    cl_mem device_out_flag;

    // Declare cl boilerplate
    cl_platform_id   platform = NULL;
    cl_device_id     device   = NULL;
    cl_command_queue queue    = NULL;
    cl_context       context  = NULL;
    cl_program       program  = NULL;
    cl_kernel        kernel   = NULL;

    // Located in utils.c -- the source is irrelevant here
    char *kernel_source = read_kernel("kernels/test.cl");

    // Initialize host_in
    int i;
    for (i = 0; i < n; i++) {
        host_in[i] = i + 1;
    }

    // Set up opencl
    cl_int error;
    error   = clGetPlatformIDs(1, &platform, NULL);
    printf("clGetPlatformIDs: %d\n", (int) error);
    error   = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, 1, &device, NULL);
    printf("clGetDeviceIDs: %d\n", (int) error);
    context = clCreateContext(NULL, 1, &device, NULL, NULL, &error);
    printf("clCreateContext: %d\n", (int) error);
    queue   = clCreateCommandQueue(context, device, 0, &error);
    printf("clCreateCommandQueue: %d\n", (int) error);
    program = clCreateProgramWithSource(context, 1,
        (const char**)&kernel_source, NULL, &error);
    printf("clCreateProgramWithSource: %d\n", (int) error);
    clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    kernel = clCreateKernel(program, "testStream", &error);
    printf("clCreateKernel: %d\n", (int) error);
    // Create the buffers
    device_in = clCreateBuffer(context, CL_MEM_READ_ONLY,
        num_bytes, NULL, NULL);
    device_out = clCreateBuffer(context,
        CL_MEM_WRITE_ONLY,
        sizeof(int) * BUFF_SIZE, NULL, NULL);
    device_out_flag = clCreateBuffer(context,
        CL_MEM_WRITE_ONLY,
        sizeof(int) * 2, NULL, NULL);

    // Write the input buffer
    clEnqueueWriteBuffer(
        queue, device_in, CL_FALSE, 0, num_bytes,
        host_in, 0, NULL, NULL);

    // Set the kernel arguments
    error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &device_in);
    error = clSetKernelArg(kernel, 1, sizeof(cl_mem), &device_out);
    error = clSetKernelArg(kernel, 2, sizeof(cl_mem), &device_out_flag);

    // Execute the kernel over the entire range of data
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, 
        (const size_t *) &n, NULL, 0, NULL, NULL);

    // Map and unmap until the flag is set to true
    int break_flag = 0;
    while(1) {
        // Map the buffers
        mapped_data = (int *) clEnqueueMapBuffer(
            queue, device_out, CL_TRUE, CL_MAP_READ, 0,
            sizeof(int) * BUFF_SIZE, 0, NULL, NULL, &error);
        mapped_flag = (int *) clEnqueueMapBuffer(
            queue, device_out_flag, CL_TRUE, CL_MAP_READ, 0,
            sizeof(int) , 0, NULL,NULL, &error);
        // Extract the data out of the buffer
        printf("mapped_flag[0] = %d\n", mapped_flag[0]);
        // Set the break_flag
        break_flag = mapped_flag[0];
        // Unmap the buffers
        error = clEnqueueUnmapMemObject(queue, device_out, mapped_data, 0,
            NULL, NULL);
        error = clEnqueueUnmapMemObject(queue, device_out_flag, mapped_flag,
            0, NULL, NULL);
        if (break_flag == 1) {break;}
        usleep(1000*1000);
    }

   return 0;
}

当我运行程序时，我得到类似于

的输出

clGetPlatformIDs: 0
clGetDeviceIDs: 0
clCreateContext: 0
clCreateCommandQueue: 0
clCreateProgramWithSource: 0
clCreateKernel: 0
mapped_flag[0] = 45366144
mapped_flag[0] = 45366144
mapped_flag[0] = 45366144
mapped_flag[0] = 45366144
mapped_flag[0] = 45366144

为什么会这样？

编辑

我在带有 Fedora 19 64 位内核 3.13.7-100.fc19.x86_64 的 HP dm1z 上运行此代码。这是 clinfo 的输出

Number of platforms:                 1
  Platform Profile:              FULL_PROFILE
  Platform Version:              OpenCL 1.2 AMD-APP (1214.3)
  Platform Name:                 AMD Accelerated Parallel Processing
  Platform Vendor:               Advanced Micro Devices, Inc.
  Platform Extensions:               cl_khr_icd cl_amd_event_callback cl_amd_offline_devices


  Platform Name:                 AMD Accelerated Parallel Processing
Number of devices:               2
  Device Type:                   CL_DEVICE_TYPE_GPU
  Device ID:                     4098
  Board name:                    AMD Radeon HD 6310 Graphics
  Device Topology:               PCI[ B#0, D#1, F#0 ]
  Max compute units:                 2
  Max work items dimensions:             3
    Max work items[0]:               256
    Max work items[1]:               256
    Max work items[2]:               256
  Max work group size:               256
  Preferred vector width char:           16
  Preferred vector width short:          8
  Preferred vector width int:            4
  Preferred vector width long:           2
  Preferred vector width float:          4
  Preferred vector width double:         0
  Native vector width char:          16
  Native vector width short:             8
  Native vector width int:           4
  Native vector width long:          2
  Native vector width float:             4
  Native vector width double:            0
  Max clock frequency:               492Mhz
  Address bits:                  32
  Max memory allocation:             134217728
  Image support:                 Yes
  Max number of images read arguments:       128
  Max number of images write arguments:      8
  Max image 2D width:                16384
  Max image 2D height:               16384
  Max image 3D width:                2048
  Max image 3D height:               2048
  Max image 3D depth:                2048
  Max samplers within kernel:            16
  Max size of kernel argument:           1024
  Alignment (bits) of base address:      2048
  Minimum alignment (bytes) for any datatype:    128
  Single precision floating point capability
    Denorms:                     No
    Quiet NaNs:                  Yes
    Round to nearest even:           Yes
    Round to zero:               Yes
    Round to +ve and infinity:           Yes
    IEEE754-2008 fused multiply-add:         Yes
  Cache type:                    None
  Cache line size:               0
  Cache size:                    0
  Global memory size:                201326592
  Constant buffer size:              65536
  Max number of constant args:           8
  Local memory type:                 Scratchpad
  Local memory size:                 32768
  Kernel Preferred work group size multiple:     32
  Error correction support:          0
  Unified memory for Host and Device:        1
  Profiling timer resolution:            1
  Device endianess:              Little
  Available:                     Yes
  Compiler available:                Yes
  Execution capabilities:                
    Execute OpenCL kernels:          Yes
    Execute native function:             No
  Queue properties:              
    Out-of-Order:                No
    Profiling :                  Yes
  Platform ID:                   0x00007fd434852fc0
  Name:                      Loveland
  Vendor:                    Advanced Micro Devices, Inc.
  Device OpenCL C version:           OpenCL C 1.2 
  Driver version:                1214.3
  Profile:                   FULL_PROFILE
  Version:                   OpenCL 1.2 AMD-APP (1214.3)
  Extensions:                    cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_3d_image_writes cl_khr_byte_addressable_store cl_khr_gl_sharing cl_ext_atomic_counters_32 cl_amd_device_attribute_query cl_amd_vec3 cl_amd_printf cl_amd_media_ops cl_amd_media_ops2 cl_amd_popcnt cl_amd_image2d_from_buffer_read_only

另外，值得注意的是，当我开始使用 OpenCL 时，我运行了一个测试程序来计算内积，但结果很奇怪。最初我虽然这是程序错误并忘记了它，但 OpenCL 实现是否可能有问题？如果有帮助，OpenGL 实现有多个错误，导致随机数据块显示在我的桌面背景上，但这也可能是 Linux 问题。

【问题讨论】：

标签： c memory opencl

【解决方案1】：

您将 NULL 作为全局工作大小传递给您的 clEnqueueNDRangeKernel 调用：

// Execute the kernel over the entire range of data
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, NULL, NULL, 0, NULL, NULL);

如果您正在检查此调用返回的错误代码（您总是应该这样做），您会得到与CL_INVALID_GLOBAL_WORK_SIZE 对应的错误代码。您始终需要指定全局工作大小，因此您的调用应如下所示：

// Execute the kernel over the entire range of data
size_t global[1] = {1};
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global, NULL, 0, NULL, NULL);
// check error == CL_SUCCESS!

您对映射和取消映射缓冲区的调用很好；我已经用上述修复测试了这段代码，它对我有用。

您修复上述问题的更新代码如下所示：

unsigned int n = 24;
...
// Execute the kernel over the entire range of data
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, 
    (const size_t *) &n, NULL, 0, NULL, NULL);

这不是将全局工作大小参数传递给内核的安全方法。例如，unsigned int n 变量可能占用 32 位，而 size_t 可能占用 64 位。这意味着，当您传递n 的地址并转换为const size_t* 时，实现将读取一个64 位的值，该值将包含n 的32 位加上其他32 个具有任意值的位。您应该将 n 分配给 size_t 变量，然后再将其传递给 clEnqueueNDRangeKernel，或者将其更改为 size_t 本身。

这可能与您遇到的问题有关，也可能无关。例如，您可能会意外启动大量个工作项，这可以解释为什么代码似乎在 CPU 上阻塞。

【讨论】：

您对 clEnqueueNDRangeKernel 没有限制是正确的，但我已修复此错误。此外，对 clEnqueueMapBuffer 和 clEnqueueUnmapMemObject 的调用都返回 CL_SUCCESS。但是，我仍然收到垃圾数据，通常是 -1，但有时是其他负数。
@user1876508 我已经在许多不同的平台上对此进行了测试，并且每次都能正常工作。你能给我们一些关于你的系统的信息 - 哪个设备、操作系统、SDK 版本等？您可以尝试用clEnqueueWriteBuffer 填充缓冲区，看看问题是否出在内核调用上。
你应该小心类型，对 cl 类型使用 cl_int 并重试。（以防万一，但我不明白为什么它不应该工作）
@jprice 查看我所做的编辑，我包含了有关我的操作系统的信息和 clinfo 的相关输出
@user1876508 clEnqueueMapBuffer 应该阻塞直到映射操作完成，因为您将 CL_TRUE 作为第三个参数传递。这可能表明先前的命令仍在执行（例如内核）。对此有一种可能的解释，我已将其添加到我的答案中。

【解决方案2】：

我想到了以下几点：

我感觉你以某种方式认为调用clEnqueueMapBuffer 会中断内核的执行。我不认为这是正确的。一旦一个命令被启动执行，它就会一直运行，直到它完成（或失败......）。可以同时启动多个命令，但是在内核仍在处理它时尝试读取一些数据将导致未定义的行为。此外，您创建命令队列的方式不会让您同时运行多个命令。您需要在创建队列时使用CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 属性以允许这样做。
我不知道在@jprice 发布后您调用内核的全局大小，但除非您只使用一个工作项执行内核，否则您将无法使用这种语句：mem_flag[0] = 1;，因为所有工作项将写入同一位置。（我猜您只发布了内核的一部分。检查您是否还有其他类似的声明……实际上，如果您发布整个内核代码，那将很有用）。
由于您总是映射和取消映射缓冲区的同一部分，并且总是检查第一个元素（mapped_flag），而且内核此时已完成其计算（见第一点），至少这是正常的你总是读到相同的值。

【讨论】：

您对我的假设也是正确的。 stackoverflow.com/questions/22246965/… 给了我这个策略。
这是您链接到的我的评论，但我从未暗示您可以映射缓冲区在 GPU 使用它时。您似乎认为 CPU 和 GPU 可以共享内存并通过它进行通信，但这不是 OpenCL 的工作方式。它是异步的，内存一次只能由一侧使用。（也就是说，直到具有 SVM 和原子内存访问的 OpenCL 2.0）