因此，我试图编写一些利用Nvidia CUDA体系结构的代码。我注意到往返于设备的复制确实损害了我的整体性能，因此现在我正尝试将大量数据移至设备上。

由于此数据用于多种功能，我希望它具有全局性。是的，我可以传递指针，但是我真的很想知道在这种情况下如何使用全局变量。

因此，我具有要访问设备分配的阵列的设备功能。

理想情况下，我可以执行以下操作：

__device__ float* global_data;

main()
{
  cudaMalloc(global_data);
  kernel1<<<blah>>>(blah); //access global data
  kernel2<<<blah>>>(blah); //access global data again
}

但是，我还没有弄清楚如何创建一个动态数组。我通过声明数组来找出解决方法，如下所示：

__device__ float global_data[REALLY_LARGE_NUMBER];

尽管不需要cudaMalloc调用，但我更喜欢动态分配方法。

解决方案

回答

花一些时间专注于NVIDIA提供的丰富文档。

从《编程指南》中：

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

这是如何分配内存的简单示例。现在，在内核中，我们应该接受一个指向浮点数的指针，如下所示：

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x]++;
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= 0.3f;
}

因此，现在我们可以像这样调用它们：

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

kernel1<<<1,128>>>(devPtr);
kernel2<<<1,128>>>(devPtr);

As this data is used in numerous
  functions, I would like it to be
  global.

使用全局变量的理由很少。这绝对不是一个。我将保留该示例的扩展范围，以将" devPtr"移至全局范围。

编辑：

好的，根本的问题是这样的：内核只能访问设备内存，并且它们只能使用的全局范围指针是GPU的。从CPU调用内核时，在后台发生的事情是，在内核执行之前，指针和基元被复制到GPU寄存器和/或者共享内存中。

因此，我可以建议的最接近的方法是：使用cudaMemcpyToSymbol()来实现目标。但是，在后台，请考虑另一种方法可能是正确的事情。

#include <algorithm>

__constant__ float devPtr[1024];

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1];
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= devPtr[2];
}

int main(int argc, char *argv[])
{
    float some_data[256];
    for (int i = 0; i < sizeof(some_data) / sizeof(some_data[0]); i++)
    {
        some_data[i] = i * 2;
    }
    cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr) ));
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    kernel2<<<1,128>>>(otherDevPtr);

    return 0;
}

在此示例中，请不要忘记使用" --host-compilation = c ++"。

回答

嗯，正是将devPtr移到全球范围的问题是我的问题。

我有一个实现该功能的实现，两个内核都有一个指向传入数据的指针。我明确地不想传入这些指针。

我已经非常仔细地阅读了文档，并访问了nvidia论坛(谷歌搜索了一个小时左右)，但是我没有找到实际运行的全局动态设备数组的实现(我尝试了几种编译并然后以新的有趣的方式失败)。

回答

这样的事情可能应该起作用。

#include <algorithm>

#define NDEBUG
#define CUT_CHECK_ERROR(errorMessage) do {                                 \
        cudaThreadSynchronize();                                           \
         cudaError_t err = cudaGetLastError();                             \
         if( cudaSuccess != err) {                                         \
                     fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n",    \
                                             errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
                     exit(EXIT_FAILURE);                                                  \
                 } } while (0)

__device__ float *devPtr;

__global__
void kernel1(float *some_neat_data)
{
    devPtr = some_neat_data;
}

__global__
void kernel2(void)
{
    devPtr[threadIdx.x] *= .3f;
}

int main(int argc, char *argv[])
{
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    CUT_CHECK_ERROR("kernel1");

    kernel2<<<1,128>>>();

    CUT_CHECK_ERROR("kernel2");

    return 0;
}

旋转一下。

回答

查看SDK随附的示例。这些示例项目中的许多都是通过示例学习的不错的方法。

回答

我继续尝试分配临时指针并将其传递给类似于kernel1的简单全局函数的解决方案。

好消息是它确实起作用:)

但是，当我尝试访问全局数据时，由于我现在得到"建议：假设全局内存空间，无法告诉指针指向什么"，因此我认为它会使编译器感到困惑。幸运的是，这种假设是正确的，但是警告很烦人。

无论如何，作为记录，我看了很多示例，并进行了nvidia练习，其中的重点是使输出说"正确！"。但是，我还没有全部查看。如果有人知道执行动态全局设备内存分配的sdk示例，我仍然想知道。

回答

As this data is used in numerous functions, I would like it to be global.

--

There are few good reasons to use globals. This definitely is not one. I'll leave it as an 
  exercise to expand this example to include moving "devPtr" to a global scope.

如果内核在由数组组成的大型const结构上运行，该怎么办？不能使用所谓的常量内存，因为它的大小非常有限。因此，我们必须将其放入全局内存中。