[Question]: extern keyword about CUDA dynamic shared memory

Accelerated Computing CUDA CUDA Programming and Performance
1

nccl/src/device/common.cu

 __shared__ ncclShmemData ncclShmem; 
 #if __CUDA_ARCH__ < 700 
   __shared__ ulong2 ncclShmemPerWarp[ncclShmemScratchWarpSize()*(NCCL_MAX_NTHREADS/WARP_SIZE)/sizeof(ulong2)]; 
 #endif

nccl/src/device/common.h

 extern __shared__ ncclShmemData ncclShmem; 
 #if __CUDA_ARCH__ >= 700 
   extern __shared__ ulong2 ncclShmemPerWarp[/*ncclShmemDynamicSize()/sizeof(ulong2)*/]; 
 #else 
   extern __shared__ ulong2 ncclShmemPerWarp[ncclShmemScratchWarpSize()*(NCCL_MAX_NTHREADS/WARP_SIZE)/sizeof(ulong2)]; 
 #endif

The extern keyword here, does it serve to declare a CUDA dynamic shared memory 「a extern _shared_ variable」, or does it tell the compiler “this variable is defined elsewhere”「a _shared_ variable from extern」? How should I distinguish between the two?

In this context, ncclShmem is both dynamic and static. Will there be a conflict?

2

What is know as dynamic shared memory is declared without an array size, i.e.
extern __shared__ ulong2 ncclShmemPerWarp[];

extern __shared__ ulong2 ncclShmemPerWarp[ncclShmemScratchWarpSize()*(NCCL_MAX_NTHREADS/WARP_SIZE)/sizeof(ulong2)]; declares a fixed size array (in shared memory) which is defined elsewhere.

NCCL seems to use a fixed size static shared memory for GPU archs older than Volta, and uses dynamic shared memory since Volta.

Below code shows different declarations of shared memory. Compile with nvcc main.cu smem.cu -rdc=true -o main

File main.cu

#include <iostream>

__device__
void printsmemsize(const char* name){
    unsigned int total_smem_size, dynamic_smem_size;

    asm(
        "mov.u32  %0, %total_smem_size;\n\t"
        "mov.u32  %1, %dynamic_smem_size; \n\t"
        :"=r"(total_smem_size),"=r"(dynamic_smem_size)
        :
    );
    printf("name %s, total_smem_size %u, dynamic_smem_size %u\n", name, total_smem_size, dynamic_smem_size);
}

//static shared memory
__global__
void kernel1(){
    __shared__ int smem1[1024];
    if(threadIdx.x == 0) printsmemsize("kernel1");
    smem1[threadIdx.x] = threadIdx.x;
    __syncthreads();
    if(threadIdx.x == 0) printf("%d\n", smem1[1023]);
}

//dynamic shared memory
__global__
void kernel2(){
    extern __shared__ int smem2[];
    if(threadIdx.x == 0) printsmemsize("kernel2");
    smem2[threadIdx.x] = threadIdx.x;
    __syncthreads();
    if(threadIdx.x == 0) printf("%d\n", smem2[1023]);
}

//static shared memory outside of kernel
__shared__ int smem3[1024];

__global__
void kernel3(){
    if(threadIdx.x == 0) printsmemsize("kernel3");
    smem3[threadIdx.x] = threadIdx.x;
    __syncthreads();
    if(threadIdx.x == 0) printf("%d\n", smem3[1023]);
}

//static shared memory outside of kernel defined in other file
extern __shared__ int smem4[1024];

__global__
void kernel4(){
    if(threadIdx.x == 0) printsmemsize("kernel4");
    smem4[threadIdx.x] = threadIdx.x;
    __syncthreads();
    if(threadIdx.x == 0) printf("%d\n", smem4[1023]);
}

//dynamic shared memory outside of kernel
extern __shared__ int smem5[];

__global__
void kernel5(){
    if(threadIdx.x == 0) printsmemsize("kernel5");
    smem5[threadIdx.x] = threadIdx.x;
    __syncthreads();
    if(threadIdx.x == 0) printf("%d\n", smem5[1023]);
}



__global__
void kernel6(){
    //static shared memory defined in other file
    extern __shared__ int smem6[1024];
    if(threadIdx.x == 0) printsmemsize("kernel6");
    smem6[threadIdx.x] = threadIdx.x;
    __syncthreads();
    if(threadIdx.x == 0) printf("%d\n", smem6[1023]);
}

int main(){
    kernel1<<<1,1>>>();
    cudaDeviceSynchronize();

    kernel2<<<1,1, 4096>>>();
    cudaDeviceSynchronize();

    kernel3<<<1,1>>>();
    cudaDeviceSynchronize();

    kernel4<<<1,1>>>();
    cudaDeviceSynchronize();

    kernel5<<<1,1,4096>>>();
    cudaDeviceSynchronize();

    kernel6<<<1,1>>>();
    cudaDeviceSynchronize();
}

File smem.cu

__shared__ int smem4[1024];
__shared__ int smem6[1024];
3

Thanks,@striker159
In NCCL, if the architecture is newer than Volta, ncclShmemPerWarp is dynamic shared memory, and ncclShmem is static shared memory. When passing the smemsize parameter to the kernel, does the smemsize specify only the dynamic shared memory, or does it include both dynamic and static shared memory?

4

The parameter passed via kernel launch for shared memory size indicates the size of dynamically allocated shared memory only.