我想知道是否有官方消息来源，为什么以下方法起作用：
#include <iostream>
struct Array{
int el[10000];
};
__device__ Array devAr;
void test(Array& ar = devAr){
for(int i=0; i<10000; i++)
ar.el[i] = i;
std::cout << ar.el[0] + ar.el[9999] << std::endl;
}
int main(){
test();
}
如果尝试直接访问devAr，但收到警告“无法在主机函数中直接读取__device__变量“ devAr””，但通过引用没有此类警告（有充分的理由）。但是在两种情况下，都可以从主机访问变量。如此看来，该变量存在一个主机实例。
我需要知道的是：我可以认为这是理所当然的吗？
其他显示指针值的测试用例：
#include <iostream>
#include <cstdio>
__device__ int devAr[2];
__global__ void foo(){
printf("Device: %p\n", &devAr);
devAr[0] = 1337;
}
int main()
{
devAr[0] = 4;
std::cout << devAr[0] << std::endl;
void* ad;
cudaGetSymbolAddress(&ad, devAr);
std::cout << ad << " " << &devAr << std::endl;
foo<<<1,1>>>();
cudaDeviceSynchronize();
int arHost[2];
cudaMemcpyFromSymbol(arHost, devAr, sizeof(arHost), 0);
std::cout << "values: " << arHost[0] << std::endl;
}
输出：
4
0x500bc0000 0x66153c
设备：0x500bc0000
值：1337
您所做的是无效的，您应该听警告：
一个
__device__变量devAr不能在主机功能被直接读
首先，让我将您的代码简化为仅显示问题所需的大小：
#include <iostream>
__device__ int devAr[1];
int main()
{
devAr[0] = 4;
std::cout << devAr[0] << std::endl;
}
现在发生了什么：
__device__ int devAr[1];在设备内存中分配固定大小的数组，并将指向该设备内存的指针存储在devAr变量内（因此出现警告）。devAr地址指向设备存储器的有效件，然而，这种地址可以即使在主机代码使用，因为主机和设备存储器使用的地址以相同的格式。但是，主机代码中devAr指向一些随机的未初始化的主机存储器。devAr[0] = 4;只是写入4主机内存中的某个随机未初始化位置。尝试运行以下代码，也许它将帮助您了解幕后发生的事情：
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
using namespace std;
__device__ int devAr[1];
__global__ void foo()
{
printf("dev: %d \n", devAr[0]);
devAr[0] = 5;
printf("dev: %d \n", devAr[0]);
}
int main()
{
cout << "host: " << devAr[0] << endl;
devAr[0] = 4;
cout << "host: " << devAr[0] << endl;
foo << <1, 1 >> >();
cudaDeviceSynchronize();
cout << "host: " << devAr[0] << endl;
}
输出将是：
host: 0
host: 4
dev: 0
dev: 5
host: 4
更新：
在阐明了您在以下注释中的要求之后，我开始研究该问题并发现了几个相关的SO线程，其中大部分引用来自答案下方的注释，它们是：
cudaMemcpy（）和cudaMemcpyFromSymbol（）：
任何静态定义的设备符号（
__device__，__constant__甚至纹理）都会导致工具链发出两个符号，一个在设备模块中，另一个在宿主对象中。CUDA运行时在这两个符号之间建立并维护动态映射。符号API调用是检索__constant__和__device__符号的映射的方式。纹理API检索纹理符号等的映射。
*PNT是__device__变量，而不是包含设备变量地址的主机变量。（令人困惑，我知道。）因此，如果您尝试像在主机上一样访问主机上的(void**)&PNT设备变量，则尝试在主机上访问它，这是不允许的。从主机代码的角度来看，它只是一个符号，因此您需要使用cudaGetSympolAddress()将设备地址存储在主机变量中，然后可以将其传递给cudaMemcpyToSymbol()@talonmies所示。
令人困惑的是，主机代码中的A和B不是有效的设备内存地址。它们是主机符号，它们提供了到运行时设备符号查找的挂钩。将它们传递给内核是非法的-如果您想要它们的设备内存地址，则必须
cudaGetSymbolAddress在运行时用于检索它。
cudaMemcpyToSymbol vs. cudaMemcpy为什么仍然存在（cudaMemcpyToSymbol）：
通过CUDA API复制到该地址的操作将失败，并带有无效的参数错误，因为它不是该API先前分配的GPU内存空间中的地址。是的，这也适用于通用
__device__指针和静态声明的设备符号。
在__device__变量上的cudaMemcpyFromSymbol：
问题的根源是不允许您使用普通主机代码中的设备变量的地址：...尽管这似乎可以正确编译，但是传递的实际地址是垃圾。要获取主机代码中设备变量的地址，我们可以使用
cudaGetSymbolAddress
基于这些证据，让我尝试从上面更新我的原始3step说明：
__device__ int devAr[1];在设备内存中分配固定大小的数组，并将“钩到运行时设备符号查找中”存储到devAr变量的主机版本中（请参阅链接的资源1和3）。devAr从主机的角度来看，该地址只是一个垃圾，仅应与符号API调用一起使用，例如cudaGetSymbolAddress（所有链接的资源似乎都支持该理论），因为该地址映射到devAr变量的设备版本。我无法提出任何“更具体的内容”，例如CUDA文档的链接，但我希望现在已经足够清楚了。总而言之，您现在似乎已经可以保证上述行为（例如，存在主机和设备版本的devAr变量），但对我而言，它似乎是您不应依赖且不应使用主机版本的实现细节。devAr用于符号API调用以外的目的的变量。
本文收集自互联网，转载请注明来源。
如有侵权，请联系 [email protected] 删除。
我来说两句