从一个内核启动到另一个内核启动的共享内存持久化

Is shared memory persistent from one kernel launch to another?

本文关键字：启动 内核 另一个 共享 内存 持久化 一个      更新时间：2023-10-16

当试图找到共享内存是否可以被多个内核访问时，我发现有时共享内存中的数据在被另一个内核访问时仍然存在，但有时不是。更重要的是，当使用cuda-gdb调试程序时，上一个内核写入共享内存的数据总是可以被下一个内核读取。

下面是一段测试代码，使用2个gpu。

    extern __shared__ double f_ds[];
    __global__ void kernel_writeToSharedMem(double* f_dev, int spd_x)
    {
       int tid_dev_x = (blockDim.x * blockIdx.x + threadIdx.x);
       int tid_dev_y = (blockDim.y * blockIdx.y + threadIdx.y);
       int tid_dev = tid_dev_y* spd_x + tid_dev_x;
       
       if(tid_dev < blockDim.x * blockDim.y * gridDim.x*gridDim.y)
          f_ds[threadIdx.y*blockDim.x+threadIdx.x] = 0.12345;
       __syncthreads()
    }
  __global__ void kernel_readFromSharedMem(double *f_dev, int dev_no, int spd_x)
    {
       int tid_dev_x = (blockDim.x * blockIdx.x + threadIdx.x);
       int tid_dev_y = (blockDim.y * blockIdx.y + threadIdx.y);
       int tid_dev = tid_dev_y* spd_x + tid_dev_x;
      
       if(tid_dev < blockDim.x * blockDim.y * gridDim.x*gridDim.y)
         {
           f_dev[tid_dev] = f_ds[threadIdx.y*blockDim.x+threadIdx.x];
           printf("threadID %d in dev [%d] is having number %fn",
                   tid_dev,dev_no,f_ds[threadIdx.y*blockDim.x+threadIdx.x]);
         }
       __syncthreads();
     }

    int main()
    {
     ...
       dim3 block_size(BLOCK_SIZE,BLOCK_SIZE);
       im3 grid_size(spd_x/BLOCK_SIZE,spd_y/BLOCK_SIZE);
       for(int i = 0; i < ngpus; i++)
         {
           cudaSetDevice(i);
           kernel_writeToSharedMem<<<grid_size,block_size,sizeof(double)*BLOCK_SIZE*BLOCK_SIZE,stream[i]>>>(f_dev[i],spd_x);
           cudaDeviceSynchronize();
           cudaThreadSynchronize();
          }
        for(int i = 0; i < ngpus; i++)
         {
           cudaSetDevice(i);
           kernel_reaFromSharedMem<<<grid_size,block_size,sizeof(double)*BLOCK_SIZE*BLOCK_SIZE,stream[i]>>>(f_dev[i], int i, spd_x);
           cudaDeviceSynchronize();
           cudaThreadSynchronize();
          }
      ...
    }

运行程序后出现四种情况:

  1. 设备0为0.12345，设备1为0;

  2. 设备0为0，设备1为0.12345;

  3. 设备0和设备1均为0;

  4. 设备0和设备1均为0.12345.

在cuda-gdb中运行时总是如此。

这是否表明共享内存的持久性只有一个内核?共享内存是否只在一个内核偶尔之后才被清除或释放?

保证共享内存的作用域只在分配给它的块的生命周期内存在。任何试图在块到块或内核启动到内核启动之间重用共享内存的行为都是完全未定义的行为，在正常的代码设计中都不应该依赖

。
相关文章：