多GPU多个流上实现复制与计算的重叠

缺乏、安全感 2022-11-17 14:41 199阅读 0赞
  • 获取应用程序可使用的 GPU 的数量
  • 激活任意可用的 GPU
  • 在多个 GPU 上分配显存
  • 在多个 GPU 上传入和转出显存数据
  • 在多个 GPU 上启动核函数

获取多个 GPU 的相关信息

如要以运行程序的方式得出可用 GPU 的数量，请使用 cudaGetDeviceCount

  1. uint64_t num_gpus;
  2. cudaGetDeviceCount(&num_gpus);

如要以运行程序的方式得到当前处于活动状态的 GPU，请使用 cudaGetDevice：

  1. uint64_t device;
  2. cudaGetDevice(&device); // `device` is now a 0-based index of the current GPU.

设置当前的 GPU

对于每个主机线程，每次只有一个 GPU 设备处于活动状态。如要将特定的 GPU 设置为活动状态，请使用 cudaSetDevice 以及所需 GPU 的索引（从 0 开始）：

  1. cudaSetDevice(0);

循环使用可用的 GPU

一种常见的模式为，遍历可用的 GPU，并为每个 GPU 执行相应操作：

  1. uint64_t num_gpus;
  2. cudaGetDeviceCount(&num_gpus);
  4. for (uint64_t gpu = 0; gpu < num_gpus; gpu++) { 
  6.     cudaSetDevice(gpu);
  7.     // Perform operations for this GPU.
  8. }

为多个 GPU 执行数据分块

与多个非默认流相同，多个 GPU 中的每个 GPU 都可处理一个数据块。我们将创建和利用数据指针数组，为每个可用的 GPU 分配显存：

  1. uint64_t num_gpus;
  2. cudaGetDeviceCount(&num_gpus);
  4. const uint64_t num_entries = 1UL << 26;
  5. const uint64_t chunk_size = sdiv(num_entries, num_gpus);
  7. uint64_t *data_gpu[num_gpus]; // One pointer for each GPU.
  9. for (uint64_t gpu = 0; gpu < num_gpus; gpu++) { 
  11.     cudaSetDevice(gpu);
  13.     const uint64_t lower = chunk_size*gpu;
  14.     const uint64_t upper = min(lower+chunk_size, num_entries);
  15.     const uint64_t width = upper-lower;
  17.     cudaMalloc(&data_gpu[gpu], sizeof(uint64_t)*width); // Allocate chunk of data for current GPU.
  18. }

为多个 GPU 复制数据
通过使用相同的循环遍历和分块技术，我们可在多个 GPU 上传入和传出数据：

  1. // ...Assume data has been allocated on host and for each GPU
  3. for (uint64_t gpu = 0; gpu < num_gpus; gpu++) { 
  5.     cudaSetDevice(gpu);
  7.     const uint64_t lower = chunk_size*gpu;
  8.     const uint64_t upper = min(lower+chunk_size, num_entries);
  9.     const uint64_t width = upper-lower;
  11.     // Note use of `cudaMemcpy` and not `cudaMemcpyAsync` since we are not
  12.     // presently using non-default streams.
  13.     cudaMemcpy(data_gpu[gpu], data_cpu+lower, 
  14.            sizeof(uint64_t)*width, cudaMemcpyHostToDevice); // ...or cudaMemcpyDeviceToHost
  15. }

为多个 GPU 启动核函数

通过使用相同的循环遍历和分块技术，我们可在多个 GPU 上启动核函数并处理数据块：

  1. // ...Assume data has been allocated on host and for each GPU
  3. for (uint64_t gpu = 0; gpu < num_gpus; gpu++) { 
  5.     cudaSetDevice(gpu);
  7.     const uint64_t lower = chunk_size*gpu;
  8.     const uint64_t upper = min(lower+chunk_size, num_entries);
  9.     const uint64_t width = upper-lower;
  11.     kernel<<<grid, block>>>(data_gpu[gpu], width); // Pass chunk of data for current GPU to work on.
  12. }

在上面，使用深度优先的方法将一部分工作传递给每个GPU。在某些情况下，尤其是在数据量极高的情况下，使用宽度优先的方法可能更有意义。这种方法上的改变并不是需要额外的CUDA知识。不过，此stack overflow的回答提供了一些使用深度优先和宽度优先方法的CUDA代码示例。

多个GPU之间进行对等内存传输，以及在多个节点上使用多个GPU， 此超级计算会议演示文稿。

例子：
多个GPU使用默认流

  1. #include <cstdint>
  2. #include <iostream>
  3. #include "helpers.cuh"
  4. #include "encryption.cuh"
  6. void encrypt_cpu(uint64_t * data, uint64_t num_entries, 
  7.                  uint64_t num_iters, bool parallel=true) {
  9.     #pragma omp parallel for if (parallel)
  10.     for (uint64_t entry = 0; entry < num_entries; entry++)
  11.         data[entry] = permute64(entry, num_iters);
  12. }
  14. __global__ 
  15. void decrypt_gpu(uint64_t * data, uint64_t num_entries, 
  16.                  uint64_t num_iters) {
  18.     const uint64_t thrdID = blockIdx.x*blockDim.x+threadIdx.x;
  19.     const uint64_t stride = blockDim.x*gridDim.x;
  21.     for (uint64_t entry = thrdID; entry < num_entries; entry += stride)
  22.         data[entry] = unpermute64(data[entry], num_iters);
  23. }
  25. bool check_result_cpu(uint64_t * data, uint64_t num_entries,
  26.                       bool parallel=true) {
  28.     uint64_t counter = 0;
  30.     #pragma omp parallel for reduction(+: counter) if (parallel)
  31.     for (uint64_t entry = 0; entry < num_entries; entry++)
  32.         counter += data[entry] == entry;
  34.     return counter == num_entries;
  35. }
  37. int main (int argc, char * argv[]) {
  39.     Timer timer;
  40.     Timer overall;
  42.     const uint64_t num_entries = 1UL << 26;
  43.     const uint64_t num_iters = 1UL << 10;
  44.     const bool openmp = true;
  46.     timer.start();
  47.     uint64_t * data_cpu;
  48.     cudaMallocHost(&data_cpu, sizeof(uint64_t)*num_entries);
  49.     // cudaMalloc    (&data_gpu, sizeof(uint64_t)*num_entries);
  51.     timer.stop("allocate memory");
  52.     check_last_error();
  54.     timer.start();
  55.     encrypt_cpu(data_cpu, num_entries, num_iters, openmp);
  56.     timer.stop("encrypt data on CPU");
  58.     overall.start();
  59.     timer.start();
  61.     int num_gpus;
  62.     cudaGetDeviceCount(&num_gpus);
  63.     const uint64_t chunk_size = sdiv(num_entries, num_gpus);
  64.     uint64_t *data_gpu[num_gpus]; // One pointer for each GPU.
  65.     for (uint64_t gpu = 0; gpu < num_gpus; gpu++) {
  67.         cudaSetDevice(gpu);
  69.         const uint64_t lower = chunk_size*gpu;
  70.         const uint64_t upper = min(lower+chunk_size, num_entries);
  71.         const uint64_t width = upper-lower;
  73.         cudaMalloc(&data_gpu[gpu], sizeof(uint64_t)*width); // Allocate chunk of data for current GPU.
  75.         cudaMemcpy(data_gpu[gpu], data_cpu+lower, 
  76.            sizeof(uint64_t)*width, cudaMemcpyHostToDevice); // ...or cudaMemcpyDeviceToHost
  78.     }
  80.     for (uint64_t gpu = 0; gpu < num_gpus; gpu++) {
  81.         cudaSetDevice(gpu);
  82.         const uint64_t lower = chunk_size*gpu;
  83.         const uint64_t upper = min(lower+chunk_size, num_entries);
  84.         const uint64_t width = upper-lower;
  85.         decrypt_gpu<<<32*80, 64>>>(data_gpu[gpu], width, num_iters); // Pass chunk of data for current GPU to work on.
  86.     }
  88.     for (uint64_t gpu = 0; gpu < num_gpus; gpu++) {
  89.         cudaSetDevice(gpu);
  90.         const uint64_t lower = chunk_size*gpu;
  91.         const uint64_t upper = min(lower+chunk_size, num_entries);
  92.         const uint64_t width = upper-lower;
  93.         cudaMemcpy(data_cpu+lower, data_gpu[gpu], 
  94.                sizeof(uint64_t)*width, cudaMemcpyDeviceToHost);
  95.     }
  98.     timer.stop("multi GPU times");
  99.     check_last_error();
  101.     overall.stop("total time on GPU");
  102.     check_last_error();
  104.     timer.start();
  105.     const bool success = check_result_cpu(data_cpu, num_entries, openmp);
  106.     std::cout << "STATUS: test " 
  107.               << ( success ? "passed" : "failed")
  108.               << std::endl;
  109.     timer.stop("checking result on CPU");
  111.     timer.start();
  112.     cudaFreeHost(data_cpu);
  113.     // cudaFree    (data_gpu);
  114.     for (uint64_t gpu = 0; gpu < num_gpus; gpu++) {
  115.       cudaFree(data_gpu[gpu]);
  116.     }
  117.     timer.stop("free memory");
  118.     check_last_error();
  119. }

在上面代码，可以看到内存传输没有重叠。 为什么会这样？
代码既不使用非默认流，也不使用“cudaMemcpyAsync”存储复制。 因此，它们阻止了操作。

在多个 GPU上实现数据复制与计算的重叠

  • 流与每个 GPU 设备是如何关联的
  • 如何为多个 GPU 创建非默认流
  • 如何在多个 GPU 上实现复制与计算的重叠

每个 GPU 都有各自的默认流。我们可以为当前处于活动状态的 GPU 设备创建、使用和销毁非默认流。切记不要在未与当前处于活动状态的 GPU 建立关联的流中启动核函数

为多个 GPU 创建多个非默认流
在多个 GPU 上使用多个非默认流时，与之前不同的是，我们不是简单地将流存储在数组中，而是将其存储于二维数组中，且数组中的每一行皆包含单个 GPU 的流：

  1. cudaStream_t streams[num_gpus][num_streams]; // 2D array containing number of streams for each GPU.
  3. // For each available GPU...
  4. for (uint64_t gpu = 0; gpu < num_gpus; gpu++) { 
  5.     // ...set as active device...
  6.     cudaSetDevice(gpu);
  7.     for (uint64_t stream = 0; stream < num_streams; stream++)
  8.         // ...create and store its number of streams.
  9.         cudaStreamCreate(&streams[gpu][stream]);
  10. }

多个 GPU 上多流的数据块大小
当在多个 GPU 上使用多个非默认流时，全局数据索引尤为棘手。为帮助实现索引，我们可以为单个流和整个 GPU 分别定义数据块大小。

  1. // Each stream needs num_entries/num_gpus/num_streams data. We use round up division for
  2. // reasons previously discussed.
  3. const uint64_t stream_chunk_size = sdiv(sdiv(num_entries, num_gpus), num_streams);
  5. // It will be helpful to also to have handy the chunk size for an entire GPU.
  6. const uint64_t gpu_chunk_size = stream_chunk_size*num_streams;

为多个 GPU 的多个流分配显存
GPU 的显存并未分配给各个流，所以此处的分配操作看起来与之前的多 GPU 任务相似，我们只需注意数据块的大小是分配给整个 GPU 的而非其中一个流的即可：

  1. ```c
  2. // For each GPU...
  3. for (uint64_t gpu = 0; gpu < num_gpus; gpu++) {
  5.     // ...set device as active...
  6.     cudaSetDevice(gpu);
  8.     // ...use a GPU chunk's worth of data to calculate indices and width...
  9.     const uint64_t lower = gpu_chunk_size*gpu;
  10.     const uint64_t upper = min(lower+gpu_chunk_size, num_entries);
  11.     const uint64_t width = upper-lower;
  13.     // ...allocate data.
  14.     cudaMalloc(&data_gpu[gpu], sizeof(uint64_t)*width);
  15. }

在多个 GPU 的多个流上实现复制与计算的重叠

  1. #include <cstdint>
  2. #include <iostream>
  3. #include "helpers.cuh"
  4. #include "encryption.cuh"
  6. void encrypt_cpu(uint64_t * data, uint64_t num_entries, 
  7.                  uint64_t num_iters, bool parallel=true) { 
  9.     #pragma omp parallel for if (parallel)
  10.     for (uint64_t entry = 0; entry < num_entries; entry++)
  11.         data[entry] = permute64(entry, num_iters);
  12. }
  14. __global__ 
  15. void decrypt_gpu(uint64_t * data, uint64_t num_entries, 
  16.                  uint64_t num_iters) { 
  18.     const uint64_t thrdID = blockIdx.x*blockDim.x+threadIdx.x;
  19.     const uint64_t stride = blockDim.x*gridDim.x;
  21.     for (uint64_t entry = thrdID; entry < num_entries; entry += stride)
  22.         data[entry] = unpermute64(data[entry], num_iters);
  23. }
  25. bool check_result_cpu(uint64_t * data, uint64_t num_entries,
  26.                       bool parallel=true) { 
  28.     uint64_t counter = 0;
  30.     #pragma omp parallel for reduction(+: counter) if (parallel)
  31.     for (uint64_t entry = 0; entry < num_entries; entry++)
  32.         counter += data[entry] == entry;
  34.     return counter == num_entries;
  35. }
  37. int main (int argc, char * argv[]) { 
  39.     Timer timer;
  40.     Timer overall;
  42.     const uint64_t num_entries = 1UL << 26;
  43.     const uint64_t num_iters = 1UL << 10;
  44.     const bool openmp = true;
  46.     timer.start();
  47.     uint64_t * data_cpu;
  48.     cudaMallocHost(&data_cpu, sizeof(uint64_t)*num_entries);
  49.     timer.stop("allocate memory");
  50.     check_last_error();
  52.     timer.start();
  53.     encrypt_cpu(data_cpu, num_entries, num_iters, openmp);
  54.     timer.stop("encrypt data on CPU");
  57.     int num_gpus;
  58.     cudaGetDeviceCount(&num_gpus);
  59.     uint64_t num_streams = 32;
  60.     uint64_t *data_gpu[num_gpus];
  61.     cudaStream_t streams[num_gpus][num_streams];
  63.     uint64_t gpu_chunk_size = sdiv(num_entries, num_gpus);
  65.     overall.start();
  67.     for (uint64_t gpu = 0; gpu < num_gpus; gpu++) { 
  69.         cudaSetDevice(gpu);
  71.         for (uint64_t s = 0; s < num_streams; s++) { 
  72.             cudaStreamCreate(&streams[gpu][s]);
  73.         }
  75.         uint64_t gpu_lower = gpu_chunk_size* gpu;
  77.         uint64_t gpu_upper = min(gpu_lower+gpu_chunk_size, num_entries);
  79.         uint64_t gpu_width = gpu_upper-gpu_lower;
  81.         cudaMalloc(&data_gpu[gpu], sizeof(uint64_t)*gpu_width);      
  82.     }
  85.     for (uint64_t gpu = 0; gpu < num_gpus; gpu++) { 
  87.         cudaSetDevice(gpu);
  89.         uint64_t gpu_lower = gpu_chunk_size * gpu;
  91.         uint64_t gpu_upper = min(gpu_lower+gpu_chunk_size, num_entries);
  93.         uint64_t gpu_width = gpu_upper - gpu_lower;
  96.         uint64_t s_chunk_size = sdiv(gpu_width, num_streams);
  98.         for (uint64_t s = 0; s < num_streams; s++) { 
  99.             uint64_t s_offset = s * s_chunk_size;
  100.             uint64_t s_lower = gpu_lower + s_offset;
  101.             uint64_t s_upper = min(s_lower + s_chunk_size, gpu_upper);
  102.             uint64_t s_width = s_upper - s_lower;
  104.             cudaMemcpyAsync(data_gpu[gpu] + s_offset, data_cpu+s_lower, 
  105.                    sizeof(uint64_t)*s_width, cudaMemcpyHostToDevice, streams[gpu][s]);
  108.             decrypt_gpu<<<80*32, 64, 0, streams[gpu][s]>>>(data_gpu[gpu]+s_offset, s_width, num_iters);
  111.             cudaMemcpyAsync(data_cpu+s_lower, data_gpu[gpu] + s_offset, 
  112.                sizeof(uint64_t)*s_width, cudaMemcpyDeviceToHost, streams[gpu][s]);
  113.         }
  114.     }
  116.     check_last_error();
  118.     for (uint64_t gpu = 0; gpu < num_gpus; gpu++) { 
  119.         cudaSetDevice(gpu);
  120.         for (uint64_t s = 0; s < num_streams; s++) { 
  121.             cudaStreamSynchronize(streams[gpu][s]);
  122.         }
  123.     }
  125.     for (uint64_t gpu = 0; gpu < num_gpus; gpu++) { 
  126.         cudaSetDevice(gpu);
  127.         for (uint64_t s = 0; s < num_streams; s++) { 
  128.            cudaStreamDestroy(streams[gpu][s]);
  129.         }
  130.     }
  132.     overall.stop("total time on GPU");
  133.     check_last_error();
  135.     timer.start();
  136.     bool success = true;
  137.     success = check_result_cpu(data_cpu, num_entries, openmp);
  138.     std::cout << "STATUS: test " 
  139.               << ( success ? "passed" : "failed")
  140.               << std::endl;
  141.     timer.stop("checking result on CPU");
  143.     timer.start();
  144.     cudaFreeHost(data_cpu);
  145.     for (uint64_t gpu = 0; gpu < num_gpus; gpu++) { 
  146.         cudaSetDevice(gpu);
  147.         cudaFree(data_gpu[gpu]);
  148.     }
  149.     timer.stop("free memory");
  150.     check_last_error();
  151. }

发表评论

表情：
评论列表 （有 0 条评论，199人围观）

还没有评论，来说两句吧...

相关阅读