设备内部的共享内存声明

Question

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

我们可以这样做吗：

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

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

sizeof(a)/sizeof(float) = 768<br> sizeof(b)/sizeof(浮动) = 1024

或者我应该只初始化一个长的一维共享数组并附加 "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

我想后者就是你要找的。

设备内部的共享内存声明

Shared memory declaration inside device

c

cuda

image

histogram