如何修复警告：函数范围的静态 __shared__ 变量不支持动态初始化？

Posted 2023-03-23

【中文标题】如何修复警告：函数范围的静态 __shared__ 变量不支持动态初始化？

【英文标题】：How to fix warning : dynamic initialization is not supported for a function-scope static __shared__ variable?

【发布时间】：2021-12-26 00:12:14

【问题描述】：

我有一个结构简单的构造函数，比如

struct TResult

    int field1, field2;

    __host__ __device__ TResult()
        : field1(0),
          field2(0)
     
;

并收集类似的代码

#define BLOCK_SIZE 128

__global__ void uniteResults(TResult *destResults, TResult *srcResults)

    __shared__ TResult sums[BLOCK_SIZE];

    sums[threadIdx.x] = TResult();
    //...

据我了解，我正确实现了共享数组的初始化。 但我得到了

warning : dynamic initialization is not supported for a function-scope static __shared__ variable within a __device__/__global__ function

对于 __shared__ 行。好的，这意味着应该调用构造函数，但是 Nvidia 的编译器不能这样做。如何正确实施，让每个人都开心？

【问题讨论】：

显而易见的问题是哪个线程应该在定义时运行构造函数？如果他们都这样做，那么你就有了记忆竞赛。答案是有一个默认的可构造类型。如果您需要初始化，请在每个块调用中定义一个线程或执行它。

您还可以切换到动态分配的共享内存（或者甚至静态分配的基本类型，如char 和reinterpret_cast）然后从单个线程执行放置new 以初始化您的对象，而无需更改任何课程代码。

你能写代码吗？ __shared__ TResult sums[];?

【参考方案1】：

正如@talonmies 建议的那样，您需要帮助 CUDA 编译器避免“哪个线程初始化数组值？”的困境。 .即使是像你这样的微不足道的构造函数也不足以让这种情况发生。只有 default 构造函数 - 当它实际上意味着非构造时 - 有效。所以：

struct TResult

    int field1, field2;

    TResult() = default;
;

这基本上就像在说：

struct TResult

    int field1, field2;
;

你将不得不动态初始化。不过幸运的是，您的动态初始化代码是 fine（只要块尺寸确实是 BLOCK_SIZE x 1 x 1）。

【参考方案2】：

以下是使用放置new 的一种可能解决方法的示例：

#define BLOCK_SIZE 128
#include <new>
struct TResult

    int field1, field2;

    __host__ __device__ TResult()
        : field1(0),
          field2(0)
     
;
__global__ void uniteResults(TResult *destResults, TResult *srcResults)

    extern __shared__ TResult sums[];
    new(sums+threadIdx.x) TResult();
    __syncthreads();

上述代码中的假设是您将使用适当的动态分配的共享内存大小声明调用内核：

uniteResults<<<blocks, BLOCK_SIZE, BLOCK_SIZE*sizeof(TResult)>>>(...);

这个想法有很多变化是可能的。