在使用动态共享内存分配的情况下正确的内核调用答案

【问题标题】：correct kernel call in case of using dynamic shared memory allocation在使用动态共享内存分配的情况下正确的内核调用
【发布时间】：2013-03-07 02:37:37
【问题描述】：

开发人员，

有人可以给我一个提示吗？我没有找到任何关于如何在同一个内核中分配常量和动态共享内存的信息，或者让我们更珍贵地问一下：如何调用内核，其中需要分配的共享内存量在编译时只是部分知道？例如，参考allocating shared memory，如何进行动态分配变得非常明显。但是让我们假设我有以下内核：

__global__ void MyKernel(int Float4ArrSize, int FloatArrSize)
{
  __shared__ float Arr1[256];
  __shared__ char  Arr2[256];
  extern __shared_ float DynamArr[];
  float4* DynamArr1 = (float4*) DynamArr;
  float* DynamArr = (float*) &DynamArr1[Float4ArrSize];

  // do something
}

内核调用：

int SharedMemorySize = Float4ArrSize + FloatArrSize;

SubstractKernel<<< numBlocks, threadsPerBlock, SharedMemorySize, stream>>>(Float4ArrSize, FloatArrSize)

我实际上无法弄清楚编译器如何将共享内存的大小仅链接到我想要动态分配的部分。或者参数“SharedMemeorySize”是否代表了每块共享内存的总量，所以我需要计算常量内存的大小（int SharedMemorySize = Float4ArrSize + FloatArrSize + 256*sizeof(float)+ 256*sizeof(char)） ?

请赐教或只是简单地指向一些代码sn-ps。提前非常感谢。

干杯格雷格

【问题讨论】：

操作！它只需要一个 __shared__ 变量。尝试将所有内容组合到一个结构中。
@SorooshBateni 不，它没有！内核调用中的参数只影响动态分配的共享内存的大小，静态分配的部分不受影响，它们的方式很好。
是的，我说动态分配单元 (:D) 只需要一个 __shared__ 变量。

标签： c++ cuda shared-memory

【解决方案1】：

引用编程指南，SharedMemorySize 指定除了静态分配的内存外，每个块动态分配的共享内存中的字节数；这个动态分配的内存被任何声明为外部数组的变量。 SharedMemorySize 是一个可选参数，默认为 0。

所以如果我明白你想要做什么，它应该看起来像

extern __shared_ float DynamArr[];
float*  DynamArr1 = DynamArr;
float4* DynamArr2 = (float4*) &DynamArr[DynamArr1_size];

请注意，我没有测试它。

Here 是非常有用的帖子。

【讨论】：

谢谢你，完美...现在我知道如何继续了！
欢迎您@GregPhil。顺便说一句，在 stackoverflow 上向某人表示感谢的最佳方式是对他/她的帖子进行投票，而且，如果存在接受最合适的答案，这是一个很好的习惯。

【解决方案2】：

来自CUDA programming guide：

[内核的]执行配置是通过插入一个表达式来指定的函数名和函数名之间的形式 >> 带括号的参数列表，其中：

Ns 是 size_t 类型，并指定共享内存中的字节数，该字节数在每个块中为此调用动态分配除了静态分配的内存；这是动态的分配的内存被任何声明为 __shared__ 中提到的外部数组； Ns 是可选的默认为 0 的参数；

所以基本上，您在内核调用期间指定的共享内存大小与动态分配的共享内存有关。您不必在共享内存中手动添加静态分配数组的大小。

【讨论】：

非常感谢，这正是我想知道的！