templates - 封装在模板类中的 CUDA 共享内存，指向相同的内存

Question

我正在尝试在模板类中的 CUDA 内核中分配共享内存:

template<typename T, int Size>
struct SharedArray {

    __device__ T* operator()(){
        __shared__ T x[Size];
        return x;
    }
};

只要没有相同类型和大小的共享内存被检索两次，这就会起作用。但是当我尝试获得相同类型和大小的两次共享内存时，第二个共享内存指向第一个:

__global__
void test() {

    // Shared array
    SharedArray<int, 5> sharedArray;
    int* x0 = sharedArray();
    int* y0 = sharedArray();
    x0[0] = 1;
    y0[0] = 0;
    printf("%i %i\n\n", x0[0], y0[0]);
    // Prints:
    // 0 0
}

一种解决方案是在每次调用共享内存类时添加一个 id，例如:

template<int ID, typename T, int Size>
struct StaticSharedArrayWithID {

    __device__ static T* shared(){
        __shared__ T x[Size];
        return x;
    }   
};

但是我必须提供一些计数器，它提供了一个非常丑陋的用户界面:

__global__
void test() {

    int& x1 = StaticSharedArrayWithID<__COUNTER__, int, 5>::shared();
    int& y1 = StaticSharedArrayWithID<__COUNTER__, int, 5>::shared();
    x1[0] = 1;
    y1[0] = 0;
    printf("%i %i\n\n", x1[0], y1[0]);
    // Prints:
    // 1 0
}

有没有人想摆脱 __COUNTER__用户界面中的宏？隐藏起来就没事了。

Answer 1

这是因为__shared__变量是 static默认情况下。同一个函数的同一个实例引用同一个变量。这种行为的最初原因是因为编译器无法推断何时可以回收内存。有一个变量 static使其与内核一样长。

一个副作用是，如果您在程序中的两个位置对同一个函数调用了两次 - 您会得到相同的结果。事实上，当多个 CUDA 线程在同一位置调用您的函数时，这就是您所期望的，不是吗？

没有一种干净的方法可以动态分配共享内存。在我的项目中，我是通过我自己的共享内存管理器完成的(前面的指针算法很丑，当心!):

typedef unsigned char byte;

/*
  Simple shared memory manager.
    With any luck if invoked with constant parameters this will not take up any register whatsoever
    Must be called uniformly by whole block which is going to use these
    sSize - amount of preallocated memory
*/
template <size_t sSize>
class SharedMemoryManager {
private:
    byte* shArray;
    byte* head;
public:

    __device__ SharedMemoryManager() {
        __shared__ byte arr[sSize];
        shArray=arr;
        head=arr;
    }

    __device__  void reset() {
        head=shArray;
    }

    __device__  byte* getHead() {return head;}
    __device__  void setHead(byte* newHead) {head=newHead;}

    template <typename T>
    __device__  T* alloc(size_t count) {
      size_t addr = head;
      size_t alignment = __alignof(T); //assuming alignment is power of 2
      addr = ((addr-1) | (alignment-1)) +1; //round up to match the alignment requirement
      head = (byte*)(addr);
      T* var = (T*)(head);
      head+=sizeof(T)*size;
      return allocAt<T>(head,count);
    }

template <typename T>
    __device__  T& alloc() {
      return *alloc<T>(1);
    }

};

您可以使用 getHead/ setHead当您知道可以回收共享内存时回收共享内存，但只能以堆栈方式回收。

当 CUDA 不是您的目标时，这种方法应该很容易抽象出非共享内存。

那么你应该能够写:

__global__
 void test() {
   SharedMemoryManager shMem<1024>();

   int& xValue = shMem.alloc<int>();
   int& yValue = shMem.alloc<int>();
   int* xArray = shMem.alloc<int>(5);
   int* yArray = shMem.alloc<int>(5);

   xArray[0] = 1;
   yArray[0] = 0;
   printf("%i %i\n\n", xArray[0], yArray[0]);
   __syncthreads();
   shMem.reset(); //memory reclaimed
   ... 
}