分配共享内存

Question

我正在尝试使用常量参数分配共享内存，但是遇到了错误。我的内核看起来像这样：

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

我收到了一个错误信息：

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

count是const类型的！为什么会出现这个错误？我该如何绕过此错误？

- lina

Answer 1

CUDA支持动态共享内存分配。如果你像这样定义内核：

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

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

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

然后它可以在运行时调整大小。请注意，运行时仅支持每个块的单个动态声明分配。如果您需要更多，则需要使用指向该单个分配内部偏移量的指针。另外，请注意在使用指针时，共享内存使用32位字，所有分配必须是32位字对齐的，无论共享内存分配的类型如何。

- talonmies

---

在sm_21上，我有一种印象，即a_size具有128字节的粒度：启动使用'a_size = 1'（需要时为16）的内核似乎运行良好，并且对于大的a_size - 127也能正常工作，而a_size - 128则会显示过小分配的影响。（我相当确信“良好”不是指“偶然”。）这是否与L1行大小= 128字节有关？ - P Marecki

@talonmies 是否可以使用模板来传递 a 长度的参数，还是这样做行不通？ - BugShotGG

@GeoPapas：是的，只要没有使用extern参数。我建议使用内核的模板实例，而不是动态分配共享内存。当数组大小在编译时已知时，编译器可以更好地优化代码。 - talonmies

谢谢。这个内存会自动释放吗？ - astrowalker

Answer 2

const 不是指 "常量"，它表示 "只读"。

常量表达式是指编译器在编译时已知其值的表达式。

- Oliver Charlesworth

---

这只回答了问题的第一部分。 - Tharindu Rusira

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);

注意：动态共享内存的指针全部具有相同的地址。我使用两个共享内存数组来说明如何手动设置共享内存中的两个数组。

- jmilloy

---

除非在内核中使用extern关键字定义共享内存分配，否则选项二不起作用。 - talonmies

@talonmies 已修复！我们还应该提到，这意味着如果您有两个具有共享内存的内核，则应为共享内存指针使用不同的标识符。这并非完全必要，但可以帮助避免一些错误。 - jmilloy

@talonmies 当我在一个内核中使用extern __shared__ int *shared，而在另一个内核中使用extern __shared__ float *shared时，nvcc编译失败。 - jmilloy

Answer 4

从“CUDA C编程指南”中：

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

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

参数说明：

• Dg 类型为 dim3，指定网格的维度和大小...
• Db 类型为 dim3，指定每个块(block)的维度和大小...
• Ns 类型为 size_t，指定共享内存中分配给每个块的动态分配内存的字节数。这些变量必须在__shared__中声明为外部数组; Ns 是可选参数，默认值为0;
• S 类型为 cudaStream_t，指定关联的流...

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

- smh

Answer 5

你不能像这样声明共享变量。

__shared__ int a[count];

尽管如果您对数组a的最大大小足够确定，那么可以直接声明如下：

__shared__ int a[100];

但在这种情况下，您应该担心程序中有多少个块，因为将共享内存固定到一个块（而没有充分利用）会导致与全局内存（高延迟）进行上下文切换，从而导致性能不佳...

对于这个问题有一个好的解决方案，就是声明

extern __shared__ int a[];

在从内存调用内核时分配内存

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

但是你在这里也应该感到困扰，因为如果你在块中使用的内存比你在内核中分配的内存还要多，那么你将会得到意想不到的结果。

- peeyush

---

为什么First是无意义的？我们能否使用某些变量声明共享内存？ - peeyush

这是无稽之谈，因为(a)共享内存按块作用域声明，每个分配都与块相关联，(b)在CUDA中并不存在“使用全局内存进行上下文切换”的情况。 - talonmies

如果每个SM有16KB的共享内存，并且我正在使用4KB的共享内存（外部共享内存），现在如果一个块被调度到SM上，则它将使用4KB的共享内存，但这还不够，好的，现在另外7个块被调度到共享内存上，每个块都使用4KB，那么这将如何处理？当有8个块要被调度时，CUDA如何管理4KB的共享内存？我同意SM处于块作用域，通过shared int [count]，我对变量计数有异议，因为变量不能用于声明共享内存.. 我错了吗？ - peeyush

GPU只会调度适合可用寄存器和共享内存空间的块数。如果一个块需要4kb的共享内存，则如果SM有16kb的共享内存，那么每个SM最多只能有4个块处于活动状态。我想我应该说你的第一段话“大多无稽之谈”，因为是的，前10个单词是正确的，但在那之后就变成了幻想。 - talonmies