Shared Memory in CUDA

Overview

  • Time: 30 min

  1. Learn about shared memory in CUDA.

  2. Understand the characteristics and usage of shared memory.

Shared memory is a type of memory that is shared among threads within the same block in CUDA. It is much faster than global memory and can be used to store data that needs to be accessed by multiple threads. Shared memory is particularly useful for reducing global memory accesses and improving performance in parallel algorithms.

Shared Memory Characteristics

Shared memory in CUDA has several key characteristics:

  1. Scope: Shared memory is accessible only to threads within the same block. Threads in different blocks cannot access each other’s shared memory.

  2. Speed: Shared memory is on-chip memory, making it much faster than global memory (DRAM). Accessing shared memory typically has a latency of around 10-50 cycles, compared to 400-800 cycles for global memory.

  3. Size: The size of shared memory is limited, typically ranging from 48 KB to 96 KB per block, depending on the GPU architecture.

  4. Coherency: Shared memory is coherent across all threads in a block, meaning that changes made by one thread are visible to others in the same block.

  5. Usage: Shared memory is often used for data that needs to be accessed frequently by multiple threads, such as intermediate results in parallel computations.

There are two main types of shared memory in CUDA:

  1. Static Shared Memory: This is defined at compile time and has a fixed size. It is declared using the __shared__ keyword in the kernel code. Static shared memory is allocated for each block and is initialized when the block is launched.

 1 __global__ void addKernel(int *input, int *output)
 2{
 3    __shared__ int sdata[256];  // dynamic shared memory
 4
 5    int tid = threadIdx.x;
 6    int idx = blockIdx.x * blockDim.x + threadIdx.x;
 7
 8    // Load data from global to shared memory
 9    sdata[tid] = input[idx];
10    __syncthreads();
11
12    // Do something with shared memory
13    sdata[tid] += 10;
14    __syncthreads();
15
16    // Store result back to global memory
17    output[idx] = sdata[tid];
18}

  1. Dynamic Shared Memory: This is allocated at runtime and can vary in size based on the needs of the kernel. It is declared using the extern __shared__ keyword in the kernel code. Dynamic shared memory allows for more flexibility, as its size can be specified when launching the kernel.

 1__global__ void addKernel(int *input, int *output)
 2{
 3    extern __shared__ int sdata[];  // dynamic shared memory
 4    int tid = threadIdx.x;
 5    int idx = blockIdx.x * blockDim.x + threadIdx.x;
 6
 7    // Load data from global to shared memory
 8    sdata[tid] = input[idx];
 9    __syncthreads();
10
11    // Do something with shared memory
12    sdata[tid] += 10;
13    __syncthreads();
14
15    // Store result back to global memory
16    output[idx] = sdata[tid];
17}
18
19int main()
20{
21    int threads = 256;
22    int sharedMemSize = threads * sizeof(int);  // Size of dynamic shared memory
23
24    ...
25    ...
26    ...
27
28    addKernel<<<1, threads, sharedMemSize>>>(d_input, d_output);
29}

Important

You declare only one extern __shared__ array inside your kernel.

Explanation

__syncthreads() is a barrier synchronization function in CUDA that ensures all threads in the same thread block:

  1. Reach the barrier, and

  2. Complete all memory accesses (reads/writes to shared memory)

before any thread in the block continues past it.

Key Points

  1. Shared memory is a fast, on-chip memory accessible by threads within the same block.

  2. It is used to reduce global memory accesses and improve performance in parallel algorithms.

  3. There are two types of shared memory: static (fixed size) and dynamic (runtime size).

  4. The __syncthreads() function is used to synchronize threads within a block, ensuring all threads reach a certain point before proceeding.