如何使用外部的CUDA设备变量

Question

我需要将代码写入多个 .cu 文件。但是在哪里定义设备变量，它们用于许多 .cu 文件。

例如

文件 common.h

__device__ int x;

File A.cu

__global__ void a() 

文件 B.cu

__global__ void b() 

a()和b()都使用了x。我该怎么办？

在C语言中，我应该这样写：

extern <strong>device</strong> int x;

然后我在另一个地方定义device int x。但是在CUDA中，我不能这样做。如果我这样做，它会告诉我“..........”之前已经声明过。

- worldterminator

Answer 1

编辑：@talonmies是正确的（像往常一样）。因此，我删除了关于CUDA 4.1的评论。

此外，我给出的编译命令不太正确。所以让我用一个明显有效且具有适当说明的答案来替换我的答案。

您需要CUDA 5.0和计算能力2.0或更高的设备才能使用此功能。

我相信可能有更好的方法，但这对我来说似乎有效：

com.h：

#ifndef DEVMAIN
extern __device__ int x;
#endif

a.cu:

#include "com.h"
__global__ void a(){

  x = -5;
}

b.cu:

#include <stdio.h>
#define DEVMAIN
#include "com.h"

extern __global__ void a();
__device__ int x;

__global__ void b(){

  x = 5;
}

int main() {
  int temp=7;
  cudaMemcpyToSymbol(x,&temp, sizeof(int));
  a<<<1,1>>>();
  cudaMemcpyFromSymbol(&temp,x,sizeof(int));
  printf("in host : %d\n",temp);
  b<<<1,1>>>();
  cudaMemcpyFromSymbol(&temp,x,sizeof(int));
  printf("in host2 : %d\n",temp);
  return 0;
}

编译中：

nvcc -arch=sm_20 -dc a.cu
nvcc -arch=sm_20 -dc b.cu
nvcc -arch=sm_20 -o ab a.o b.o

输出：

$ ./ab
in host : -5
in host2 : 5
$

很抱歉之前我的错误。

- Robert Crovella

---

值得指出的是，这可能只在具有适当设备代码链接器的CUDA 5中合法。 - talonmies

@RobertCrovella 我试了你的代码，你能否试试这段代码？它应该是1,-1，但实际上是1,1。`int main() { int temp; b<<<1,1>>>(); cudaMemcpyFromSymbol(&temp,x,sizeof(int),0,D_T_H); printf("在主机上 : %d\n",temp); a<<<1,1>>>(); cudaMemcpyFromSymbol(&temp,x,sizeof(int),0,D_T_H); printf("在主机上2 : %d\n",temp); cudaDeviceSynchronize(); return 0; }` - worldterminator

@RobertCrovella和D_T_H是#define D_T_H cudaMemcpyDeviceToHost。 - worldterminator

@worldterminator 我之前的回答有几个错误，我已经修正了（希望如此），并且编辑了我的回答以反映这些更改。 - Robert Crovella

没错，它可以工作！顺便问一下，-dc是什么意思？和-c有什么区别？ - worldterminator

您可以在此处找到答案：http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#code-changes-for-separate-compilation。使用CUDA 5，我们现在可以分离链接设备代码。为了做到这一点，编译的设备代码必须以可重定位格式交付给链接器。 -dc指示编译器生成可重定位设备代码，稍后可以将其链接。 - Robert Crovella