因此,我试图编写一些利用Nvidia的CUDA架构的代码。我注意到往返于设备的复制确实损害了我的整体性能,所以现在我正尝试将大量数据移至设备上。
由于此数据用于多种功能,我希望它具有全局性。是的,我可以传递指针,但是我真的很想知道在这种情况下如何使用全局变量。
因此,我具有要访问设备分配的阵列的设备功能。
理想情况下,我可以执行以下操作:
1 2 3 4 5 6 7 8
| __device__ float* global_data;
main()
{
cudaMalloc(global_data);
kernel1<<<blah>>>(blah); //access global data
kernel2<<<blah>>>(blah); //access global data again
} |
但是,我还没有弄清楚如何创建一个动态数组。我通过声明数组来找出解决方法,如下所示:
1
| __device__ float global_data[REALLY_LARGE_NUMBER]; |
尽管不需要cudaMalloc调用,但我更喜欢动态分配方法。
类似的事情可能应该起作用。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44
| #include
#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.\
", \\
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;
} |
旋转一下。
我继续尝试分配临时指针并将其传递给类似于kernel1的简单全局函数的解决方案。
好消息是它确实起作用:)
但是,当我尝试访问全局数据时,由于我现在得到"建议:假设全局内存空间,无法告诉指针指向什么",我认为它会使编译器感到困惑。幸运的是,这个假设恰好是正确的,但是警告很烦人。
无论如何,为了记录-我查看了许多示例,并进行了nvidia练习,其中的重点是使输出说"正确!"。但是,我还没有全部查看。如果有人知道他们在其中进行动态全局设备内存分配的sdk示例,我仍然想知道。
花一些时间专注于NVIDIA提供的丰富文档。
摘自《编程指南》:
1 2 3
| float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr)); |
这是如何分配内存的简单示例。现在,在您的内核中,您应该接受一个指向浮点数的指针,如下所示:
1 2 3 4 5 6 7 8 9 10 11
| __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;
} |
所以现在您可以像这样调用它们:
1 2 3 4 5 6
| 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()来实现您的目标。但是,在后台,请考虑另一种方法可能是正确的事情。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34
| #include
__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'。
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结构上运行,该怎么办?不能使用所谓的常量内存,因为它的大小非常有限。因此,您必须将其放入全局内存中。
签出SDK随附的示例。这些示例项目中的许多都是通过示例进行学习的一种不错的方法。
Erm,正是将devPtr移到全局范围的问题是我的问题。
我有一个实现该功能的实现,两个内核都有一个指向数据的指针。我明确地不想传入这些指针。
我已经非常仔细地阅读了文档,并访问了nvidia论坛(谷歌搜索了一个小时左右),但是我没有找到实际运行的全局动态设备数组的实现(我已经尝试了几个编译,然后以新的有趣方式失败。)