具有动态共享内存的模板化 CUDA 内核答案

【问题标题】：Templated CUDA kernel with dynamic shared memory具有动态共享内存的模板化 CUDA 内核
【发布时间】：2015-02-18 15:43:43
【问题描述】：

我想在一个程序中调用具有动态分配共享内存的模板化 CUDA 内核的不同实例。我的第一个天真的方法是写：

template<typename T>
__global__ void kernel(T* ptr)
{
  extern __shared__ T smem[];
  // calculations here ...                                                                                                                                          
}

template<typename T>
void call_kernel( T* ptr, const int n )
{
  dim3 dimBlock(n), dimGrid;
  kernel<<<dimGrid, dimBlock, n*sizeof(T)>>>(ptr);
}

int main(int argc, char *argv[])
{
  const int n = 32;
  float *float_ptr;
  double *double_ptr;
  cudaMalloc( (void**)&float_ptr, n*sizeof(float) );
  cudaMalloc( (void**)&double_ptr, n*sizeof(double) );

  call_kernel( float_ptr, n );
  call_kernel( double_ptr, n ); // problem, 2nd instantiation

  cudaFree( (void*)float_ptr );
  cudaFree( (void*)double_ptr );
  return 0;
}

但是，无法编译此代码。 nvcc 给了我以下错误信息：

main.cu(4): error: declaration is incompatible with previous "smem"
(4): here
          detected during:
            instantiation of "void kernel(T *) [with T=double]"
(12): here
            instantiation of "void call_kernel(T *, int) [with T=double]"
(24): here

我知道我遇到了名称冲突，因为共享内存被声明为 extern。然而，据我所知，如果我想在运行时定义它的大小，那是没有办法的。

所以，我的问题是：有什么优雅的方式来获得所需的行为吗？优雅是指没有代码重复等。

【问题讨论】：

可能是 CUDA 编译器的疏忽，因为 C++ 中允许这样做（没有 __shared__ 限定符）。

标签： c++ cuda

【解决方案1】：

动态分配的共享内存实际上只是一个大小（以字节为单位）和一个为内核设置的指针。所以这样的事情应该可以工作：

替换这个：

extern __shared__ T smem[];

用这个：

extern __shared__ __align__(sizeof(T)) unsigned char my_smem[];
T *smem = reinterpret_cast<T *>(my_smem);

您可以在programming guide 中查看其他重新转换动态分配的共享内存指针的示例，这些示例可以满足其他需求。

编辑：更新了我的答案以反映@njuffa 的评论。

【讨论】：

保守地说，不应该质疑可能的指针对齐问题吗？我通常将my_smem[] 声明为double2 类型以确保16 字节对齐，然后将指针转换为T 类型。
这引发了一个问题，即传递给线程块的动态共享内存分配过程创建的指针是否是 16 字节对齐的（或有任何对齐）。在我看来它可能会，但由于我不知道它是指定的，我同意你的方式似乎更好。当然，全局内存分配具有定义的对齐方式，甚至超过任何向量类型的对齐方式。修改了我的答案。
根据 CUDA 文档，我不知道任何对齐保证，这就是为什么我一直使用 double2 方法作为保守方法。当然，__align__ 属性的使用也应该有效，并且可以说更干净。
我认为这行不通，因为模板的不同实例化的对齐规范是不同的。另外，我认为这是不必要的，因为（开始）共享内存应该很好地对齐。还有@njuffa，你怎么看？
@einpoklum 我想说的都说完了：我的方法是将动态共享内存声明为double2 类型以强制16 字节对齐，然后转换double2 指针指向程序需要的任何T 的指针。当你说“我认为这行不通”时，不清楚“这个”指的是什么；你需要更具体。

【解决方案2】：

_{（@RobertCrovella 的 answer 的变体）}

NVCC 不愿意接受两个同名但类型不同的 extern __shared__ 数组——即使它们从未在彼此的范围内。我们需要通过让我们的模板实例在底层使用相同类型的共享内存来满足 NVCC，同时让使用它们的内核代码看到它喜欢的类型。

所以我们替换这条指令：

extern __shared__ T smem[];

用这个：

auto smem = shared_memory_proxy<T>();

地点：

template <typename T>
__device__ T* shared_memory_proxy()
{
    // do we need an __align__() here? I don't think so...
    extern __shared__ unsigned char memory[];
    return reinterpret_cast<T*>(memory);
}

在某些设备端代码包含文件中。

优点：

使用现场的单线。
更容易记住的语法。
关注点分离 - 阅读内核的人不必考虑他/她为什么会看到 extern、对齐说明符或重新解释转换等。

编辑：这是作为我的 CUDA kernel author's tools 仅标头库的一部分实现的：shared_memory.cuh（它被命名为 shared_memory::dynamic::proxy()）。

【讨论】：