在 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
虽然我知道成员变量不能是
__shared__
,但我想我有一个错误的假设或者我想做的与__shared__
限定符特征不匹配。感谢您的帮助。
看来您对 CUDA 中
__shared__
访问说明符的实际作用有一些误解,再加上一个相当棘手的模板,该模板旨在在模板化内核实例中使用 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); }
};
这是一个包装器,您可以使用它来“隐藏”指针和偏移量,以便“索引”自由访问指针,例如:
#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
用线程索引包装在包装器实例中,运算符重载让您无需通常的显式索引操作即可访问特定的数组元素。也许你可以修改它以满足你的需要。