我想在shared memory中存放一维数组表示的哈希表,这个哈希表可以被多个__global__声明的函数修改,我是用以下方法直接定义的__shared__ hash hashtable[sizeof(hash)*512];
global void lookup(int search_key){
…
int value=hashtable[search_key].value;
…
}
global void insert(hash insert_item){
…
hashtable[slot].key=insert_item.key;
hashtable[slot].value=insert_item.value;
…
}
我使用CUDA10.1,参数是-arch=sm_30,sizeof(hash)是64位,使用512个线程计算
但是使用shared_memory的速度相比于global memory几乎没有提升,
并且如果我把数组的长度增大至1024,会提示我too much shared data,但是数组的大小远远小于shared memory的48KB大小;
如果我将数组定义在__global__声明的函数中就不会报以下错误
ptxas error : Entry function ‘_Z20gpu_hashtable_lookupP8KeyValuej’ uses too much shared data (0x10000 bytes, 0xc000 max)
ptxas error : Entry function ‘_Z20gpu_hashtable_insertPK8KeyValuej’ uses too much shared data (0x10000 bytes, 0xc000 max)
ptxas error : Entry function ‘_Z14init_hashtablev’ uses too much shared data (0x10000 bytes, 0xc000 max)
我看到论坛中有些回复说的是因为在计算能力为1.x的设备中shared memory会传递内核参数,因此会占用部分shared memory,但是我使用的并非1.x的设备。
因此我想请教一下为什么会出现shared memory内存不足的问题,以及为何速度相比于global memory并没有明显增加?