首页 / 知识

关于nvidia：CUDA全局(分配给C)动态数组分配给设备内存

2023-04-16 08:11:00

CUDA global (as in C) dynamic arrays allocated to device memory

因此，我试图编写一些利用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论坛(谷歌搜索了一个小时左右)，但是我没有找到实际运行的全局动态设备数组的实现(我已经尝试了几个编译，然后以新的有趣方式失败。)

内存动态数组数据损害