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

Question

我正在尝试使用 CUDA 5 进行单独编译。出于这个原因，我在 Visual Studio 2010 中将“生成可重定位设备代码”设置为“是（-rdc=true）”。程序编译没有错误，但是， 当我尝试使用 cudaMemcpyToSymbol 初始化设备常量时出现无效设备符号错误。

即我有以下常量

__constant__ float gdDomainOrigin[2];

并尝试使用

对其进行初始化

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

这会导致错误。当我在没有上述选项集的情况下将所有内容编译为一个整体时，不会发生该错误。有人可以帮我吗？

Answer 1

我无法重现此内容。如果从两个 .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;
}

在带有 Kepler 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 工具链会发生什么。