xjt*_*c55 1 cuda constants declaration gpu-shared-memory
我正在运行 1024 个矩阵的适应度函数,每个矩阵都有自己的块并且大小相同。每个块都有n*n
线程(矩阵的维度)并且需要具有n*n
共享内存,以便我可以轻松进行求和缩减。然而,n
所有矩阵的维数在运行前都是可变的(即可以手动更改,但始终是 2 的幂,因此求和很简单)。这里的问题是共享内存必须使用常量分配,但我还需要将值从主机传递到内核。我在哪里声明维度,n
以便它对 CPU 可见(用于传递给内核)并可用于声明共享内存的大小(在内核内)?
我的代码结构如下:
从main.cu
我调用内核:
const int num_states = 1024
const int dimension = 4
fitness <<< num_states, dimension * dimension >>> (device_array_of_states, dimension, num_states, device_fitness_return);
Run Code Online (Sandbox Code Playgroud)
然后kernel.cu
我有:
__global__ void fitness(
int *numbers,
int dimension,
int num_states,
int *fitness_return) {
__shared__ int fitness[16]; <<-- needs to be dimension * dimension
//code
}
Run Code Online (Sandbox Code Playgroud)
numbers
是表示 1024 个矩阵的数组,dimension
是行和列长度,num_states
是 1024,fitness_return
是长度为 1024 的数组,用于保存每个矩阵的适应度值。在内核中,共享内存是用 的平方进行硬编码的dimension
(dimension
本例中是 4)。
我在哪里以及如何声明dimension
它可以用于分配共享内存以及调用内核,这样我只需dimension
在一处进行更新?感谢您的帮助。
所有块中分配的共享内存量是统一的。您可能在每个块中使用不同数量的共享内存,但它仍然全部可用。而且,共享内存的数量无论如何都是相当有限的,因此 n*n 个元素不能超过最大空间量(通常为 48KiB);对于float
-type 元素(每个 4 个字节),这意味着 n < 340 左右。
现在,有两种分配共享内存的方法:静态和动态。
静态分配是您给出的示例,这是行不通的:
__shared__ int fitness[16];
Run Code Online (Sandbox Code Playgroud)
在这些情况下,必须在编译时(在设备端代码编译时)知道大小- 但您的情况并非如此。
使用动态共享内存分配,您无需在内核代码中指定大小 -将其留空并添加前缀extern
:
extern __shared__ int fitness[];
Run Code Online (Sandbox Code Playgroud)
相反,您在启动内核时指定数量,并且不同块的线程不一定知道它是什么。
但就您而言,线程确实需要知道 n 是什么。好吧,只需将其作为内核参数传递即可。所以,
__global__ void fitness(
int *numbers,
int dimension,
int num_states,
int *fitness_return,
unsigned short fitness_matrix_order /* that's your n*/)
{
extern __shared__ int fitness[];
/* ... etc ... */
}
Run Code Online (Sandbox Code Playgroud)
nVIDIA 的Parallel-for-all博客有一篇很好的文章,其中更深入地介绍了如何使用共享内存,其中专门介绍了静态和动态共享内存分配。
归档时间: |
|
查看次数: |
2132 次 |
最近记录: |