What is the fastest memory in CUDA?

Registers are the fastest memory in CUDA, followed by Shared Memory. Global Memory is the slowest.

What is __shared__ memory used for?

Shared memory is used for inter-thread communication within a block and as a programmable cache to reduce global memory accesses.

CUDA Memory Management

Chapter 3 — Global, shared, constant, and local memory

The GPU Memory Hierarchy: Latency and Bandwidth

Understanding memory is the single most important factor for writing fast CUDA code. The GPU has multiple memory types, each with different size, speed, and scope:

Registers — fastest, private to each thread (a few KB)
Local memory — per-thread, backed by DRAM (slow), used when registers spill
Shared memory — on-chip scratchpad shared by all threads in a block (~48–164 KB, very fast)
Global memory — large DRAM accessible by all threads (GBs, slowest)
Constant memory — read-only, cached, 64 KB
Texture memory — read-only with a specialized cache for spatial locality

Speed comparison

Registers: ~1 cycle | Shared memory: ~5 cycles | Global memory: ~200–800 cycles. That's a 100× difference! Optimizing memory access is critical.

Global Memory: cudaMalloc and Data Transfers

Global memory is the primary way to move data between host and device. You allocate with cudaMalloc and transfer with cudaMemcpy:

cuda
123456789101112131415161718192021222324
#include <stdio.h>

int main() {
    int N = 1024;
    float *h_data = (float*)malloc(N * sizeof(float));
    float *d_data;
    size_t bytes = N * sizeof(float);

    // Allocate on GPU
    cudaMalloc(&d_data, bytes);

    // Host → Device
    cudaMemcpy(d_data, h_data, bytes, cudaMemcpyHostToDevice);

    // Device → Host
    cudaMemcpy(h_data, d_data, bytes, cudaMemcpyDeviceToHost);

    printf("Successfully allocated and copied %d bytes\n", (int)bytes);

    // Free GPU memory
    cudaFree(d_data);
    free(h_data);
    return 0;
}

Shared Memory: Fast Inter-Thread Communication

Shared memory is an on-chip scratchpad that threads within the same block can use to communicate and collaborate. It's declared with __shared__:

shared_example.cu
123456789101112131415161718192021222324252627282930313233343536373839404142
#include <stdio.h>

__global__ void sharedMemReverse(float *input, float *output, int n) {
    // Dynamically sized shared memory or static
    __shared__ float cache[256];

    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + threadIdx.x;

    // Load from global memory into shared memory
    if (gid < n) {
        cache[tid] = input[gid];
    }

    // Synchronize — ensure all threads have loaded their data
    __syncthreads();

    // Now every thread in the block can read any cache[] element
    // Example: reverse order within block
    if (gid < n) {
        output[gid] = cache[blockDim.x - 1 - tid];
    }
}

int main() {
    const int N = 256;
    float h_in[N], h_out[N];
    for (int i = 0; i < N; i++) h_in[i] = i;

    float *d_in, *d_out;
    cudaMalloc(&d_in, N * sizeof(float));
    cudaMalloc(&d_out, N * sizeof(float));

    cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
    sharedMemReverse<<<1, N>>>(d_in, d_out, N);
    cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);

    printf("h_out[0] (should be 255): %f\n", h_out[0]);
    
    cudaFree(d_in); cudaFree(d_out);
    return 0;
}

Always sync before reading shared memory

If you read shared memory written by other threads without calling__syncthreads() first, you will get undefined (garbage) values.

Constant Memory: Read-Only Caching for Parameters

Use constant memory for small, read-only data that every thread reads (e.g., lookup tables, filter coefficients). It's cached and broadcast-optimized:

cuda
123456789101112131415161718192021222324252627282930313233
#include <stdio.h>

__constant__ float coefficients[8];  // stored in fast constant cache

__global__ void applyFilter(float *data, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        data[i] *= coefficients[i % 8];  // all threads broadcast same coeff
    }
}

int main() {
    // Simple box-filter weights
    float h_coeffs[8] = {0.5f, 0.75f, 1.0f, 0.75f, 0.5f, 0.25f, 0.125f, 0.0625f};
    cudaMemcpyToSymbol(coefficients, h_coeffs, sizeof(h_coeffs));

    const int N = 64;
    float h_data[N], h_out[N];
    for (int i = 0; i < N; i++) h_data[i] = 1.0f;

    float *d_data;
    cudaMalloc(&d_data, N * sizeof(float));
    cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);

    applyFilter<<<1, N>>>(d_data, N);

    cudaMemcpy(h_out, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);
    printf("h_out[0]=%.4f (coeff[0]=0.5), h_out[2]=%.4f (coeff[2]=1.0)\n",
           h_out[0], h_out[2]);

    cudaFree(d_data);
    return 0;
}

Unified Memory (Managed Memory)

CUDA 6+ introduced Unified Memory, which simplifies programming by creating a single address space accessible from both CPU and GPU:

cuda
1234567891011121314
float *data;
cudaMallocManaged(&data, N * sizeof(float));

// Use on CPU
for (int i = 0; i < N; i++) data[i] = i;

// Use on GPU — no explicit copy needed!
myKernel<<<blocks, threads>>>(data, N);
cudaDeviceSynchronize();

// Read results on CPU — automatically migrated back
printf("result: %f\n", data[0]);

cudaFree(data);

When to use Unified Memory

Great for prototyping and simple programs. For maximum performance, explicitcudaMemcpy with pinned memory and streams gives you more control.

Pinned (Page-Locked) Host Memory for Faster Transfers

Regular malloc memory can be paged out by the OS, slowing transfers. Pinned memory stays in physical RAM and enables faster DMA transfers:

cuda
1234567
float *h_pinned;
cudaMallocHost(&h_pinned, bytes);  // pinned allocation

// Transfers with pinned memory are 2-3× faster
cudaMemcpy(d_data, h_pinned, bytes, cudaMemcpyHostToDevice);

cudaFreeHost(h_pinned);

Summary

Global memory is large but slow — minimize accesses and use coalesced patterns
Shared memory is fast on-chip memory for intra-block communication — always sync before reading
Constant memory is ideal for small, read-only, broadcast data
Unified Memory simplifies coding but explicit management gives more performance control
Pinned memory on the host accelerates host↔device transfers

Mastering memory is key to performance. In the next chapter, we'll look atKernels & Executionto understand how to launch and configure your kernels optimally.

Common Mistakes in CUDA Memory Management

Memory Leakage: Forgetting to call cudaFree on device pointers.
Uncoalesced Access: Accessing global memory in a non-contiguous pattern (explained in Chapter 7).
Race Conditions: Multiple threads writing to the same shared memory location without synchronization.

Practice Exercises

Try to implement a kernel that swaps two halves of an array using Shared Memory.
Compare the performance of cudaMallocManaged (Unified Memory) vs. explicit cudaMalloc and cudaMemcpy for a large vector addition.

Try it yourself

For an array of n floats, compute bytes = n * sizeof(float). Trace every allocation, host-to-device copy, kernel access, device-to-host copy, and free. Change the input to 1,003 elements and compare every output against a CPU reference.

Common mistake to avoid

A pointer value does not tell you whether its memory is accessible from the host or device. Track ownership and lifetime explicitly. Shared memory belongs to a block and must be initialized before use; allocation alone does not initialize device memory.

Check your understanding

When is shared memory worth using?

When threads in one block can reuse data enough to offset loading it, coordinating access, and consuming limited per-block resources. Measure the result; copying data to shared memory without reuse can add work.

Continue your reading in the official NVIDIA CUDA Programming Guide. Toolkit and GPU capabilities vary; check the documentation for your environment.

← PreviousProgramming Model Next →Kernels & Execution

Practice These Concepts

Reinforce what you just learned with hands-on GPU coding challenges.

Use CUDA constant memory to store a sine lookup table. Compute sin(x) for 1024 values using the lookup table instead of calling sin() directly. This demonstrates constant memory caching for read-only data accessed by all threads.