分配共享内存

我想通过使用一个常量参数分配共享内存，但得到一个错误。 我的内核看起来像这样：

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

我得到一个错误说

  error: expression must have a constant value 

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

Answer 1

const并不意味着“常量”，它的意思是“只读”。

常量expression式是编译器在编译时已知的值。

Answer 2

CUDA支持dynamic共享内存分配。 如果你像这样声明内核：

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

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

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

那么可以在运行时resize。 请注意运行时只支持每个块的一个dynamic声明的分配。 如果你需要更多，你需要使用指向偏移量的指针。 当使用共享内存使用32位字的指针时，请注意，不pipe共享内存分配的types如何，所有分配都必须是32位字alignment的。

Answer 3

选项一：用常量声明共享内存（与const不一样）

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

选项二：在内核启动configuration中dynamic声明共享内存：

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

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

Answer 4

从“CUDA C编程指南”：通过插入expression式来指定执行组态i：

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

其中：

• Dg是dim3types，并指定网格的尺寸和大小…
• Dbtypes为dim3 ，指定每个块的尺寸和大小…
• Ns的types为size_t，并指定共享内存中的字节数，除了静态分配的内存以外，每个块为此调用dynamic分配的字节数。 这个dynamic分配的内存被任何声明为外部数组的variables使用，如__shared__ ; Ns是可选参数，默认为0;
• S的types是cudaStream_t并指定关联的stream…

所以通过使用dynamic参数Ns，用户可以指定一个内核函数可以使用的共享内存的总大小，而不pipe内核中有多less个共享variables。

Answer 5

你不能像这样声明共享variables

 __shared__ int a[count]; 

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

 __shared__ int a[100]; 

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

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

 extern __shared__ int a[]; 

并从内存中调用内核时分配内存

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

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