CUDA 5.0“生成可重定位的设备代码”导致无效的设备符号错误

我想单独编译使用CUDA 5.为此，我将“生成可重定位设备代码”设置为“是（-rdc = true）”在Visual Studio 2010.该程序编译时无错误，但是，当我尝试使用cudaMemcpyToSymbol初始化设备常量时，出现无效的设备符号错误。CUDA 5.0“生成可重定位的设备代码”导致无效的设备符号错误

即我有以下的不断

__constant__ float gdDomainOrigin[2];

，并尝试与

cudaMemcpyToSymbol(gdDomainOrigin, mDomainOrigin, 2*sizeof(float));

从而导致错误进行初始化。如果没有前面提到的选项集，编译所有内容时都不会发生错误。有人可以帮我吗？

来源

2013-03-18 scttrbrn

根据您的代码结构，您[可能需要]使用（http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#code-changes-for-separate-compilation ）'extern'或'static'来声明符号的可见性。您还没有提供足够的信息，说明哪些模块正在声明符号以及哪些模块正在引用它们。 – 2013-03-18 02:52:44

感谢您的回复。上述两个代码片段都在相同的.cu文件中。此外，该符号仅在此文件中引用。但是，我想用extern声明一些曾经是.cu文件一部分的内核，并在不同的.cu文件中定义它们以使我的代码更易读。但是，这些内核不能访问上面提到的符号。 – scttrbrn 2013-03-18 03:39:10

我不能重现这一点。如果从两个.cu文件构建应用程序，一个包含__constant__符号和一个简单内核，另一个包含运行时API咒语来填充该常量内存并调用内核，则仅在启用可重定位设备代码时才能工作，即：

__constant__ float gdDomainOrigin[2]; 

__global__ 
void kernel(float *inout) 
{ 
    inout[0] = gdDomainOrigin[0]; 
    inout[1] = gdDomainOrigin[1]; 
}

和

#include <cstdio> 

extern __constant__ float gdDomainOrigin; 
extern __global__ void kernel(float *); 

inline 
void gpuAssert(cudaError_t code, char * file, int line, bool Abort=true) 
{ 
    if (code != 0) { 
     fprintf(stderr, "GPUassert: %s %s %d\n", 
          cudaGetErrorString(code),file,line); 
     if (Abort) exit(code); 
    }  
} 
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 

int main(void) 
{ 
    const float mDomainOrigin[2] = { 1.234f, 5.6789f }; 
    const size_t sz = sizeof(float) * size_t(2); 

    float * dbuf, * hbuf; 

    gpuErrchk(cudaFree(0)); 
    gpuErrchk(cudaMemcpyToSymbol(gdDomainOrigin, mDomainOrigin, sz)); 
    gpuErrchk(cudaMalloc((void **)&dbuf, sz)); 

    kernel<<<1,1>>>(dbuf); 
    gpuErrchk(cudaPeekAtLastError()); 

    hbuf = new float[2]; 
    gpuErrchk(cudaMemcpy(hbuf, dbuf, sz, cudaMemcpyDeviceToHost)); 

    fprintf(stdout, "%f %f\n", hbuf[0], hbuf[1]); 

    return 0; 
}

编译和用开普勒GPU 64位Linux系统上在CUDA 5运行这些产生以下：

$ nvcc -arch=sm_30 -o shared shared.cu shared_dev.cu 
$ ./shared 
GPUassert: invalid device symbol shared.cu 23 

$ nvcc -arch=sm_30 -rdc=true -o shared shared.cu shared_dev.cu 
$ ./shared 
1.234000 5.678900

您可以看到，在第一次编译时，没有可重定位的GPU代码生成，找不到符号。在第二种情况下，重新定位的GPU代码生成，它被发现，并在目标文件中ELF头看起来就像你所期望的：

$ nvcc -arch=sm_30 -rdc=true -c shared_dev.cu 
$ cuobjdump -symbols shared_dev.o 

Fatbin elf code: 
================ 
arch = sm_30 
code version = [1,6] 
producer = cuda 
host = linux 
compile_size = 64bit 
identifier = shared_dev.cu 

symbols: 
STT_SECTION  STB_LOCAL .text._Z6kernelPf 
STT_SECTION  STB_LOCAL .nv.constant3 
STT_SECTION  STB_LOCAL .nv.constant0._Z6kernelPf 
STT_CUDA_OBJECT STB_LOCAL _param 
STT_SECTION  STB_LOCAL .nv.callgraph 
STT_FUNC   STB_GLOBAL _Z6kernelPf 
STT_CUDA_OBJECT STB_GLOBAL gdDomainOrigin 

Fatbin ptx code: 
================ 
arch = sm_30 
code version = [3,1] 
producer = cuda 
host = linux 
compile_size = 64bit 
compressed 
identifier = shared_dev.cu 
ptxasOptions = --compile-only

也许你可以试试我的代码和编译/诊断步骤，看看你的Windows工具链会发生什么？

来源

2013-03-18 15:09:04 talonmies

感谢您的努力并发布了这些示例文件，它帮助我解决了这个问题！其实我（愚蠢的）错误是编译计算能力2.0而不是3.0 – scttrbrn 2013-03-24 05:11:09

CUDA 5.0“生成可重定位的设备代码”导致无效的设备符号错误

回答

相关问题