5赞

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

作者：携手相约幸福 | 2023-09-11 12:41

如何解决《从主机访问CUDA全局设备变量》经验，为你挑选了1个好方法。

我想知道是否有官方消息来源,为什么以下工作:

#include 

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但是通过引用没有这样的警告(有充分理由),则会收到警告"a __device__ variable"devAr"无法在主机函数中直接读取".但在这两种情况下都可以从主机访问变量.所以看来,该变量有一个主机实例.

我需要知道的是:我认为这是理所当然的吗？

其他测试用例显示指针的值:

#include 
#include 

__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

Michal Hosal.. 13

你在做什么是无效的,你应该听警告:

一个__device__变量devAr不能在主机功能被直接读

首先,让我简化您的代码,只显示显示问题所需的大小:

#include 

__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 

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()vs cudaMemcpyFromSymbol():

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

在CUDA中使用全局与常量内存:

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

CUDA常量内存错误:

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

cudaMemcpyToSymbol与cudaMemcpy为何仍然存在(cudaMemcpyToSymbol):

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

cudaMemcpyFromSymbol对一个__device__变量:

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

基于这些证据,我试着从上面更新我原来的3步说明:

__device__ int devAr[1];在设备内存中分配固定大小的数组,并将"钩子运行到运行时设备符号中"存储到devAr变量的主机版本中(参见链接的资源1和3).

该devAr地址是单从主机的角度垃圾,应该只用符号API调用,如使用cudaGetSymbolAddress(所有链接的资源似乎支持这一理论),因为其映射到的设备版本devAr的变量.

我无法提出任何"更具体"的内容,例如链接到CUDA文档,但我希望现在已经足够清楚了.总而言之,您似乎现在可以保证上述行为(即存在devAr变量的主机和设备版本),但对我来说,它更像是一个您不应该依赖的实现细节,不应该使用主机版本的devAr比象征API调用其他目的变量.

1> Michal Hosal..：