42

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

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

我收到一个错误消息

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

计数是常量！为什么我会收到此错误？我怎样才能解决这个问题？

4

5 回答 5

93

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

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

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

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

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

于 2011-04-03T18:40:42.510 回答
39

const并不意味着“恒定”，而是意味着“只读”。

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

于 2011-04-03T17:36:06.693 回答
21

选项一：用常量值声明共享内存（不一样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);

注意：指向动态共享内存的指针都被赋予了相同的地址。我使用两个共享内存数组来说明如何在共享内存中手动设置两个数组。

于 2011-04-04T16:57:39.480 回答
4

来自“CUDA C 编程指南”：

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

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

在哪里：

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

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

于 2017-08-09T03:10:09.430 回答
1

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

__shared__ int a[count];

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

__shared__ int a[100];

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

这个问题有一个很好的解决方案来声明 

extern __shared__ int a[];

并在从内存中调用内核时分配内存，例如

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

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

于 2011-04-04T07:00:21.417 回答