想把shared memory配置到64KB,不知道是不是把L1修改成Non-caching模型就可以了,即添加编译选项 -Xptxas -dlm=cg。
求大大们帮忙,谢谢啦!
如果把64KB都设置为shared memory。在kernel中,我定义一个__shared__ float aa[2048],刚好把64kb全部占完,这是否可行?是否有问题。
楼主您好,未能在凌晨2点给您回复,深表歉意。
现在说明您的几个疑问:
(1)L1 cache无法被关闭,以及添加-dlcm=cg只是在访问global memory的时候bypass L1, 直接访问L2.
(2)即使您使用了-Xptxas -dlcm=cg, L1 cache依然存在,并为您的local traffic服务。
(3)关于您的shared memory设置问题,您无法设置它为64KB的。您的可选shared memory配置为:16KB shared + 48 KB L1, 48KB shared + 16KB L1 或者 32KB shared + 32KB L1(这种需要3.x计算能力)。
以及,上文说的最大48KB的shared memory您可以通过:
cudaDeviceSetCacheConfig(cudaFuncCachePreferShared);得到。
以及,因为您实际无法使用64KB的shared memory, 所以您如何定义__shared__数组来占据整个64KB的问题的前置条件不成立,所以无需回答。
感谢您的来访。
以及,需要说明的是,在3.x上,访问global memory总是bypass L1的,无论您的配置。此时L1 data cache只为local traffic服务。
谢谢版主的细心回答
看了版主的回答才知道之前对可配置的64kb理解有错的。
您的每一丝进步都让我们感到欣慰, 感谢您选择CUDA技术。
其实,为什么不可以有0+64KB的设置参数呢?感觉上是设计上的缺陷吧。有时这种极端配置效率最高,为什么不给用户更多自由选择的权利?
您可以将此建议发送给NVIDIA。如果被接受,则您可能在计算能力4.x中看到。
您好,官方文档指出无法配置为0+64KB形式,我亦无法回答此间原因。
如果您觉得有必要支持0+64KB的形式,可以向原厂支持建议。当然威胁老黄去烤肉可能是更快捷的方法XD
祝您编码愉快~