CUDA 中的全局变量

Question

如何在 CUDA 中创建全局变量？ 可以举个例子吗？

例如，如何在 CUDA 函数中创建数组

__global__ void test()
{
  int *a = new int[10];
}

或者我怎样才能创建一个全局数组并通过这个函数访问它。例如

__device__ int *a;
__global__ void test()
{
  a[0] = 2;
}

或者我怎样才能像下面这样使用..

__global__ void ProcessData(int img)
{
   int *neighborhood = new int[8]; 
   getNeighbourhood(img, neighbourhood);
}

---

我仍然有一些问题。我发现与

相比

__device__

如果我定义

"__device__ __constant__" (read only)

将改善内存访问。 但我的问题是我在主机内存中有一个数组说

 float *arr = new float[sizeOfTheArray]; 

我想将它作为设备中的变量数组，我需要在设备内存中修改它，我需要将它复制回主机。我该怎么做？？

Answer 1

CUDA 4.0 的计算能力 2.0 和 2.1（即 Fermi）支持 C++ new 运算符，因此您可以使用 new 将全局内存分配到设备符号上，尽管前两个代码 sn -ps 是它在实践中的实现方式。

在较旧的硬件和/或 CUDA 4.0 之前的工具包上，标准方法是在主机代码中使用 cudaMemcpyToSymbol API：

__device__ float *a;

int main()
{
    const size_t sz = 10 * sizeof(float);

    float *ah;
    cudaMalloc((void **)&ah, sz);
    cudaMemcpyToSymbol("a", &ah, sizeof(float *), size_t(0),cudaMemcpyHostToDevice);
}

将动态分配的设备指针复制到可以直接在设备代码中使用的符号上。

---

编辑：回答这个问题有点像击中一个移动的目标。对于您现在似乎感兴趣的常量记忆案例，这是一个完整的工作示例：

#include <cstdio>

#define nn (10)

__constant__ float a[nn];

__global__ void kernel(float *out)
{
    if (threadIdx.x < nn)
        out[threadIdx.x] = a[threadIdx.x];

}

int main()
{
    const size_t sz = size_t(nn) * sizeof(float);
    const float avals[nn]={ 1., 2., 3., 4., 5., 6., 7., 8., 9., 10. };
    float ah[nn];

    cudaMemcpyToSymbol("a", &avals[0], sz, size_t(0),cudaMemcpyHostToDevice);

    float *ad;
    cudaMalloc((void **)&ad, sz);

    kernel<<<dim3(1),dim3(16)>>>(ad);

    cudaMemcpy(&ah[0],ad,sz,cudaMemcpyDeviceToHost);

    for(int i=0; i<nn; i++) {
        printf("%d %f\n", i, ah[i]);
    }
}

这显示了将数据复制到一个常量内存符号上，并在内核中使用该数据。