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

Correct kernel call in case of using dynamic shared memory allocation

本文关键字：情况下 内核 调用 分配 动态 共享 内存      更新时间：2023-10-16

我没有找到任何关于如何在同一内核中分配静态和动态共享内存的信息，或者让我们更准确地问：

如何调用在编译时仅部分知道需要分配的共享内存量的内核？

例如，提到分配共享内存，如何进行动态分配变得非常明显。

但是让我们假设我有以下内核：

__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)
实际上，我

无法弄清楚编译器如何将共享内存的大小仅链接到我想要动态分配的部分。还是参数SharedMemorySize表示每个块的共享内存总量，所以我需要计算静态共享内存的大小(int SharedMemorySize = Float4ArrSize + FloatArrSize + 256*sizeof(float)+ 256*sizeof(char)(？

请启发我或只是简单地指出一些代码片段。提前非常感谢。

干杯格雷格

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

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

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

请注意，我没有测试它。

这是非常有用的帖子。

来自 CUDA 编程指南：

[内核] 的执行配置是通过插入表达式函数名称和<<< Dg, Db, Ns, S >>>括号参数列表，其中：

  • [...]
  • Ns 的类型为 size_t，除了静态分配的内存外，还指定了共享内存中每个块为此调用动态分配的字节数; 此动态分配的内存由声明为外部数组的任何变量使用，如__shared__中所述; Ns 是一个可选参数，默认为 0;

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

相关文章：