SpoderDude's blog
All posts

3 min read

PMPP Chapter 2 Notes

PMPP Chapter 2 Notes

notes, pmpp

Architecture

  • CUDA C++ Programs

    • Pseudo C++
    • Extends C++ compiler (rn supports up to C++20 afaik)
      • Compiles into both host (i.e. CPU) and device (i.e. GPU) code
      • Device code can be packaged for multiple target GPU architectures
        • Native cubins allow architecture-specific compilation, while embedded PTX can be JIT-compiled for a compatible GPU

  • How a CUDA kernel starts

Host program -> CUDA runtime/driver -> enqueue launch in a stream -> GPU dispatch -> thread block scheduled to an SM -> SM warp scheduler
  • Work submitted to a stream is asynchronous with respect to the CPU. Operations within one stream are ordered, so if the CPU enqueues kernel B after kernel A in the same stream, B becomes eligible to start after A completes without another acknowledgement from the CPU.
  • Work from different streams may overlap when dependencies and hardware resources allow, but concurrency is not guaranteed.
  • Stream synchronization, stream queries, CUDA events, and device synchronization are the main mechanisms for waiting on work or steering ordering. CUDA does not expose the underlying architecture behind stream communication and synchronization beyond those abstractions, so the implementation can change over time without breaking backwards compatibility.

Definitions

  • Data parallelism
    • Some operation can be applied in parallel across partitions of the data (i.e. adding matrices)
  • Task parallelism
    • Heterogeneous routines that can be done simultaneously (i.e. vector addition vs matrix multiplication)
  • CUDA Kernel Launch
    • Creates a grid of thread blocks
      • Since Hopper, a grid can group blocks into Thread Block Clusters for co-scheduling and distributed shared memory
      • Each block contains threads
        • Threads are grouped into Warps (32 threads at a time) to be run on SMs
  • PTX - Parallel Thread Execution
    • In between compiler IR and assembly ISA
    • Unlike x86 machine code, PTX is a virtual instruction set that can be compiled for different GPU generations
      • A fat binary can package target-specific cubins alongside PTX, allowing the runtime to load a compatible cubin or JIT-compile PTX
      • This allows GPU architectures to change more dramatically over time while preserving compatibility
  • nvcc - NVIDIA CUDA compiler driver
  • nvlink (not the physical NvLink connection) - device code linker
  • cuobjdump - objdump for gpu binaries
  • SASS - NVIDIA's architecture-specific native GPU instruction set
  • cubin - binary containing machine code for a particular GPU architecture
  • fatbin - container that can hold multiple cubins and PTX

Syntax

  • <<<grid dimensions, block dimensions>>>
  • cudaMalloc(&f_device_ptr, n_bytes)
  • cudaFree(f_device_ptr)
  • cudaMemcpy(dest_ptr, src_ptr, n_bytes, {cuda copy direction})
  • cudaError_t
  • Row-major access - CUDA does not prescribe which index is a row or column. The kernel chooses the mapping; for a 2D launch, a common convention is the x dimension for columns and the y dimension for rows.

Numbers

  • A thread block can contain at most 1024 threads on current GPUs; individual x/y/z dimension limits also apply