以前在 CUDA C++ 中定义的外部全局设备变量

extern global device variable previously defined in cuda c++

本文关键字：全局 外部 变量 定义 CUDA C++      更新时间：2023-10-16

我有三个文件：a.cu、b.cu 和 c.h。我想让所有编译单元都可以访问调试变量，所以我在 c.h 中将它们声明为 extern，并且仅在 a.cu 中定义，根据此处的答案 多个文件中的全局变量：

a.cu

#include "c.h"
#include <stdio.h>
#include <iostream>
bool h_debug;                                                                       //works fine
//__device__ bool d_debug;                     //POINT A
int main(int argc, char* argv[])
{
    h_debug = (argc > 1 ? true : false);
    std::cout<<"Host: "<<(h_debug ? "true" : "false")<<std::endl;
    cudaMemcpyToSymbol(d_debug, &h_debug, sizeof(bool));  
    std::cout<<cudaGetErrorString(cudaGetLastError())<<std::endl; 
    cudaDeviceSynchronize();
}

b.cu

#include "c.h"
#include "stdio.h"
__global__ 
void myKernel(){
    if(d_debug){
        printf("device debug onn");
    }
    else{
        printf("device debug offn");
    }
}

C.H

#ifndef MAIN_H
#define MAIN_H
extern __device__ bool d_debug;
extern bool h_debug;
#endif  /* MAIN_H */

在标记为 POINT A 的行中，如果我对其进行注释，代码将编译，但在运行时出现 cuda 错误，正如人们所期望的那样： 

$ nvcc a.cu b.cu -o globalTest
$ globalTest
Host: true
invalid device symbol

如果我取消注释该行以定义d_debug，则会出现一个对我来说没有意义的编译器错误...... 

$ nvcc a.cu b.cu -o globalTest
a.cu:1:32: warning: unknown option after ‘#pragma GCC diagnostic’ kind [-Wpragmas]
a.cu:6:13: error: redefinition of ‘bool d_debug’
c.h:4:13: error: ‘bool d_debug’ previously declared here

为什么它不像全局主机变量那样工作？我应该如何创建一个所有编译单元都可以访问的全局设备变量？

来自罗伯特·克罗维拉的评论和这里的答案： CUDA 常量内存值不正确， 问题是 a.cu 和 b.cu 有单独的d_debug副本。为了防止这种情况，必须使用单独的编译和链接;为此，请创建可重定位的设备代码（而不是可执行的设备代码）。

--

relocatable-device-code {true|false}
-rdc 启用（禁用）可重定位设备代码的生成。如果禁用，则可执行设备 代码已生成。在可重定位设备代码之前必须链接 可以执行。

此选项允许的值：真、假。

默认值：假

将 nvcc 参数更改为： 

nvcc -rdc=true a.cu b.cu -o globalTest

解决了问题。

相关文章：