为啥我的内核的共享内存似乎初始化为零？

Posted 2023-04-15

【中文标题】为啥我的内核的共享内存似乎初始化为零？

【英文标题】：Why does my kernel's shared memory seems to be initialized to zero?为什么我的内核的共享内存似乎初始化为零？

【发布时间】：2014-04-06 01:01:41

【问题描述】：

正如Shared Memory Array Default Value 问题中提到的，共享内存是未初始化的，即可以包含任何值。

#include <stdio.h>

#define BLOCK_SIZE 512

__global__ void scan(float *input, float *output, int len) 
    __shared__ int data[BLOCK_SIZE];

    // DEBUG
    if (threadIdx.x == 0 && blockIdx.x == 0)
    
        printf("Block Number: %d\n", blockIdx.x);
        for (int i = 0; i < BLOCK_SIZE; ++i)
        
            printf("DATA[%d] = %d\n", i, data[i]);
        
    



int main(int argc, char ** argv) 
    dim3 block(BLOCK_SIZE, 1, 1);
    dim3 grid(10, 1, 1);
    scan<<<grid,block>>>(NULL, NULL, NULL);
    cudaDeviceSynchronize();
    return 0;

但是为什么在这段代码中它是不正确的并且我经常将共享内存归零？

DATA[0] = 0
DATA[1] = 0
DATA[2] = 0
DATA[3] = 0
DATA[4] = 0
DATA[5] = 0
DATA[6] = 0
...

我使用 Release 和 Debug Mode 进行了测试：“-O3 -arch=sm_20”、“-O3 -arch=sm_30”和“-arch=sm_30”。结果总是一样的。

【问题讨论】：

你在release和debug模式下测试了吗？在某些项目中，我观察到在调试模式下共享内存被初始化为 0，但不是在发布模式下，也不是在所有项目中都通用。这不是@CygnusX1 在您的链接问题中回答的定义行为。您必须自己初始化共享内存！

如果它可以包含任何值而不是可以包含零，不是吗？系统有时可能仍需要重新初始化内存以防止进程之间的信息泄漏（安全性）。

是的，我测试过。使用“-arch=sm_30”和“-O3 -arch=sm_30”选项，也可以使用“-arch=sm_20”。结果是相同的 - 共享内存归零。

是的，它也可以包含和归零，但奇怪的是共享内存中没有任何其他值，意味着它是专门归零的。

零在“任意值”的子集中。

【参考方案1】：

tl;dr: 共享内存未初始化为 0

我认为您对初始化为0 的共享内存的猜想是有问题的。试试下面的代码，这是对你的稍作修改。在这里，我调用内核两次并更改data 数组的值。第一次启动内核时，data 的“未初始化”值将全部为0。第二次启动内核时，data 的“未初始化”值将与0 的完全不同。

我认为这取决于共享内存是 SRAM 的事实，它展示了data remanence。

#include <stdio.h>

#define BLOCK_SIZE 32

__global__ void scan(float *input, float *output, int len) 

    __shared__ int data[BLOCK_SIZE];

    if (threadIdx.x == 0 && blockIdx.x == 0)
    
        for (int i = 0; i < BLOCK_SIZE; ++i)
        
            printf("DATA[%d] = %d\n", i, data[i]);
            data[i] = i;
        

    


int main(int argc, char ** argv) 
    dim3 block(BLOCK_SIZE, 1, 1);
    dim3 grid(10, 1, 1);
    scan<<<grid,block>>>(NULL, NULL, NULL);
    scan<<<grid,block>>>(NULL, NULL, NULL);
    cudaDeviceSynchronize();
    getchar();
    return 0;

【讨论】：

是的，你完全正确！感谢这个很好的解释，现在很清楚为什么我们应该手动初始化共享内存！