1

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

__shared__ int a[123];

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

extern __shared__ int b[];

这是 nVIDIA 明确允许/禁止的吗？（我在编程指南第 B.2.3 节中没有看到__shared__）所有大小一起计入共享内存限制，还是一次可能使用的最大值？还是其他什么规则？

这可以被认为是这个问题的后续问题。

4

1 回答 1

3

共享内存分为两部分：静态分配和动态分配。第一部分是在编译期间计算的，每个声明都是一个实际的分配 - 在编译期间激活 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 ;
}

此代码使用 输出 ptxas 信息384 bytes smem，即一个数组用于全局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

内存的共享空间在此处的 PTX 文档中定义：

共享 (.shared) 状态空间是每个 CTA 的内存区域，供 CTA 中的线程共享数据。CTA 中的任何线程都可以读取和写入共享内存中的地址。使用 ld.shared 和 st.shared 访问共享变量。

尽管没有关于运行时的详细信息。这里的编程指南中有一个词，没有进一步详细说明两者的混合。

在 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 在评论中所建议的那样，这是实验性的，而不是规范/API 定义的一部分。

于 2016-05-01T11:57:47.553 回答