Niko1221 5a30c2c4e8 Images (I2): NVIDIA mma.sync flash attention for the native runner, fused q/k/v GEMM ...

- k_attn_mma (sm_80+): m16n8k16 f16 -> f32, 8 warps x 16 queries, 64-key K/V tiles double-buffered with
  cp.async, ldmatrix (.trans for V), fp32 online softmax.  Bench on the RTX 3060 at 4096 x 4150 keys x 32
  heads: 23.7 TFLOPS vs ggml's flash_attn_ext_f16 15.4 (0.376 vs 0.577 s/step); CPU check 7e-5.
  STRATA_QI_GGML_ATTN=1 keeps ggml's.
- q, k and v of a block with one weight type load into one [3C x C] tensor: one GEMM and one activation
  quantize instead of three (STRATA_QI_NO_QKV_FUSE=1 keeps three).
- A float4 LayerNorm path with tanh(gate) once per step, off by default (STRATA_QI_FAST_LN=1): no measured gain.

RTX 3060, Q8_0 1024, seed 42, 3 interleaved rounds: steady step 2.09 s vs stock 2.61 (1.25x), denoise
17.93 s vs 22.13 (1.23x).

2026-10-13 05:53:21 +02:00

..

sdcpp_native

Images (I2): NVIDIA mma.sync flash attention for the native runner, fused q/k/v GEMM

2026-10-13 05:53:21 +02:00

install.py

Images (I2): setup builds the native DiT runner into sd-cli on CUDA / HIP; --image-native on|off (off by default) and the config's "native" add --model-args qwen_image_2_1_native=1

2026-10-13 01:07:28 +02:00

qimage.py

Images (I2): setup builds the native DiT runner into sd-cli on CUDA / HIP; --image-native on|off (off by default) and the config's "native" add --model-args qwen_image_2_1_native=1

2026-10-13 01:07:28 +02:00