将指向多个设备的指针数组传递给 Cuda C 内核答案

【问题标题】：Pass array of pointers to multiple devices to Cuda C Kernel将指向多个设备的指针数组传递给 Cuda C 内核
【发布时间】：2020-01-01 01:05:46
【问题描述】：

我有一个需要处理的一维数组，但它对于单个 GPU 来说太大了。因此，我将数组传递给多个 GPU 以存储在内存中，其数量将根据问题的大小而变化。如果我将一组指针传递给不同 GPU 中的数组，我将无法从我的 Cuda C 内核访问其他数组。

我尝试通过内核调用将一个简单的设备指针数组传递给每个设备，但是当我尝试访问这些数组时，代码似乎中断了。即使是运行内核的设备也无法访问自己内存中的数组。

数据结构：

typedef struct ComplexArray
{
   double *real;
} ComplexArray;

typedef struct ComplexArrayArray
{
   ComplexArray* Arr;
} ComplexArrayArray;

马洛克：

ComplexArrayArray stateVector;
stateVector.Arr = (ComplexArray*)malloc(sizeof(ComplexArray*) * numberOfGPU));

for (int dev = 0; dev < numberOfGPI; dev++)
{
    ...
    cudaMalloc(&(stateVector.Arr[dev].real), numberOfElements * sizeof(*(stateVector.Arr[dev].real)) / numberOfGPU);
    ...
}

内核：

__global__ void kernel(..., ComplexArrayArray stateVector, ...)
{
   // Calculate necessary device
   int device_number = ...;
   int index = ...;

   double val = stateVector.Arr[device_number].real[index];
   ...
}

当我尝试以这种方式访问数组时，内核似乎“中断”了。没有错误信息，但很明显数据没有被读取。此外，在数据访问之后，我没有达到任何 printf 语句。

关于将指向设备内存的指针数组传递给 Cuda C 内核的最佳方法有什么想法吗？

【问题讨论】：

您的stateVector 定义ComplexArrayArray stateVector ... 不是指针。给它分配一个指针值怎么可能有意义？这甚至可以编译吗？事实上，这一行代码可能存在 3 个不同的问题。
你是对的，错字。我尝试简化它来自的代码并且写错了。我正在使用 malloc 为结构中的指针数组分配。

标签： cuda gpgpu multi-gpu

【解决方案1】：

您尝试使用带有指向结构数组的指针的结构，每个结构数组都有一个嵌入的指针，这将使cudaMalloc 的实现变得非常复杂。如果您使用cudaMallocManaged，它可能会更简单一些，但仍然不必要地复杂。复杂性的出现是因为cudaMalloc 在特定设备上分配空间，并且（默认情况下）任何其他设备都无法访问该数据，并且还因为您的嵌入式指针创建了各种“深拷贝”的必要性。这是一个有效的例子：

$ cat t1492.cu
#include <iostream>
#include <stdio.h>
typedef struct ComplexArray
{
   double *real;
} ComplexArray;

typedef struct ComplexArrayArray
{
   ComplexArray* Arr;
} ComplexArrayArray;

__global__ void kernel(ComplexArrayArray stateVector, int dev, int ds)
{
   // Calculate necessary device
   int device_number = dev;
   int index = blockIdx.x*blockDim.x+threadIdx.x;
   if (index < ds){
     double val = stateVector.Arr[device_number].real[index] + dev;
     stateVector.Arr[device_number].real[index] = val;
   }
}
const int nTPB = 256;
int main(){
  int numberOfGPU;
  cudaGetDeviceCount(&numberOfGPU);
  std::cout << "GPU count: " << numberOfGPU << std::endl;
  ComplexArrayArray *stateVector = new ComplexArrayArray[numberOfGPU];
  const int ds = 32;
  double *hdata = new double[ds]();
  ComplexArray *ddata = new ComplexArray[numberOfGPU];
  for (int i = 0; i < numberOfGPU; i++){
    cudaSetDevice(i);
    cudaMalloc(&(stateVector[i].Arr), sizeof(ComplexArray) * numberOfGPU);
    cudaMalloc(&(ddata[i].real), (ds/numberOfGPU)*sizeof(double));
    cudaMemcpy(ddata[i].real, hdata + i*(ds/numberOfGPU), (ds/numberOfGPU)*sizeof(double), cudaMemcpyHostToDevice);}
  for (int i = 0; i < numberOfGPU; i++){
    cudaSetDevice(i);
    cudaMemcpy(stateVector[i].Arr, ddata, sizeof(ComplexArray)*numberOfGPU, cudaMemcpyHostToDevice);}
  for (int i = 0; i < numberOfGPU; i++){
    cudaSetDevice(i);
    kernel<<<((ds/numberOfGPU)+nTPB-1)/nTPB,nTPB>>>(stateVector[i], i, (ds/numberOfGPU));}
  for (int i = 0; i < numberOfGPU; i++){
    cudaSetDevice(i);
    cudaMemcpy(hdata + i*(ds/numberOfGPU), ddata[i].real,  (ds/numberOfGPU)*sizeof(double), cudaMemcpyDeviceToHost);}
  for (int i = 0; i < ds; i++)
    std::cout << hdata[i] << " ";
  std::cout << std::endl;
}
$ nvcc -o t1492 t1492.cu
$ cuda-memcheck ./t1492
========= CUDA-MEMCHECK
GPU count: 4
0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 2 2 2 2 2 2 2 2 3 3 3 3 3 3 3 3
========= ERROR SUMMARY: 0 errors
$

但是，如果您想将主机阵列划分为每个 GPU 的一个块，则不需要那种复杂程度。这是一个更简单的例子：

$ cat t1493.cu
#include <iostream>
#include <stdio.h>
typedef struct ComplexArray
{
   double *real;
} ComplexArray;

typedef struct ComplexArrayArray
{
   ComplexArray* Arr;
} ComplexArrayArray;

__global__ void kernel(ComplexArray stateVector, int dev, int ds)
{
   int index = blockIdx.x*blockDim.x+threadIdx.x;
   if (index < ds){
     double val = stateVector.real[index] + dev;
     stateVector.real[index] = val;
   }
}
const int nTPB = 256;
int main(){
  int numberOfGPU;
  cudaGetDeviceCount(&numberOfGPU);
  std::cout << "GPU count: " << numberOfGPU << std::endl;
  ComplexArray *stateVector = new ComplexArray[numberOfGPU];
  const int ds = 32;
  double *hdata = new double[ds]();
  for (int i = 0; i < numberOfGPU; i++){
    cudaSetDevice(i);
    cudaMalloc(&(stateVector[i].real), (ds/numberOfGPU)*sizeof(double));
    cudaMemcpy(stateVector[i].real, hdata + i*(ds/numberOfGPU), (ds/numberOfGPU)*sizeof(double), cudaMemcpyHostToDevice);}
  for (int i = 0; i < numberOfGPU; i++){
    cudaSetDevice(i);
    kernel<<<((ds/numberOfGPU)+nTPB-1)/nTPB,nTPB>>>(stateVector[i], i, (ds/numberOfGPU));}
  for (int i = 0; i < numberOfGPU; i++){
    cudaSetDevice(i);
    cudaMemcpy(hdata + i*(ds/numberOfGPU), stateVector[i].real,  (ds/numberOfGPU)*sizeof(double), cudaMemcpyDeviceToHost);}
  for (int i = 0; i < ds; i++)
    std::cout << hdata[i] << " ";
  std::cout << std::endl;
}
$ nvcc -o t1493 t1493.cu
$ cuda-memcheck ./t1493
========= CUDA-MEMCHECK
GPU count: 4
0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 2 2 2 2 2 2 2 2 3 3 3 3 3 3 3 3
========= ERROR SUMMARY: 0 errors
$

请注意，您的问题似乎提到了您将数据分成块的想法，并且每个内核都可能访问所有块。这将需要管理内存使用或系统可以支持 GPU 之间的 P2P 访问的知识。这增加了更多的复杂性，超出了我在这里回答的范围，重点是关于内核无法访问“它自己的”数据的问题。

由于我们应该能够设置可以参与的 GPU 数量的上限（让我们将其设置为最多 8 个），因此我们可以避免第一种方法的深度复制，同时仍然允许所有 GPU 拥有所有指针。这是一个修改过的例子：

$ cat t1495.cu
#include <iostream>
#include <stdio.h>
const int maxGPU=8;
typedef struct ComplexArray
{
   double *real[maxGPU];
} ComplexArray;

__global__ void kernel(ComplexArray stateVector, int dev, int ds)
{
   int index = blockIdx.x*blockDim.x+threadIdx.x;
   if (index < ds){
     double val = stateVector.real[dev][index] + dev;
     stateVector.real[dev][index] = val;
   }
}
const int nTPB = 256;
int main(){
  int numberOfGPU;
  cudaGetDeviceCount(&numberOfGPU);
  std::cout << "GPU count: " << numberOfGPU << std::endl;
  ComplexArray stateVector;
  const int ds = 32;
  double *hdata = new double[ds]();
  for (int i = 0; i < numberOfGPU; i++){
    cudaSetDevice(i);
    cudaMalloc(&(stateVector.real[i]), (ds/numberOfGPU)*sizeof(double));
    cudaMemcpy(stateVector.real[i], hdata + i*(ds/numberOfGPU), (ds/numberOfGPU)*sizeof(double), cudaMemcpyHostToDevice);}
  for (int i = 0; i < numberOfGPU; i++){
    cudaSetDevice(i);
    kernel<<<((ds/numberOfGPU)+nTPB-1)/nTPB,nTPB>>>(stateVector, i, (ds/numberOfGPU));}
  for (int i = 0; i < numberOfGPU; i++){
    cudaSetDevice(i);
    cudaMemcpy(hdata + i*(ds/numberOfGPU), stateVector.real[i],  (ds/numberOfGPU)*sizeof(double), cudaMemcpyDeviceToHost);}
  for (int i = 0; i < ds; i++)
    std::cout << hdata[i] << " ";
  std::cout << std::endl;
}
$ nvcc -o t1495 t1495.cu
$ cuda-memcheck ./t1495
========= CUDA-MEMCHECK
GPU count: 4
0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 2 2 2 2 2 2 2 2 3 3 3 3 3 3 3 3
========= ERROR SUMMARY: 0 errors
$

【讨论】：

感谢您抽出宝贵时间回答我的问题。我正在使用 P100 并打算使用 P2P 内存访问从不同的 GPU 读取特定元素。但是，我遇到的问题不是您概述的问题。我的问题是（以您更简单的示例）将整个“stateVector”数组传递给内核，而不仅仅是与设备关联的元素。在所有阵列分配之后，我是否需要将此阵列 cudaMemcpy 到每个设备上？我已经尝试通过在每个设备上创建一个“stateVector”副本并 Memcpy 对原始设备进行此操作，但我似乎总是遇到段错误。
是的，这实际上是我回答中的第一个示例代码所做的。如果您必须处理任意数量的 GPU，它将需要深拷贝。但它可以通过例如简化最多支持 8 个 GPU。您唯一需要添加的是在 GPU 之间启用 P2P。
我添加了另一个示例，它允许将所有数据指针传递给所有 GPU，但避免了深层复制的复杂性。
太好了，谢谢。第三个例子以最干净的方式解决了我的问题。标记为已解决。