Skip to content

Add NVIDIA Turing support to FlashMMA kernel - #1246

Merged
lightvector merged 2 commits into
lightvector:masterfrom
nlzy:nlzy/turing
Aug 30, 2026
Merged

Add NVIDIA Turing support to FlashMMA kernel#1246
lightvector merged 2 commits into
lightvector:masterfrom
nlzy:nlzy/turing

Conversation

@nlzy

@nlzy nlzy commented Aug 28, 2026

Copy link
Copy Markdown
Contributor

Summary

The flash attention kernel in cpp/neuralnet/cudaflashmma.cuh needed sm_80+, because it used mma.sync.aligned.m16n8k16 and cp.async.

On NVIDIA Turing GPUs (sm_75), transformer models fallback to a much slower attention path. cuDNN graph SDPA also needs sm_80, so it is not available either.

This PR added Turing support to cudaflashmma.cuh, with a +70.3% speedup compared to original CUDA backend on transformer models.

Environment

  • Intel Xeon CPU E5-2620 v4 @ 2.10GHz
  • NVIDIA RTX 2080 Ti (22GB mod)
  • Ubuntu 24.04 with Linux 6.8 kernel
  • CUDA 12.9, cuDNN 9.10.2, TensorRT 10.12.0
  • Model: kata1-tf3-b11c768-s11001M-d5973M.bin.gz

Testing

  • kata-raw-nn on the empty board (komi 7.5), mma on vs off: whiteWin 0.655850 vs 0.656328, whiteLead 0.876 vs 0.882.
  • GTP genmove sanity check passes.
  • ./katago runtests: all pass.
  • ./katago runnnlayertests: all pass.
  • cpp/rungpuerrortest.sh: all pass.
  • cpp/runcudaopttests.sh with EXPECT_FUSED_FFN=0 EXPECT_SDPA=0: all pass.

Performance

./katago benchmark -config gtp.cfg -model <tf3 model> -v 5000 -n 5

CUDA builds use numNNServerThreadsPerModel = 2. TensorRT use 1.

Build Best nnEvals/s visits/s Config (batch / threads) vs Origin (CUDA)
Origin (CUDA) 400.13 661.56 16 / 32
Origin (TRT) 589.89 994.32 16 / 32 +47.4%
This PR (CUDA) 681.60 1161.40 16 / 32 +70.3%

This PR is also +15.5% faster than Origin TensorRT backend in nnEvals/s.

Full sweep (nnEvals/s):

Build batch t=8 t=16 t=24 t=32
Origin (CUDA) 8 327.11 383.63 391.62 391.25
Origin (CUDA) 16 325.12 376.99 394.48 400.13
Origin (CUDA) 32 324.59 375.04 388.73 395.49
Origin (TRT) 8 466.90 542.44 541.67 542.60
Origin (TRT) 16 488.71 544.74 569.59 589.89
Origin (TRT) 32 484.76 551.48 569.79 581.18
This PR (CUDA) 8 564.38 641.43 657.90 654.25
This PR (CUDA) 16 556.40 624.94 659.96 681.60
This PR (CUDA) 32 556.75 623.48 652.76 672.68

@lightvector
lightvector merged commit e85340f into lightvector:master Aug 30, 2026
@lightvector

Copy link
Copy Markdown
Owner

Thanks!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants