CUDA P2P存储器访问和__constant__存储器
我在任何地方都找不到答案，我可能忽略了这一点，但似乎不能使用__constant__内存(以及cudaMemcpyToSymbol)和UVA的点对点访问。
我已经尝试过simpleP2P nvidia示例代码，它在我的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)来定义它，它会工作得很好。
所以我想知道，这是真的还是我做错了什么？还有其他类型的内存也不能以这种方式访问吗(例如纹理内存)？
Stack Overflow用户
发布于 2020-03-10 23:33:10
你的问题为什么“结果基本上是零”和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为中央处理器和图形处理器存储器分配一个地址空间，使得通过使用具有kind cudaMemcpyDefault的cudaMemcpy，可以容易地访问主机和多个设备的全局存储器之间的存储器复制。
此外，设备之间的P2P通信允许在多个设备的全局存储器之间直接访问和传输数据。
与上面的__constant__示例类似，当声明像texture <float> some_texture这样的纹理时，将为每个可见设备定义some_texture，但是在使用多个设备时，需要将some_texture显式绑定到每个设备上的纹理引用。
https://stackoverflow.com/questions/60606113
复制相似问题
腾讯云开发者
