通过结构和重载运算符包装 CUDA 共享内存定义和访问

Question

在代码 here 中，我遇到了一个用于共享内存定义和用法的结构。我将分配修改为静态，并在如下测试程序中使用它：

#include <stdio.h>

template<class T, uint bDim>
struct SharedMemory
{
     __device__ inline operator T *() {
        __shared__ T __smem[ bDim ];
        return (T*) (void *) __smem;
    }
     __device__ inline operator const T *() const {
        __shared__ T __smem[ bDim ];
        return (T*) (void *) __smem;
    }
};

template <uint bDim>
__global__ void myKernel() {
    SharedMemory<uint, bDim> myShared;
    myShared[ threadIdx.x ] = threadIdx.x;
    __syncthreads();
    printf("%d\tsees\t%d\tat two on the circular right.\n", threadIdx.x,     myShared[ ( threadIdx.x + 2 ) & 31 ]);
}

int main() {
    myKernel<32><<<1, 32>>>();
    cudaDeviceSynchronize();
    return 0;
}

正如预期的那样工作正常。但是，我对这种用法有几个问题：

我不明白 sharedMemory 结构中运算符重载部分使用的语法。它是否重载了取消引用运算符 *？如果是，如何通过方括号访问转化为解引用指针？另外，为什么将 __device__ inline operator T *() { 行更改为 __device__ inline T operator *() { 会产生编译器错误？

我想通过重载赋值运算符或定义成员函数来简化包装器的使用，以便每个线程更新与其线程索引对应的共享内存位置。因此，例如，写下 myShared = 47; 或 myShared.set( 47 ); 在幕后转化为 myShared[threadIdx.x] = 47;。但是我这样做没有成功。它编译得很好，但共享内存缓冲区被读取所有 0（我认为这是调试模式下的默认共享内存初始化）。你能告诉我我哪里做错了吗？这是我的尝试：

template<class T, uint bDim>
struct SharedMemory
{
     __device__ inline operator T*() {
        __shared__ T __smem[ bDim ];
        return (T*) (void *) __smem;
    }
     __device__ inline operator const T *() const {
        __shared__ T __smem[ bDim ];
        return (T*) (void *) __smem;
    }
    __device__ inline T& operator=( const T& __in ) {
        __shared__ T __smem[ bDim ];
        __smem[ threadIdx.x ] = __in;
        return (T&) __smem[ threadIdx.x ];
    }
    __device__ inline void set( const T __in ) {
        __shared__ T __smem[ bDim ];
        __smem[ threadIdx.x ] = __in;
    }

};

对于成员函数，编译器给出警告：

variable "__smem" was set but never used

虽然我知道 member variables cannot be __shared__，但我想我有一个错误的假设，或者我想做的事情与 __shared__ 限定符特征不匹配。感谢您的帮助。

Answer 1

看来您对 __shared__ 访问说明符在 CUDA 中的实际作用有一些误解，再加上一个相当棘手的模板，该模板旨在在 extern __shared__ 内存的情况下欺骗编译器在模板化的内核实例中使用，让你走上了一条盲路。

如果我没有正确理解您的需求，那么您真正想要的是这样的：

template<typename T>
struct wrapper
{
    T * p;
    unsigned int tid;

    __device__ wrapper(T * _p, unsigned int _tid) : p(_p), tid(_tid) {}
    __device__ const T* operator->() const { return p + tid; }
    __device__ T& operator*() { return *(p + tid); }
    __device__ const T& operator*() const { return *(p + tid); }
};

这是一个包装器，您可以使用它来 "hide" 一个指针和一个偏移量以 "indexing" 自由访问该指针，例如：

#include <cstdio>

// structure definition goes here

void __global__ kernel(float *in)
{
    __shared__ float _buff[32];
    wrapper<float> buff(&_buff[0], threadIdx.x);

    *buff = in[threadIdx.x + blockIdx.x * blockDim.x];
    __syncthreads();

    for(int i=0; (i<32) && (threadIdx.x == 0); ++i) { 
        printf("%d %d %f\n", blockIdx.x, i, _buff[i]);
    }
}

int main()
{
    float * d = new float[128];
    for(int i=0; i<128; i++) { d[i] = 1.5f + float(i); }

    float * _d;
    cudaMalloc((void **)&_d, sizeof(float) * size_t(128));
    cudaMemcpy(_d, d, sizeof(float) * size_t(128), cudaMemcpyHostToDevice);

    kernel<<<4, 32>>>(_d);
    cudaDeviceSynchronize();
    cudaDeviceReset();

    return 0;
}

在示例内核中，共享内存数组 _buff 在包装器实例中用线程索引包装，运算符重载让您无需通常的显式索引操作即可访问特定的数组元素。也许您可以修改它以满足您的需要。

通过结构和重载运算符包装 CUDA 共享内存定义和访问

Wrapping CUDA shared memory definition and accesses by a struct and overloading operators

c++

struct

cuda

gpu

class