Modern Architecture Features

cuda::pipeline / cuda::memcpy_async with multi-stage pipelines enable double/multi-buffering: load tile k+1 while computing on tile k.

28.1 Asynchronous copy global → shared (Ampere+)

cuda::pipeline / cuda::memcpy_async with multi-stage pipelines enable double/multi-buffering: load tile k+1 while computing on tile k.

#include <cuda/pipeline>
#include <cooperative_groups/memcpy_async.h>

__global__ void asyncTile(const float* in, float* out) {
    __shared__ float tile[256];
    auto block = cg::this_thread_block();
    cg::memcpy_async(block, tile, in + blockIdx.x * 256, sizeof(float) * 256);   // bypasses registers
    cg::wait(block);                                                             // wait for completion
    out[blockIdx.x * 256 + threadIdx.x] = tile[threadIdx.x] * 2.0f;
}

28.2 Thread block clusters & distributed shared memory (Hopper+)

Clusters group blocks scheduled together on a GPC, enabling fast block-to-block communication via distributed shared memory (DSMEM).

__global__ void __cluster_dims__(2, 1, 1) clusterKernel(float* data) {
    namespace cg = cooperative_groups;
    cg::cluster_group cluster = cg::this_cluster();
    __shared__ float smem[256];
    unsigned rank = cluster.block_rank();
    smem[threadIdx.x] = rank;
    cluster.sync();
    float* peer = cluster.map_shared_rank(smem, rank ^ 1);   // access another block's shared memory
    data[blockIdx.x * 256 + threadIdx.x] = peer[threadIdx.x];
    cluster.sync();
}

28.3 Tensor Memory Accelerator (TMA, Hopper+)

Hardware unit that copies multi-dimensional tiles between global and shared memory asynchronously, described by a tensor map (cuTensorMapEncodeTiled), with mbarrier synchronization — used by CUTLASS 3 and FlashAttention-3 for maximum bandwidth with minimal thread involvement.

28.4 Other recent features

Warp-specialized kernels (producer warps load data, consumer warps compute).

mbarrier / cuda::barrier split arrive/wait barriers.

FP8 Transformer Engine (Hopper), FP4 and 2nd-gen Transformer Engine (Blackwell).

Confidential computing (H100), MIG partitions, NVLink switch systems (NVL72).

Programmatic dependent launch (overlap the next kernel's prologue with the current kernel's tail).

Green contexts (SM partitioning within a process), CUDA 12.x/13 driver features.