CUDA中是否有官方认可的__global__函数间复用__shared__数据的方法？

是否存在官方认可的在CUDA的__global__函数之间复用__shared__数据的方法？参考以下示例代码：

#include <cuda.h>
#include <stdio.h>

__global__ void kernel_A(int** s) {
  __shared__ int sdata;
  sdata = clock();
  *s = &sdata;
  printf("sdata: %p = %i\n", &sdata, sdata);
}

//somewhat surprisingly, inserting a device function with __shared__ data
//does not move kernel_B's shared data. 
//Maybe this is just luck?
__device__ int& make_shared() {
  __shared__ int bdata;
  bdata = 1;
  auto& result = bdata;
  return result;
}

__global__ void kernel_B(int** ptr_sdata) {
  auto& bdata = make_shared();
  printf("bdata: %p = %i\n", &bdata, bdata);
  __shared__ int sdata;
  printf("A.sdata: %p = %i, B.sdata: %p = %i\n", *ptr_sdata, *ptr_sdata[0], &sdata, sdata);
}

int main() {
  int** ptr_sdata;
  cudaMalloc(&ptr_sdata, sizeof(int*));
  kernel_A<<<1,1>>>(ptr_sdata);
  kernel_B<<<1,1>>>(ptr_sdata);
  cudaDeviceSynchronize();
}

运行输出：

sdata: 0x7d52c5000000 = -594894858
bdata: 0x7d52c5000004 = 1
A.sdata: 0x7d52c5000000 = -594894858, B.sdata: 0x7d52c5000000 = -594894858

已知kernel_B可能不会分配到与kernel_A相同的SM，但若假设每个SM仅运行一个block，请问能否可靠复用kernel_A的__shared__内存内容？

结论：没有官方认可的跨全局函数复用__shared__内存的方法，即使假设单SM单Block也不可靠 #

• __shared__内存的生命周期绑定线程块：CUDA规范明确规定，__shared__内存的生命周期与启动它的线程块完全绑定。当线程块执行完毕后，该Block占用的__shared__内存会被CUDA运行时回收或重置，后续任何Block（即使在同一个SM上）都不能依赖该内存的内容保持不变。

• SM资源调度存在不确定性：即便假设每个SM只运行一个Block，CUDA运行时仍有权在Block结束后重新初始化SM上的共享内存区域，或将该区域分配给其他用途。你示例中的输出只是特定硬件、驱动版本下的偶然现象，不属于可依赖的稳定行为。

• 官方不支持跨Kernel共享__shared__内存：CUDA编程模型中，不同__global__函数（Kernel）之间的状态隔离是基本设计原则。跨Kernel共享数据的官方方法只有使用全局内存（Global Memory）或统一内存（Unified Memory），__shared__内存从设计上就不是用来实现跨Kernel共享的。

内容的提问来源于stack exchange，提问作者Johan