分配共享内存

Question

我试图通过使用常量参数来分配共享内存，但出现错误。我的内核看起来像这样：

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

并且我收到一条错误消息

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

count is const！为什么我会收到此错误？我该如何解决这个问题？

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，用户可以指定一个内核函数可以使用的共享内存的总大小，无论有多少个共享变量这个内核里有。

Answer 5

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

__shared__ int a[count];

虽然如果你对数组 a 的最大大小足够确定那么你可以直接声明，

__shared__ int a[100];

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

这个问题有一个很好的解决方案，可以

extern __shared__ int a[];

在从内存调用内核时声明和分配内存类似

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

，但您也应该在这里感到困扰，因为如果您在块中使用的内存多于在内核中分配的内存，您将得到意想不到的结果。