PMPP Chapter 1 Notes

Definitions

• Device
  • __device__ code to be executed on the device (i.e. GPU) only and only called by GPU code (i.e. __global__ or __device__ code)
• Host
  • __host__ code to be executed on the host (i.e. CPU) only
• Global
  • __global__ declares a kernel entry point that executes on the GPU. It is typically launched from host code with a grid and block configuration; with dynamic parallelism, it can also be launched from GPU code.
    • Worth noting, __host__ __device__ foo() is not the same as __global__ foo(). The former causes the compiler to generate separate host and device versions, while the latter generates a GPU kernel entry point.
• Warp
  • A group of 32 threads that an SM schedules together for execution
• SM
  • Streaming Multiprocessor
    • My understanding is that it's a hub for some shared register files (to be allocated dynamically to threads), a shared "L1" cache, and a bunch (~128 ish? Architecture dependent) of ALUs which support executing threads in parallel
    • Worth noting, multiple warps can occur at the same time on an SM

Notes

• GPUs execute simple programs in parallel across massive numbers of threads. The tradeoffs include:
  • No out-of-order execution in a single thread (AFAIK), i.e. decreased smartness
    • Worth noting, warps can still be parked while other warps execute
  • Increased latency
  • More programmatic decomposition (i.e. via hand or compiler) needed
  • Before Volta, all 32 threads in a warp shared a single program counter, while an active mask tracked which threads participated in the current instruction. Volta introduced per-thread execution state through Independent Thread Scheduling.
• Historically, GPUs became popular because of video games. Turns out they are also very relevant to ML workloads.
• Reducing latency can be much more expensive than increasing throughput in terms of power and chip area
  • CPUs focus highly on reducing latency