问 CUDA P2P存储器访问和__constant__存储器

EN

你的问题为什么“结果基本上是零”和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);