使用了共享内存最终时间没有没有减少

  • 5 replies
  • 273 views
使用了共享内存最终时间没有没有减少
« 于: 十月 19, 2022, 03:23:15 pm »
程序中 kernel 1得到u值,kernel 2和kernel 3都要使用这个u值,因此我申请了全局可见的共享内存用来存储u,以减少2、3步中读取u的延时,可是为什么使用共享内存优化之后runtime没减少呢?
如下图,可以明显看出访问全局内存的次数明显变少,但时间却从1.96ms变成了1.97ms【这是使用nsight compute记录的核函数执行时间】

Re: 使用了共享内存最终时间没有没有减少
« 回复 #1 于: 十月 31, 2022, 04:42:51 pm »
程序中 kernel 1得到u值,kernel 2和kernel 3都要使用这个u值,因此我申请了全局可见的共享内存用来存储u,以减少2、3步中读取u的延时,可是为什么使用共享内存优化之后runtime没减少呢?
如下图,可以明显看出访问全局内存的次数明显变少,但时间却从1.96ms变成了1.97ms【这是使用nsight compute记录的核函数执行时间】

Shared Memory的使用,不能跨kernel生效。你在kernel 1中计算U,和在Kernel2、3中又使用到U,这种不能通过Shared memory进行缓冲(因为shared memory中的数据有效生存期是1次kernel启动中的1个block,不能跨block,更不能跨多次kernel启动)。

这种情况依赖于GPU的L1和L2来缓冲即可。正常情况下,他们可以跨多个kernel启动有效。特殊情况下(例如6.1上的某些情况中),至少L2会生效。

Re: 使用了共享内存最终时间没有没有减少
« 回复 #2 于: 十一月 01, 2022, 10:39:51 am »
Shared Memory的使用,不能跨kernel生效。你在kernel 1中计算U,和在Kernel2、3中又使用到U,这种不能通过Shared memory进行缓冲(因为shared memory中的数据有效生存期是1次kernel启动中的1个block,不能跨block,更不能跨多次kernel启动)。

这种情况依赖于GPU的L1和L2来缓冲即可。正常情况下,他们可以跨多个kernel启动有效。特殊情况下(例如6.1上的某些情况中),至少L2会生效。
感谢回复!
但是代码中我是把共享内存U声明为了全局可见, kernel 1中用计算结果初始化共享数组U,kernel2和kernel3 从共享内存U中取数据来完成自己的计算,结果表示确实正确得到并使用了kernel 1的计算结果。相较于之前数组U是存放在global memory,现在存到了shared memory,这难道不能起到加速作用吗?【另外这是不是也说明了全局可见的共享内存可以跨kernel生效呀?】

Re: 使用了共享内存最终时间没有没有减少
« 回复 #3 于: 十一月 01, 2022, 01:34:45 pm »
感谢回复!
但是代码中我是把共享内存U声明为了全局可见, kernel 1中用计算结果初始化共享数组U,kernel2和kernel3 从共享内存U中取数据来完成自己的计算,结果表示确实正确得到并使用了kernel 1的计算结果。相较于之前数组U是存放在global memory,现在存到了shared memory,这难道不能起到加速作用吗?【另外这是不是也说明了全局可见的共享内存可以跨kernel生效呀?】

并不存在真正"全局可见"的shared memory的,例如以下代码,会给你造成幻觉:
__shared__ int fake_global_visible_shared_array[100];
__global__ void your_kernel()
{
   ....//use fake_global_visible_shared_array[] here
}

这种依然等效于在your_kernel的内部,定义一个shared memory的数组的。虽然它看起来像是"全局的shared memory", 然而它不是。

所以并不能跨kernel多次使用(例如你的一个kernel写入,另外一个kernel读取的)。

其次,因为每次kernel启动前, shared memory中的数据并不清零,你存在概率会读到之前的某kernel的某个block的残留内容的,但这种并不可靠(你也可以读不到,或者读到错误的内容)。

所以综合这两点,这种写法不可取。

Re: 使用了共享内存最终时间没有没有减少
« 回复 #4 于: 十一月 01, 2022, 09:15:40 pm »
并不存在真正"全局可见"的shared memory的,例如以下代码,会给你造成幻觉:
__shared__ int fake_global_visible_shared_array[100];
__global__ void your_kernel()
{
   ....//use fake_global_visible_shared_array[] here
}

这种依然等效于在your_kernel的内部,定义一个shared memory的数组的。虽然它看起来像是"全局的shared memory", 然而它不是。

所以并不能跨kernel多次使用(例如你的一个kernel写入,另外一个kernel读取的)。

其次,因为每次kernel启动前, shared memory中的数据并不清零,你存在概率会读到之前的某kernel的某个block的残留内容的,但这种并不可靠(你也可以读不到,或者读到错误的内容)。

所以综合这两点,这种写法不可取。
那意思就是如果想要借助共享内存来减少存取延时,必须要在一个核函数范围内了吗? 
假设想要把这3个kernel合并在一个kernel中,因为寄存器的原因导致launch失败【寄存器数量超出硬件限制】,即不能合并成一个kernel,那这种情况下只能把 U数组 存在global memory了吗?

Re: 使用了共享内存最终时间没有没有减少
« 回复 #5 于: 十一月 02, 2022, 12:23:24 pm »
那意思就是如果想要借助共享内存来减少存取延时,必须要在一个核函数范围内了吗? 
假设想要把这3个kernel合并在一个kernel中,因为寄存器的原因导致launch失败【寄存器数量超出硬件限制】,即不能合并成一个kernel,那这种情况下只能把 U数组 存在global memory了吗?

(1)首先说,使用global memory并不一定会慢的(或者说使用shared memory一定会快是错觉)。
(2)其次是,如果你非要这样做的话也可以。同时计算能力5.0以上的硬件(最近8年)并不存在"寄存器超出硬件数量限制的问题", 因为编译器最多会为5.0+的卡生成使用最多255个常规寄存器的代码,而5.0+的硬件本身最多一个线程也可以分配255个寄存器(但此时可能最大的block形状受到限制,例如256个线程的block)。当然,你如果还在使用3.X的老卡,那就没办法了. :(