c - 为什么我的内核的共享内存似乎被初始化为零？

正如在这个 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模式 进行了测试:“-O3 -arch=sm_20”、“-O3 -arch=sm_30”和“-arch=sm_30”。结果总是一样的。

最佳答案

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;
}

关于c - 为什么我的内核的共享内存似乎被初始化为零？，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/22172881/