分配共享内存

Question

我试图通过使用常量参数但获得错误来分配共享内存.我的内核看起来像这样:

__global__ void Kernel(const int count)
{
    __shared__ int a[count];
}

我收到一个错误说

错误:表达式必须具有常量值

伯爵是常量!为什么我收到此错误？我怎么能绕过这个？

Answer 1

CUDA支持动态共享内存分配.如果你像这样声明内核:

__global__ void Kernel(const int count)
{
    extern __shared__ int a[];
}

然后传递所需的字节数作为内核启动的第三个参数

Kernel<<< gridDim, blockDim, a_size >>>(count)

然后它可以在运行时调整大小.请注意,运行时仅支持每个块的单个动态声明分配.如果需要更多,则需要在单个分配中使用指向偏移的指针.使用指针时,请注意共享内存使用32位字,并且所有分配必须是32位字对齐,而不管共享内存分配的类型如何.

Answer 2

const 并不意味着"不变",它意味着"只读".

常量表达式是编译器在编译时知道其值的东西.

Answer 3

选项一:声明具有常量值的共享内存(不相同const)

__global__ void Kernel(int count_a, int count_b)
{
    __shared__ int a[100];
    __shared__ int b[4];
}

选项二:在内核启动配置中动态声明共享内存:

__global__ void Kernel(int count_a, int count_b)
{
    extern __shared__ int *shared;
    int *a = &shared[0]; //a is manually set at the beginning of shared
    int *b = &shared[count_a]; //b is manually set at the end of a
}

sharedMemory = count_a*size(int) + size_b*size(int);
Kernel <<<numBlocks, threadsPerBlock, sharedMemory>>> (count_a, count_b);

注意:动态共享内存指针是所有给予相同的地址.我使用两个共享内存数组来说明如何在共享内存中手动设置两个数组.

Answer 4

摘自《CUDA C编程指南》：

执行配置通过插入以下形式的表达式来指定：

<<<Dg, Db, Ns, S>>>

在哪里：

• Dg的类型为dim3，指定网格的维度和大小...
• Db的类型为dim3，指定每个块的维度和大小...
• Ns的类型为size_t，指定除了静态分配的内存之外，为此调用每块动态分配的共享内存中的字节数。此动态分配的内存可供声明为外部数组的任何变量使用，如__shared__中所述；Ns 是可选参数，默认为 0；
• S 的类型为cudaStream_t并指定关联的流...

因此，通过使用动态参数 Ns，用户可以指定一个内核函数可以使用的共享内存的总大小，无论该内核中有多少个共享变量。