从主机访问CUDA全局设备变量

Question

我想知道是否有官方来源，为什么以下有效：

#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: %pn", &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

Answer 1

您正在做的是无效的，您应该听取警告：

CCD_ 1变量CCD_

首先，让我把你的代码简化一点，只显示问题所需的大小：

#include <iostream>
__device__ int devAr[1];
int main()
{
    devAr[0] = 4;
    std::cout << devAr[0] << std::endl;
}

现在发生了什么：

1. __device__ int devAr[1];在设备内存中分配固定大小的数组，并将指向该设备内存的指针存储在devAr变量中（因此发出警告）
2. devAr地址指向有效的设备内存，但是，即使在主机代码中也可以使用这样的地址，因为主机和设备内存使用相同格式的地址。然而，在主机代码中，devAr指向一些随机的未初始化的主机内存
3. 基于以上，可以说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线程，大多数引用都来自答案下面的评论，它们是：

1. cudaMemcpy（）与cudaMemcpyFromSymbol（）：

   任何静态定义的设备符号（__device__、__device__0，甚至纹理）都会导致工具链发出两个符号，一个在设备模块中，另一个在主机对象中。CUDA运行时设置并维护这两个符号之间的动态映射。符号API调用是检索__constant__和__device__符号的映射的方式。纹理API检索纹理符号等的映射。

2. CUDA中全局内存与常量内存的使用：

   *PNT是__device__变量，而不是包含设备变量地址的主机变量。（我知道这很困惑。）因此，如果你像(void**)&PNT一样试图在主机上访问它，你就是在试图从主机读取一个设备变量，这是不允许的。从主机代码的角度来看，它只是一个符号，因此您需要使用cudaGetSympolAddress()将设备地址存储在一个主机变量中，然后可以将其传递给cudaMemcpyToSymbol()，如@talonmies所示。

3. CUDA恒定内存错误：

   有些令人困惑的是，主机代码中的A和B不是有效的设备内存地址。它们是主机符号，为运行时设备符号查找提供挂钩。将它们传递给内核是非法的——如果你想要它们的设备内存地址，你必须在运行时使用cudaGetSymbolAddress来检索它。

4. cudaMemcpyToSymbol与cudaMemcpy为什么它仍然存在（cudaMemcpyToSymbol）：

   通过CUDA API复制到该地址将失败，并出现无效参数错误，因为它不是API之前分配的GPU内存空间中的地址。是的，这也适用于一般的__device__指针和静态声明的设备符号。

5. devAr0变量上的cudaMemcpyFromSymbol：

   问题的根源是不允许在普通主机代码中获取设备变量的地址：。。。尽管这似乎是正确编译的，但实际传递的地址是垃圾。要在主机代码中获取设备变量的地址，我们可以使用cudaGetSymbolAddress

基于这些证据，让我尝试更新我最初的3步解释：

1. __device__ int devAr[1];在设备内存中分配固定大小的数组，并将"挂钩到运行时设备符号查找"存储到devAr变量的主机版本中（请参阅链接资源1和3）
2. 从主机的角度来看，devAr地址只是一个垃圾，只应与符号API调用一起使用，如cudaGetSymbolAddress（所有链接的资源似乎都支持这一理论），因为它映射到devAr变量的设备版本

我没能想出任何"更具体"的东西，比如CUDA文档的链接，但我希望现在已经足够清楚了。总而言之，您现在似乎对上述行为有了保证（即存在devAr变量的主机和设备版本），但对我来说，它只是一个实现细节，您不应该依赖它，也不应该将devAr变量的主机版本用于符号API调用之外的目的。