同时使用动态分配和静态分配的共享内存

Question

假设我有两个 __device__ CUDA 函数，每个函数都有以下局部变量：

__shared__ int a[123];

和另一个函数（假设它是我的内核，即 __global__ 函数），具有：

extern __shared__ int b[];

这是 nVIDIA 明确 allowed/forbidden 吗？（我在 __shared__ 的 programming guide B.2.3 节中没有看到它）是否将所有大小一起计入共享内存限制，或者它是一次可能使用的最大值？或者其他规则？

_{这可以被视为 this one 的后续问题。}

Answer 1

共享内存分为两部分：静态分配和动态分配。第一部分是在编译期间计算的，每个声明都是实际分配 - 在编译期间激活 ptxas 信息说明如下：

  ptxas info    : Used 22 registers, 384 bytes smem, 48 bytes cmem[0]

在这里，我们有 384 个字节，这是 3 个 32 个整数的数组。（请参阅下面的示例代码）。

自 Kepler 以来，您可以将指向共享内存的指针传递给另一个允许设备子函数访问另一个共享内存声明的函数。

然后是动态分配的共享内存，其保留大小在内核调用期间声明。

这是几个函数中的一些不同用途的示例。注意每个共享内存区域的指针值。

__device__ void dev1()
{
    __shared__ int a[32] ;
    a[threadIdx.x] = threadIdx.x ;

    if (threadIdx.x == 0)
        printf ("dev1 : %x\n", a) ;
}

__device__ void dev2()
{
    __shared__ int a[32] ;
    a[threadIdx.x] = threadIdx.x * 5 ;

    if (threadIdx.x == 0)
        printf ("dev2 : %x\n", a) ;
}

__global__ void kernel(int* res, int* res2)
{
    __shared__ int a[32] ;
    extern __shared__ int b[];

    a[threadIdx.x] = 0 ;
    b[threadIdx.x] = threadIdx.x * 3 ;

    dev1();
    __syncthreads();
    dev2();
    __syncthreads();

    res[threadIdx.x] = a[threadIdx.x] ;
    res2[threadIdx.x] = b[threadIdx.x] ;

    if (threadIdx.x == 0)
        printf ("global a : %x\n", a) ;
    if (threadIdx.x == 0)
        printf ("global b : %x\n", b) ;
}

int main()
{
    int* dres  ;
    int* dres2 ;

    cudaMalloc <> (&dres, 32*sizeof(int)) ;
    cudaMalloc <> (&dres2, 32*sizeof(int)) ;

    kernel<<<1,32,32*sizeof(float)>>> (dres, dres2);

    int hres[32] ;
    int hres2[32] ;

    cudaMemcpy (hres, dres, 32 * sizeof(int), cudaMemcpyDeviceToHost) ;
    cudaMemcpy (hres2, dres2, 32 * sizeof(int), cudaMemcpyDeviceToHost) ;

    for (int k = 0 ; k < 32 ; ++k)
    {
        printf ("%d -- %d \n", hres[k], hres2[k]) ;
    }
    return 0 ;
}

此代码使用 384 bytes smem 输出 ptxas 信息，即一个数组用于全局 a 数组，第二个用于 dev1 方法 a 数组，第三个用于 dev2 方法 a数组。总计 3*32*sizeof(float)=384 bytes.

当运行具有动态共享内存的内核等于32*sizeof(float)时，指向b的指针就在这三个数组之后开始。

编辑： 此代码生成的 ptx 文件包含静态定义的共享内存的声明，

.shared .align 4 .b8 _ZZ4dev1vE1a[128];
.shared .align 4 .b8 _ZZ4dev2vE1a[128];
.extern .shared .align 4 .b8 b[];

方法主体中定义的入口点除外

// _ZZ6kernelPiS_E1a has been demoted

内存的共享space在PTX文档中定义here:

The shared (.shared) state space is a per-CTA region of memory for threads in a CTA to share data. An address in shared memory can be read and written by any thread in a CTA. Use ld.shared and st.shared to access shared variables.

虽然没有关于运行时的详细信息。编程指南here中有一个词没有进一步详细说明两者的混合。

在 PTX 编译期间，编译器可能知道静态分配的共享内存量。可能会有一些补充魔法。查看 SASS，第一条指令使用 SR_LMEMHIOFF

1             IADD32I R1, R1, -0x8;
2             S2R R0, SR_LMEMHIOFF;
3             ISETP.GE.U32.AND P0, PT, R1, R0, PT;

并以相反的顺序调用函数为静态分配的共享内存分配不同的值（看起来非常像一种 stackalloc 形式）。

我相信 ptxas 编译器会计算在可能调用所有方法的最坏情况下可能需要的所有共享内存（当不使用其中一种方法并使用函数指针时，b 地址不会更改，并且永远不会访问未分配的共享内存区域）。

最后，正如 einpoklum 在评论中建议的那样，这是实验性的，而不是 norm/API 定义的一部分。

同时使用动态分配和静态分配的共享内存

Using both dynamically-allocated and statically-allocated shared memory

cuda

gpu-shared-memory