Please enable JavaScript.

Coggle requires JavaScript to display documents.

QServe - Coggle Diagram

- - - - int4->int8, dequant might overflow, but saturation instruction is 1/3 as slow, use 119 instead of 127 as bound.
      - Compared to DGQ, this fuses int4->int8(CUDA core) in GEMM kernel (Tensor Core)
    - - directly dequantize wu4g weights to wf16
  - - - shift quant from key to query, similar as SmoothQuant shift from activations to weights
      - select the lambda as max(|Ki|)^alpha
      - fuse lambda to WK and WQ: consider ROPE, where channel i and i+D/2 are paired, we let them have same lambda
  - - - Cur rotate - 1. QKV Projection 2. FFN up layer
        Prev rev rot- 1. FFN down 2. ATTN out
        
        Use scaled Hadamard to suppress outliers
    - - cur smooth - 1. FFN down 2. ATTN out
        prev smooth rev - 1. FFN up 2. ATTN v
    - - to make channels with similar sailence grouped together for quantization, similar to reserve fp16 for sailent weights
    - - finding clip ratio, by minimize MSE of layer/block output bef/after quant
- - - - Make ldmatrix usable by reorder each 32x32 tiles for int4 weights zero points and scales
    - - Register-Level Parallelism
        
        uint4 -> uint8: w0w16, w1w17..., 3 ops
      - Substraction after mul
        
        uint8 -> sint8: -xzs after gemm
    - - subs before mul: int4->int8 -> >int8 fail, overflow
        subs after mul: int4 -> int8 -> int8 okay
    - - Multi-stage pipelining + async mem copy
      - bank conflict: L1 shared mem
      - tile partition, weight reuse
      - K parallel for small L
  - - - per-head dynamic kv4 is slow on A100
      - dequant kv4 to fp32 is slow CUDA core on A100
    - - Dequant to fp16 instead of fp32
      - reducing dequant ops from 5 to 2
      - simplify control logic, prefetch scale and zero
  - - - ATTN: fp16 in CUDA core
      - Whole LLM block: fp16 -> fp16
      - GEMM: w4a8 -> fp16
    - - QKV: prev layernorm
      - FFN up: prev layernorm
      - FFN down: prev act
      - ATTN out: insert a quant node ahead
    - - Paged Cache
        
        sfp16, zu4 for each head, on-th-fly
        
        in-flight batching