优化主机到 GPU 的传输答案

【问题标题】：Optimising Host to GPU transfer优化主机到 GPU 的传输
【发布时间】：2014-12-21 13:56:34
【问题描述】：

我正在使用 OpenCL（矩阵乘法的一种变体）将工作卸载到 GPU。矩阵代码本身工作得非常好，但是将数据移动到 GPU 的成本是令人望而却步的。

我已从使用 clEnqueueRead/clEnqueueWrite 转移到内存映射缓冲区，如下所示：

d_a  = clCreateBuffer(context,  CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR,
                    sizeof(char) * queryVector_size,
                    NULL, NULL);
checkErr(err,"Buf A");

d_b  = clCreateBuffer(context,  CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR,
                    sizeof(char) * segment_size,
                     NULL, NULL);

checkErr(err,"Buf B");




err  = clSetKernelArg(ko_smat, 0, sizeof(cl_mem), &d_c);
checkErr(err,"Compute Kernel");
err = clSetKernelArg(ko_smat, 1, sizeof(cl_mem), &d_a);
checkErr(err,"Compute Kernel");
err = clSetKernelArg(ko_smat, 2, sizeof(cl_mem), &d_b);
checkErr(err,"Compute Kernel");

  query_vector = (char*) clEnqueueMapBuffer(commands, d_a, CL_TRUE,CL_MAP_READ, 0, sizeof(char) * queryVector_size, 0, NULL, NULL, &err);
 checkErr(err,"Write A");

 segment_data = (char*) clEnqueueMapBuffer(commands, d_b, CL_TRUE,CL_MAP_READ, 0, sizeof(char) * segment_size, 0, NULL, NULL, &err);
    checkErr(err,"Write B");

     // code which initialises buffers using ptrs (segment_data and queryV)

  err = clEnqueueUnmapMemObject(commands,
                             d_a,
                      query_vector, 0, NULL, NULL);
 checkErr(err,"Unmap Buffer");

  err = clEnqueueUnmapMemObject(commands,
                       d_b,
                      segment_data, 0, NULL, NULL);
 checkErr(err,"Unmap Buff");
 err = clEnqueueNDRangeKernel(commands, ko_smat, 2, NULL, globalWorkItems, localWorkItems, 0, NULL, NULL);

 err = clFinish(commands);
 checkErr(err, "Execute Kernel");

     result = (char*) clEnqueueMapBuffer(commands, d_c, CL_TRUE,CL_MAP_WRITE, 0, sizeof(char) * result_size, 0, NULL, NULL, &err);
     checkErr(err,"Write C");

  printMatrix(result, result_row, result_col);

当我使用 ReadEnqueue/WriteEnqueue 方法并通过它初始化 d_a、d_b、d_c 时，此代码工作正常，但是当我使用 MappedBuffers 时，由于 d_a 和 d_b 为空，结果为 0 运行内核时。

映射/取消映射缓冲区的适当方法是什么？

编辑：核心问题似乎来自这里

  segment_data = (char*) clEnqueueMapBuffer(commands, d_b, CL_TRUE,CL_MAP_READ, 0, sizeof(char) * segment_width * segment_length, 0, NULL, NULL, &err);

  // INITIALISE

  printMatrix(segment_data, segment_length, segment_width);

  // ALL GOOD    

   err = clEnqueueUnmapMemObject(commands,
                           d_b,
                          segment_data, 0, NULL, NULL);
  checkErr(err,"Unmap Buff");

   segment_data = (char*) clEnqueueMapBuffer(commands, d_b, CL_TRUE,CL_MAP_READ, 0, sizeof(char) * segment_width * segment_length, 0\
, NULL, NULL, &err);

   printMatrix(segment_data, segment_length, segment_width);

   // ALL ZEROs again

第一个 printMatrix() 返回正确的输出，一旦我取消映射并重新映射它，segment_data 就变成全 0（它是初始值）。我怀疑我在某处使用了不正确的标志？我不知道在哪里。

【问题讨论】：

标签： opencl gpu

【解决方案1】：

来自 OpenCL 1.2 规范：

5.4.3 访问内存对象的映射区域

...

如果内存对象当前被映射以供读取，则应用程序必须确保在写入此内存对象或其任何关联内存对象（子缓冲区或 1D 图像缓冲区）的任何入队内核或命令之前取消映射内存对象对象）或其父对象（如果内存对象是子缓冲区或一维图像缓冲区对象）开始执行；否则行为未定义。

因此，您需要在将内核入队之后映射results 缓冲区。同样，您需要在将内核排入队列之前取消映射输入缓冲区。映射/取消映射缓冲区的时间线大致如下：

Create input buffers
Create output buffers
Map input buffers
Write input data
Unmap input buffers
Enqueue kernel
Map output buffers
Read output data
Unmap output buffers

【讨论】：

同意。此外，如果您有多个内核要运行，另一个胜利是将数据传输与内核计算重叠。许多高端 GPU 具有双 DMA 引擎，可以同时进行上传、下载和计算。通过这样的重叠操作，您只需为最昂贵的操作付费。
当我这样做时（这是我最初所做的），我系统地将输出设为 null
@user1018513 你的意思是clEnqueueMapBuffer 返回NULL？如果是，它返回的错误代码是什么？
我已经编辑了上面的代码。抱歉措辞不佳。我的内核在任何地方都将 d_a 和 d_b 视为 0（我在取消映射之前打印出 query_vector 和 segment_data 的内容，并且它们已正确初始化），这反过来导致结果为 0。任何地方都没有抛出错误。如果我手动将结果强制为内核中的任意值，我可以读回正确的值。同样，如果我取消映射/重新映射 d_a，那么第二次 d_a 都是 0。

【解决方案2】：

显然，加快代码速度的最佳方法是使用映射缓冲区。您可以使用 CL_MEM_ALLOC_HOST_PTR 创建缓冲区，这基本上可以通过启动 DMA 传输来减轻 CPU 的传输负担。

以下是使用映射缓冲区的示例：

// pointer to hold the result
int * host_ptr = malloc(size * sizeof(int));

d_mem = clCreateBuffer(context,CL_MEM_READ_WRITE|CL_MEM_ALLOC_HOST_PTR,
                       size*sizeof(cl_int), NULL, &ret);

int * map_ptr = clEnqueueMapBuffer(command_queue,d_mem,CL_TRUE,CL_MAP_WRITE,
                                   0,size*sizeof(int),0,NULL,NULL,&ret);
// initialize data
for (i=0; i<size;i++) {
  map_ptr[i] = i;
}

ret = clEnqueueUnmapMemObject(command_queue,d_mem,map_ptr,0,NULL,NULL); 

//Set OpenCL Kernel Parameters
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&d_mem);

size_t global_work[1]  = { size };
//Execute OpenCL Kernel
ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, 
                             global_work, NULL, 0, 0, NULL);

map_ptr = clEnqueueMapBuffer(command_queue,d_mem,CL_TRUE,CL_MAP_READ,
                             0,size*sizeof(int),0,NULL,NULL,&ret);
// copy the data to result array 
for (i=0; i<size;i++){
  host_ptr[i] = map_ptr[i];
} 

ret = clEnqueueUnmapMemObject(command_queue,d_mem,map_ptr,0,NULL,NULL);        

// cl finish etc

取自 this 帖子。

【讨论】：

这是我最初做的，但当我这样做时，map_ptr 的内容系统地为空。

【解决方案3】：

  query_vector = (char*) clEnqueueMapBuffer(commands, d_a, CL_TRUE,CL_MAP_READ, 0, sizeof(char) * queryVector_size, 0, NULL, NULL, &err);
 checkErr(err,"Write A");

 segment_data = (char*) clEnqueueMapBuffer(commands, d_b, CL_TRUE,CL_MAP_READ, 0, sizeof(char) * segment_size, 0, NULL, NULL, &err);
    checkErr(err,"Write B");

缓冲区被映射为 CL_MAP_READ 但写入它们。与缓冲区创建不同，这些标志不采用内存的设备视图，而是主机视图，因此它们应该使用 CL_MAP_WRITE 标志进行映射，否则任何更改都将在未映射时被丢弃

【讨论】：