设备内的共享内存声明

Question

在CUDA设备内核中允许多少共享内存声明？

我们能做类似这样的事情吗：

extern __shared__ float a[];
extern __shared__ float b[];

我希望有两个不同大小的数组。例如，在1024x768的图像中，我可以通过先在行上最小化，然后在列上最小化来进行并行最小化。因此，为了存储中间值，我将需要

sizeof(a)/sizeof(float) == 768

并且

sizeof(b)/sizeof(float) == 1024

还是我应该初始化一个长的一维共享数组并附加 a 和 b？

- harpribot

Answer 1

您可以拥有任意数量的共享内存声明。然而，运行时只分配一个共享内存缓冲区，每个共享内存数组将被分配相同的地址（即共享内存分配的起始地址）。因此，例如：

#include <cstdio>

extern __shared__ int a[];
extern __shared__ int b[];
extern __shared__ int c[];

__global__
void kernel(void)
{
    int * a0 = &a[0];
    int * b0 = &b[0];
    int * c0 = &c[0];

    printf("a0 = %#x \n", a0);
    printf("b0 = %#x \n", b0);
    printf("c0 = %#x \n", c0);
}

int main()
{
    kernel<<<1,1,1024>>>();
    cudaDeviceReset();

    return 0;
}

这个做什么：

$ nvcc -arch=sm_30 -run extshm.cu 
a0 = 0x1000000 
b0 = 0x1000000 
c0 = 0x1000000 

如果想要有两个共享数组，在任何支持（即计算能力>=2.0）的GPU上，可以像这样做：

#include <cstdio>

extern __shared__ int a[];

__global__
void kernel(void)
{
    int * a0 = &a[0];
    int * b0 = &a[1024];
    int * c0 = &a[1024+768];

    printf("a0 = %#x \n", a0);
    printf("b0 = %#x \n", b0);
    printf("c0 = %#x \n", c0);
}

int main()
{
    kernel<<<1,1,1024+768+512>>>();
    cudaDeviceReset();

    return 0;
}

这将会得到：

nvcc -arch=sm_30 -run extshm2.cu 
a0 = 0x1000000 
b0 = 0x1001000 
c0 = 0x1001c00 

我认为后者是您正在寻找的内容。

- talonmies

---

在编程指南中也给出了一个例子。 - Robert Crovella