CUDA多gpu:运行相同内核(双芯片设备)

Question

我正在编写必须在GTX690（双芯片）板上运行相同内核的代码。由于计算是非常可分离的，我不需要在设备之间进行数据交换，我只需要在CPU上合并结果。我了解如何在每个设备上运行代码，也了解如何在每台设备的内存空间中提供I/o数据空间。

我在尝试设置常量时遇到了问题，在两台设备上运行的内核都使用常量。我需要为每个设备制作一个阵列吗？如果是这样，我如何在内核中检查哪个设备正在运行，以便从数组中访问正确的常量？

这是我正在使用的一些代码。当我只使用1个芯片（通过设置numDev = 1）时，它可以正常工作，但不能同时使用这两个芯片。

__constant__ float d_cellSizeZ;  
std::vector<int*> d_cell;  
.................  
bool Init(int cellsN_, float size_){  
  bool res = true;
  if(cudaSuccess != cudaGetDeviceCount(&numDev))  
    return false;  
  //numDev = 1;  
  d_cl.resize(numDev);  
  for(int i = 0; i < numDev; ++i){  
    res &= (cudaSuccess == cudaSetDevice(i));  
    if(!res)  
      break;  
    res &= (cudaSuccess == cudaMalloc((void**)&d_cell[i], cellsN_*sizef(int)));  
  };  
  res &= (cudaSuccess == cudaMemcpyToSymbol(d_cellSizeZ, &size_, sizeof(float)));  
  if(!res)  
    Cleanup();  
  return res;  
}

在内核中，我只使用d_cellSizeZ常量。那么，我应该如何为每个设备设置一个常数，以便从同一内核中很好地使用它呢？

另一个问题是：如果我尝试在设备之间交换数据，它会通过PCI总线还是在双芯片板上存在一些内部路径？

Answer 1

您只在最后一个设备上初始化d_cellSizeZ，所以在其他设备上它将是未定义的。您需要在每个设备上初始化d_cellSizeZ，最简单的方法是在循环中进行初始化，正如Greg在评论中建议的那样：

for(int i = 0; i < numDev; ++i)
{  
    checkCudaErrors(cudaSetDevice(i));
    checkCudaErrors(cudaMalloc((void**)&d_cell[i], cellsN_*sizef(int)));
    checkCudaErrors(cudaMemcpyToSymbol(d_cellSizeZ, &size_, sizeof(float)));
};

重复使用d_cellSizeZ符号确实有点奇怪。幕后有一点小聪明，但本质上cudaMemcpyToSymbol()函数会查找当前活动设备上的符号，因此每次都会复制到正确的设备。