0

我无法在任何地方找到答案，我可能忽略了它，但似乎无法使用__constant__内存（连同cudaMemcpyToSymbol）和 UVA 的点对点访问。

我已经尝试了 simpleP2P nvidia 示例代码，它在我拥有的带有 nvlink 的 4 NV100 上运行良好，但只要我在内核中将因子 2 声明为：

__constant__ float M_; // in global space

float M = 2.0;
cudaMemcpyToSymbol(M_, &M, sizeof(float), 0, cudaMemcpyDefault);

结果基本为零。如果我使用 C 预处理器（例如 #define M_ 2.0）定义它，它工作正常。

所以我想知道，这是真的还是我做错了什么？是否还有其他类型的内存也不能以这种方式访问​​（例如纹理内存）？

4

1 回答 1

1

我不清楚为什么“结果基本上为零”的问题与使用 UVA 的 P2P 访问之间的关系。

这是真的还是我做错了什么？

很难说，因为您的问题有点含糊，并且没有显示完整的示例。

__constant__ float M_在所有CUDA 可见设备M_的常量内存上 分配一个变量。为了在多个设备上设置值，您应该执行以下操作：

__constant__ float M_; // <= This declares M_ on the constant memory of all CUDA visible devices

__global__ void showMKernel() {
    printf("****** M_ = %f\n", M_);
}

int main()
{

float M = 2.0;

 // Make sure that the return values are properly checked for cudaSuccess ...

int deviceCount = -1;
cudaGetDeviceCount(&deviceCount);

// Set M_ on the constant memory of each device:
for (int i = 0; i < deviceCount; i++) {
  cudaSetDevice(i);
  cudaMemcpyToSymbol(M_, &M, sizeof(float), 0, cudaMemcpyDefault);
}

// Now, run a kernel to show M_:
for (int i = 0; i < deviceCount; i++) 
{
  cudaSetDevice(i);
  printf("Device %g :\n", i);
  showMKernel<<<1,1>>>();
  cudaDeviceSynchronize();
}

}

返回：

Device 0 :
****** M = 2.000000
Device 1 :
****** M = 2.000000
// so on for other devices

现在，如果我更换

// Set M_ on the constant memory of each device:
for (int i = 0; i < deviceCount; i++) {
  cudaSetDevice(i);
  cudaMemcpyToSymbol(M_, &M, sizeof(float), 0, cudaMemcpyDefault);
}

和

cudaMemcpyToSymbol(M_, &M, sizeof(float), 0, cudaMemcpyDefault);

这只会M_在活动设备上设置 的值，因此返回

Device 0 :
****** M = 2.000000
Device 1 :
****** M = 0.000000 // <= I assume this is what you meant by 'the results are basically zero'
// M = 0 for other devices too

是否还有其他类型的内存也无法以这种方式访问​​（例如纹理内存）？

同样，我不完全确定这种方式是什么。我认为通常您无法从任何其他设备访问一个设备的常量内存或纹理内存，尽管我不是 100% 确定。

UVA 为 CPU 和 GPU 内存分配一个地址空间，以便通过使用with kind可以轻松访问主机和多个设备的全局内存之间的内存复制。cudaMemcpycudaMemcpyDefault

此外，设备之间的 P2P 通信允许在多个设备的全局内存之间直接访问和传输数据。

与上面的示例类似__constant__，当您声明类似 的纹理时texture <float> some_texture，将为每个可见设备定义，但是在使用多个设备时，some_texture您需要显式绑定到每个设备上的纹理引用。some_texture

于 2020-03-10T15:33:10.193 回答