Skip to content

[Ops] Rewrite NPU chunk_scaled_dot_kkt_fwd#1023

Open
OsirisDuan wants to merge 2 commits into
fla-org:mainfrom
OsirisDuan:260714/chunk_scaled_dot_kkt_fwd_kernel_perf_opti
Open

[Ops] Rewrite NPU chunk_scaled_dot_kkt_fwd#1023
OsirisDuan wants to merge 2 commits into
fla-org:mainfrom
OsirisDuan:260714/chunk_scaled_dot_kkt_fwd_kernel_perf_opti

Conversation

@OsirisDuan

@OsirisDuan OsirisDuan commented Jul 14, 2026

Copy link
Copy Markdown
Contributor

summary

  • Replace 2D (NT, B*HV) grid with 1D (num_core,) + tl.range task-loop to improve
    Ascend core occupancy.
  • Compute the full [BT,BT] A matrix in a single pass instead
    of iterating BC=16 subchunks, reducing tl.dot launches.
  • Switch from UB-aware BK
    sizing to @triton.autotune over BK={32,64,128}.
  • Permute beta/g to (HV,B,T) for
    unit-stride block pointers. Drop _launch_kkt_kernel and _get_bk helpers.

benchmark

python benchmarks/ops/run.py --op chunk_gdn

===========================================================================================================================
  Machine: Ascend910B3 | NPU: CANN 9.0.0-beta.2 | PyTorch 2.7.1+cpu
===========================================================================================================================
  fwd        B      T    H    D  op                                   main[ccb0ff94](ms) 260714/c...[3800d902](ms)  speedup
          -----------------------------------------------------------------------------------------------------------------
             1   8192   96  128  chunk_gdn                                      1043.344                   813.850    1.28x
          -----------------------------------------------------------------------------------------------------------------
             2  16384   16  128  chunk_gdn                                       694.422                   545.404    1.27x
          -----------------------------------------------------------------------------------------------------------------
             4   2048   16  128  chunk_gdn                                       172.372                   134.716    1.28x
          -----------------------------------------------------------------------------------------------------------------
             4   4096   64  128  chunk_gdn                                      1390.365                  1086.051    1.28x
          -----------------------------------------------------------------------------------------------------------------
             8   1024    8   64  chunk_gdn                                        52.917                    34.514    1.53x
          -----------------------------------------------------------------------------------------------------------------
             8   2048   32  256  chunk_gdn                                      1296.023                  1141.177    1.14x
===========================================================================================================================
  fwdbwd     B      T    H    D  op                                   main[ccb0ff94](ms) 260714/c...[3800d902](ms)  speedup
          -----------------------------------------------------------------------------------------------------------------
             1   8192   96  128  chunk_gdn                                      4063.395                  3843.020    1.06x
          -----------------------------------------------------------------------------------------------------------------
             2  16384   16  128  chunk_gdn                                      2764.333                  2613.041    1.06x
          -----------------------------------------------------------------------------------------------------------------
             4   2048   16  128  chunk_gdn                                       686.859                   643.030    1.07x
          -----------------------------------------------------------------------------------------------------------------
             4   4096   64  128  chunk_gdn                                      5460.384                  5139.979    1.06x
          -----------------------------------------------------------------------------------------------------------------
             8   1024    8   64  chunk_gdn                                       180.857                   162.939    1.11x
          -----------------------------------------------------------------------------------------------------------------
             8   2048   32  256  chunk_gdn                                      6179.791                  6033.952    1.02x
===========================================================================================================================

@OsirisDuan
OsirisDuan force-pushed the 260714/chunk_scaled_dot_kkt_fwd_kernel_perf_opti branch from 7fa48f0 to 6135b98 Compare July 16, 2026 09:35
@OsirisDuan
OsirisDuan force-pushed the 260714/chunk_scaled_dot_kkt_fwd_kernel_perf_opti branch from 6135b98 to 3800d90 Compare July 17, 2026 10:08
@OsirisDuan

Copy link
Copy Markdown
Contributor Author

@zheliuyu Please review this PR.

Comment on lines +33 to +42
@triton.autotune(
configs=[
triton.Config({'BK': BK}, num_warps=num_warps, num_stages=num_stages)
for BK in [32, 64, 128]
for num_warps in [2, 4, 8]
for num_stages in [2, 3, 4]
],
key=['H', 'HV', 'K', 'BT', 'IS_VARLEN'],
**autotune_cache_kwargs,
)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

blocking: full [BT,BT] + autotune BK≤128 drops the UB-aware sizing used elsewhere (#1001). Fine if faster, but please clamp/filter configs or confirm the UB envelope on 910B/C so we don’t hit intermittent UB/compile failures outside the bench shapes.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

nit: on triton-ascend, num_warps/num_stages are ignored (see Ascend docs /
"[WARNING] Please DO NOT tune arg ..."). Prefer autotuning BK only (and maybe
multibuffer) to avoid a 3×3 wasted search.

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