设备内部共享内存声明

Question

在CUDA中的设备内核中允许多少共享内存decalrations？

我们可以做这样的事情:

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

我希望有2个不同大小的数组.例如,在1024*768图像中.我可以通过首先最小化跨行然后跨列来进行并行最小化.所以要存储我需要的中间值

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

或者我应该初始化一个长1D共享arrray并附加"a"和"b" ??

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 

我想,后者正是你要找的.