PMPP Chapter 5 Notes

• $AI_{\text{algorithmic}} = \frac{\frac{\text{FLOP}}{\text{second}}}{\frac{\text{logically requested bytes}}{\text{second}}}$
• $AI_{\text{HBM}} = \frac{\frac{\text{FLOP}}{\text{second}}}{\frac{\text{bytes crossing L2} \leftrightarrow \text{HBM}}{\text{second}}}$
• $AI_{\text{L2}} = \frac{\frac{\text{FLOP}}{\text{second}}}{\frac{\text{bytes crossing SM/L1} \leftrightarrow \text{L2 boundary}}{\text{second}}}$

• memory physical layer - connects memory controller to external memory.
  • needed as digital (memory controller) to memory (analog) connection
• Examples
  • CPU memory controller (digital) -> PCIe PHY -> analogue waves over PCIe physical infra -> RAM PCIe PHY -> digital signals to RAM memory controller, etc
  • GPU memory controller -> HBM PHY -> send analog waves to HBM

• individual physical wire within data bus used to transfer raw bits of data back/forth between memory stack and processor

• logic die <-(wires consisting of CA and DQ)-> memory dies
• CA bus is unidirectional (memory controller to memory die)
• DQ bus is bidirectional

• backed by global memory
• software visibility to threads only
• statically allocated arrays, spilled registers, etc

• Execution barrier prevents execution beyond it and synchronize all threads before proceeding
• Memory fence is about ordering

• Shared bus
• Crossbar
• Tree
• Ring
• Mesh
• Multistage network
• Some mixture

• load $2N^{2}$ matrices, and store a $N^{2}$ matrix -> $3n_bN^{2}$ of logical memory movement at RAM $\leftrightarrow$ HBM interface
• note, that memory movement at L2 $\leftrightarrow$ SM interface and HBM $\leftrightarrow$ L2 can be different

• register < local memory (non spilled over) < shared memory < global memory = local memory (spilled over)

Roofline Plot

• X axis can be thought of as how efficient am I w/ each unit of data that I load
• Y axis is "how much more is actually getting done per second"?
• One can see that if one is not very efficient w/ each unit of data load, i.e. load a unit of data then do minimal compute on it then transfer it back, one could imagine that the compute units are sitting idle while the memory loading path is bottlenecked. So the amount of compute done is relatively low.
• Note that there is sort of a bottom bound on how inefficient one could actually be w/ loading data - at worst you do one add or whatever operation per datum loaded, so in reality the roofline plot has a hard cutoff on the x axis at the left - the line will never go to 0.
• Worth noting that $$\frac{\text{FLOP}}{\text{s}} = \frac{\text{Byte}}{\text{s}} \frac{\text{FLOP}}{\text{Byte}}$$ inside the memory bound region. $\frac{\text{Byte}}{\text{s}}$ is how much bytes we can transfer, $\frac{\text{FLOP}}{\text{Byte}}$ is how much compute we can do per byte transferred. We assume here compute is not a bottleneck which is why there is a direct linear relationship between bytes transferred and compute done.
• After a certain amount of data efficiency the compute is the bottleneck. When that is the case the amount of FLOP done per unit time becomes constant because there exists a max throughput of $\tilde{\text{executors}} \frac{\text{FLOP}}{\text{Executor} \cdot \text{second}}$ throughput.

HBM data lanes are NOT duplex - i.e.

• The data wires support bidirectional but are buses
• i.e. lower bound for time to transfer data is $\frac{\text{Bytes read} + \text{Bytes stored}}{\text{HBM bandwidth}}$
• this is different from PCIe (or some other duplex hardware), where lower bound for time to transfer data is $\min\left(\frac{\text{Bytes read}}{\text{read bandwidth}}, \frac{\text{Bytes written}}{\text{write bandwidth}}\right)$

Tiling can very much cause low occupancy due to not enough shared memory if not done with consideration to max supported shared memory per thread.

It is possible to configure CUDA to allow SM to have more shared memory at the expense of other on chip memory resources

Rough memory hierarchy

• HBM stacks (stacked DRAM on package)
  • memory controllers and PHYs
    • Banked GPU-wide L2 (SRAM, on GPU die)
      • on chip fabric/crossbar
        • SM owned shared memory
          • registers

When do we want to turn something into analog?

• Note that PHY is needed for core to core and/or long distance communication because digital signals, afaik, have significant interference (why tho?). The factors at play are
  • distance
  • synchronous clock domain
  • low capacitance and voltage
• in a CPU core (i.e. core asking L1 cache for data) it's purely digital because low distance, synchronous clock domain and low capacitance and voltage.

Block admits into SMs are all or nothing - either all warps are resident or none are.

• Without this __syncthreads() can block for very long periods of time, or possibly even deadlock. Consider a block w/ more threads than the SM has to offer. It will block indefinitely on sync threads.

Compute done/byte of memory transferred ratio can be estimated as ~GPU flops / GPU max memory bandwidth. Ballpark estimate for kernels' required arithmetic intensity cutoff for compute vs memory bound

Energy consumed from accessing register at least order of mag smaller than DRAM

• scope w/in thread block

• scope is global

• if you do naive dot product like

  for (int k = 0; k < Width; ++k) {
      Pvalue += M[row * Width + k] * N[k * Width + col];
  }

• you end up w/ 2 global loads (8 bytes) per 2 flops (1 multiplication, 1 add)
  • which is extremely inefficient (0.25 FLOP/byte)
  • less than 1% of max throughput of H100!
  • if taking into account tensor cores (peak throughput 969 TFLOP/second) - less than 0.01% of the peak!!!