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

程序中 kernel 1得到u值，kernel 2和kernel 3都要使用这个u值，因此我申请了全局可见的共享内存用来存储u，以减少2、3步中读取u的延时，可是为什么使用共享内存优化之后runtime没减少呢？
如下图，可以明显看出访问全局内存的次数明显变少，但时间却从1.96ms变成了1.97ms【这是使用nsight compute记录的核函数执行时间】

程序中 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会生效。

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生效呀？】

感谢回复！
但是代码中我是把共享内存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的残留内容的，但这种并不可靠（你也可以读不到，或者读到错误的内容）。

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

并不存在真正"全局可见"的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了吗？

那意思就是如果想要借助共享内存来减少存取延时，必须要在一个核函数范围内了吗？ 
假设想要把这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的老卡，那就没办法了.