C 中的统一内存和流

Question

我正在尝试在 C 中使用具有 CUDA 6 和统一内存的流。我之前的流实现如下所示：

for(x=0; x<DSIZE; x+=N*2){

 gpuErrchk(cudaMemcpyAsync(array_d0, array_h+x, N*sizeof(char), cudaMemcpyHostToDevice, stream0));
 gpuErrchk(cudaMemcpyAsync(array_d1, array_h+x+N, N*sizeof(char), cudaMemcpyHostToDevice, stream1));


gpuErrchk(cudaMemcpyAsync(data_d0, data_h, wrap->size*sizeof(int), cudaMemcpyHostToDevice, stream0));
gpuErrchk(cudaMemcpyAsync(data_d1, data_h, wrap->size*sizeof(int), cudaMemcpyHostToDevice, stream1));

searchGPUModified<<<N/128,128,0,stream0>>>(data_d0, array_d0, out_d0 );
searchGPUModified<<<N/128,128,0,stream1>>>(data_d1, array_d1, out_d1);

gpuErrchk(cudaMemcpyAsync(out_h+x, out_d0 , N * sizeof(int), cudaMemcpyDeviceToHost, stream0));
gpuErrchk(cudaMemcpyAsync(out_h+x+N, out_d1 ,N *  sizeof(int), cudaMemcpyDeviceToHost, stream1));

} 

但我找不到使用相同技术将大块数据发送到 GPU 的流和统一内存的示例。因此我想知道是否有办法做到这一点？

Answer 1

您应该阅读编程指南的section J.2.2（最好是所有附录 J）。

使用统一内存，使用 cudaMallocManaged 分配的内存默认附加到所有流（“全局”），我们必须修改它以便有效地使用流，例如用于计算/复制重叠。我们可以使用cudaStreamAttachMemAsync 函数来做到这一点，如 J.2.2.3 节中所述。通过以这种方式将每个内存“块”与流相关联，UM 子系统可以就何时传输每个数据项做出明智的决策。

以下示例演示了这一点：

#include <stdio.h>
#include <time.h>
#define DSIZE 1048576
#define DWAIT 100000ULL
#define nTPB 256

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

typedef int mytype;

__global__ void mykernel(mytype *data){

  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  if (idx < DSIZE) data[idx] = 1;
  unsigned long long int tstart = clock64();
  while (clock64() < tstart + DWAIT);
}

int main(){

  mytype *data1, *data2, *data3;
  cudaStream_t stream1, stream2, stream3;
  cudaMallocManaged(&data1, DSIZE*sizeof(mytype));
  cudaMallocManaged(&data2, DSIZE*sizeof(mytype));
  cudaMallocManaged(&data3, DSIZE*sizeof(mytype));
  cudaCheckErrors("cudaMallocManaged fail");
  cudaStreamCreate(&stream1);
  cudaStreamCreate(&stream2);
  cudaStreamCreate(&stream3);
  cudaCheckErrors("cudaStreamCreate fail");
  cudaStreamAttachMemAsync(stream1, data1);
  cudaStreamAttachMemAsync(stream2, data2);
  cudaStreamAttachMemAsync(stream3, data3);
  cudaDeviceSynchronize();
  cudaCheckErrors("cudaStreamAttach fail");
  memset(data1, 0, DSIZE*sizeof(mytype));
  memset(data2, 0, DSIZE*sizeof(mytype));
  memset(data3, 0, DSIZE*sizeof(mytype));
  mykernel<<<(DSIZE+nTPB-1)/nTPB, nTPB, 0, stream1>>>(data1);
  mykernel<<<(DSIZE+nTPB-1)/nTPB, nTPB, 0, stream2>>>(data2);
  mykernel<<<(DSIZE+nTPB-1)/nTPB, nTPB, 0, stream3>>>(data3);
  cudaDeviceSynchronize();
  cudaCheckErrors("kernel fail");
  for (int i = 0; i < DSIZE; i++){
    if (data1[i] != 1) {printf("data1 mismatch at %d, should be: %d, was: %d\n", i, 1, data1[i]); return 1;}
    if (data2[i] != 1) {printf("data2 mismatch at %d, should be: %d, was: %d\n", i, 1, data2[i]); return 1;}
    if (data3[i] != 1) {printf("data3 mismatch at %d, should be: %d, was: %d\n", i, 1, data3[i]); return 1;}
    }
  printf("Success!\n");
  return 0;
}

上面的程序使用clock64()创建了一个人工长时间运行的内核，以便为我们提供计算/复制重叠的模拟机会（模拟计算密集型内核）。我们正在启动这个内核的 3 个实例，每个实例都在一个单独的“块”数据上运行。

当我们对上述程序进行剖析时，会看到以下内容：

首先，请注意第三次内核启动以黄色突出显示，它在第二次内核启动以紫色突出显示之后立即开始。启动此第三个内核的实际cudaLaunch 运行时 API 事件在运行时 API 行中由鼠标指针指示，也以黄色突出显示（前面是前 2 个内核的 cudaLaunch 事件）。由于此启动发生在第一个内核的执行期间，并且从该点到第三个内核启动之前没有中间的“空白空间”，我们可以观察到第三个内核启动的数据传输（即data3 ) 发生在内核 1 和 2 正在执行时。因此我们有复制和计算的有效重叠。 （我们可以对内核 2 进行类似的观察）。

虽然我没有在这里展示它，但如果我们省略 cudaStreamAttachMemAsync 行，程序仍然可以正确编译和运行，但是如果我们对其进行分析，我们会观察到 cudaLaunch 事件和内核之间的不同关系。整体配置文件看起来相似，并且内核正在背靠背执行，但是整个 cudaLaunch 进程现在开始并在第一个内核开始执行之前结束，并且在内核执行期间没有 cudaLaunch 事件。这表明（因为所有 cudaMallocManaged 内存都是全局的）所有数据传输都发生在第一次内核启动之前。该程序无法将“全局”分配与任何特定内核相关联，因此必须在第一次内核启动之前传输所有此类分配的内存（即使该内核仅使用data1）。