在代码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;
}
它像预测的那样工作正常。但是,我对这种用法有几个问题:
1.我不理解sharedMemory
结构中运算符重载部分使用的语法。是否重载解引用运算符*
?如果是,如何通过方括号访问转换为解引用指针?另外,为什么将__device__ inline operator T *() {
行更改为__device__ inline T operator *() {
会产生编译器错误?
1.我想通过重载赋值操作符或定义成员函数来简化 Package 器的使用,以便每个线程更新与其线程索引相对应的共享内存位置。因此,例如,写下myShared = 47;
或myShared.set( 47 );
在窗帘后面转换为myShared[threadIdx.x] = 47;
。但我一直没有成功地做到这一点。它编译得很好,但是共享内存缓冲区被读取了所有0
(我认为这是Debug模式下的默认共享内存初始化)。你能告诉我哪里做错了吗?以下是我的尝试:
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__
限定符特征不匹配。谢谢你的帮助
1条答案
按热度按时间nhaq1z211#
看起来你对
__shared__
访问说明符在CUDA中的实际作用有一些误解,再加上一个相当棘手的模板,旨在欺骗编译器,因为在模板化的内核示例中使用了extern __shared__
内存,这让你走上了一条盲目的道路。如果我对你的需求理解正确的话,你真正想要的是这样的:
这是一个 Package 器,您可以使用它来“隐藏”指针和偏移量,以便“索引”自由访问指针,例如:
在示例内核中,共享内存数组
_buff
被 Package 在 Package 器示例中的线程索引中,操作符重载允许您访问特定的数组元素,而无需通常的显式索引操作。也许你可以修改它以满足你的需要。