Skip to content

[Triton/Gluon] : Change RotatingBuffersBench API and update any existing benchmarks - #7121

Merged
rahulbatra85 merged 7 commits into
mainfrom
triton_gen_rot_buf_updates
Nov 3, 2026
Merged

rahulbatra85 merged 7 commits into
mainfrom
triton_gen_rot_buf_updates

Conversation

@rahulbatra85

@rahulbatra85 rahulbatra85 commented Oct 29, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

--Currently, only tensors are accepted for building rotating buffers pool. Allow list/dict/tuple of tensors in generate_rotating_buffers_pool
--Separate out RotatingBuffer and do_bench_aiter_triton into it's own file
--pool target_mb is not set based on the arch/device

PR Self-Review

What behavior does this PR change, and why is the existing behavior insufficient?
-separate rotating buffers and do_bench_aiter_triton into it's own file
-new API using class called RotatingBufferBench
-use defaults for target_mb based on arch.
-determine max_gpu_mb based on free device memory available when creating the pool

What did you validate, and what remains unverified? Include commands, hardware/configuration and results.
-Just tried the change with benchmark that needs it.

What behavior would each changed op test fail on before this PR, and why is that coverage not already provided by an existing canonical test?
--This adds functionality for Triton benchmarks. There are no test for benchmarks.

For every changed op_tests/ file, what behavior does it cover and why is it not redundant? Write None if no op test changed.
-This adds functionality for Triton benchmarks. There are no test for benchmarks.

@rahulbatra85
rahulbatra85 requested a review from a team October 29, 2026 00:48
@github-actions

Copy link
Copy Markdown
Contributor

🏷️ CI Guide

Runs automatically on every PR:

  • ✅ Pre-checks (submodule verification, code formatting)
  • ✅ Aiter op tests (gfx942 + gfx950)
  • ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label Tests
ci:gfx1250-ffm-triton Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x Run an additional Triton test job on MI300X in PRs (added automatically when gfx942 configs change); main branch always runs both MI35X and MI300X
ci:triton-355 Run the full Triton test suite on MI35X, not only the tests the change affects
ci:windows Windows JIT op smoke test on the gfx1201 ROCm runner
ci:multi-gpu Aiter multi-GPU tests on the 8-GPU runner
ci:sglang SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm vLLM benchmark: DeepSeek-V4-Pro, MiniMax-M3-MXFP4
ci:all All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 7121 --add-label <label>

One backend per PR:
A PR changes one kernel backend: [Triton/Gluon] (Triton and Gluon count as one), [HIP], [ASM], [CK], [OPUS] or [FlyDSL]. If the title ends up with two backend tags, split the PR -- as stacked pull requests when one part cannot merge without the other.

PR title tags & labels:
Component tags ([Triton/Gluon], [HIP], [CK], [ASM], ...) are added to the PR title and as PR labels automatically from the changed files and re-synced on every push — change-type tags like [fix]/[Perf], op tags like [MLA], and human labels (ci:*) are left untouched. Add the no-auto-title label to stop the title rewrites; labels stay in sync either way.

@github-actions github-actions Bot changed the title [Triton/Gluon]: Add exclude list arg and allow list/dict/tuple of tensors in generate_rotating_buffers_pool [Triton/Gluon] : Add exclude list arg and allow list/dict/tuple of tensors in generate_rotating_buffers_pool Oct 29, 2026
@Boss2002n
Boss2002n requested a balanced review from Copilot October 29, 2026 01:33

Copilot AI left a comment

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.

🟡 Changes recommended

The sizing loop retains its last tensor sample during pool allocation, increasing peak GPU memory and potentially causing an otherwise fitting pool to run out of memory.

1 open issue
What changed in this PR

This PR extends the Triton/Gluon benchmark helper to size rotating buffer pools from nested tensor inputs and exclude selected fields from sizing.

Changes:

  • Adds exclude_from_sizing to omit named fields from the memory estimate.
  • Recursively counts tensors in dictionaries, lists, and tuples while preserving storage deduplication.
File Description
op_tests/​op_benchmarks/​triton/​utils/​benchmark_utils.py Adds field exclusions and recursive tensor counting for rotating pool sizing.

Review effort: Balanced

@rahulbatra85
rahulbatra85 force-pushed the triton_gen_rot_buf_updates branch from 31f3ded to ecf63c2 Compare October 29, 2026 16:45
@aiter-bot

aiter-bot commented Oct 29, 2026 •

Copy link
Copy Markdown

NOT READY FOR HUMAN REVIEW

PR Self-Review

Revision 88f4014d83b0 · mode shadow

Waiting for the PR author. The readiness report will appear after the self-review answers are complete for this revision.

Questions to answer in the PR description

  1. What behavior does this PR change, and why is the existing behavior insufficient?
  2. What did you validate, and what remains unverified? Include commands, hardware/configuration and results.
  3. What behavior would each changed op test fail on before this PR, and why is that coverage not already provided by an existing canonical test?
  4. For every changed op_tests/ file, what behavior does it cover and why is it not redundant? Write None if no op test changed.

Edit the PR Self-Review section in the PR description. Include the current revision SHA; this comment refreshes automatically.

Missing or invalid answers

  • FAIL — missing ### Reviewed revision in the PR description
  • FAIL — missing ### Behavior and intent in the PR description
  • FAIL — missing ### Validation and unverified scope in the PR description
  • FAIL — missing ### PR-specific question in the PR description
  • FAIL — missing ### Changed op tests in the PR description
  • FAIL — Reviewed revision must contain at least the first 7 characters of the current head SHA
  • FAIL — Behavior and intent has no substantive answer
  • FAIL — Validation and unverified scope has no substantive answer
  • FAIL — PR-specific question has no substantive answer
  • FAIL — changed op test not explained in Changed op tests: op_tests/op_benchmarks/triton/AGENTS.md
  • FAIL — changed op test not explained in Changed op tests: op_tests/op_benchmarks/triton/README.md
  • FAIL — changed op test not explained in Changed op tests: op_tests/op_benchmarks/triton/bench_all_reduce.py
  • FAIL — changed op test not explained in Changed op tests: op_tests/op_benchmarks/triton/bench_all_reduce_rmsnorm_quant.py
  • FAIL — changed op test not explained in Changed op tests: op_tests/op_benchmarks/triton/bench_batched_gemm_a16wfp4.py
  • FAIL — changed op test not explained in Changed op tests: op_tests/op_benchmarks/triton/bench_batched_gemm_a8w8.py
  • FAIL — changed op test not explained in Changed op tests: op_tests/op_benchmarks/triton/bench_batched_gemm_a8w8_a_per_token_group_prequant_w_per_batched_tensor_quant.py
  • FAIL — changed op test not explained in Changed op tests: op_tests/op_benchmarks/triton/bench_batched_gemm_afp4wfp4.py
  • FAIL — 48 additional changed op test(s) are not explained
  • FAIL — Changed op tests must explain why each changed op test is necessary and non-redundant

Shadow mode reports what would block but keeps the status successful. Set AITER_PR_SELF_REVIEW_MODE=enforce only after reviewing real false positives.

@zufayu
zufayu requested review from a team and azaidy October 30, 2026 01:28
@Boss2002n Boss2002n assigned rahulbatra85 and unassigned azaidy Oct 30, 2026
@rahulbatra85
rahulbatra85 force-pushed the triton_gen_rot_buf_updates branch from ecf63c2 to cb4a98e Compare October 31, 2026 21:45

@github-actions github-actions Bot left a comment

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.

Remaining comments which cannot be posted as a review comment to avoid GitHub Rate Limit

ruff

⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
from aiter.ops.triton.gemm.fused.fused_gemm_a8w8_blockscale_a16w16 import (
fused_gemm_a8w8_blockscale_a16w16,
)
from op_tests.op_benchmarks.triton.bench_fused_gemm_commons import (
metric_to_scalar,
parse_fused_args,
run_fused_benchmark,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8w8_blockscale import (
generate_gemm_a8w8_blockscale_inputs,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w16 import (
generate_gemm_a16w16_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
from aiter.ops.triton.gemm.fused.fused_gemm_a8w8_blockscale_mul_add import (
fused_gemm_a8w8_blockscale_mul_add,
)
from op_tests.op_benchmarks.triton.bench_fused_gemm_commons import (
make_mul_add_ab,
metric_to_scalar,
parse_fused_args,
run_fused_benchmark,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8w8_blockscale import (
generate_gemm_a8w8_blockscale_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
from aiter.ops.triton.gemm.fused.fused_gemm_afp4wfp4_a16w16 import (
fused_gemm_afp4wfp4_a16w16,
)
from op_tests.op_benchmarks.triton.bench_fused_gemm_commons import (
metric_to_scalar,
parse_fused_args,
run_fused_benchmark,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w16 import (
generate_gemm_a16w16_inputs,
)
from op_tests.triton_tests.gemm.basic.test_gemm_afp4wfp4 import (
generate_gemm_afp4wfp4_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
from aiter.ops.triton.gemm.fused.fused_gemm_afp4wfp4_mul_add import (
fused_gemm_afp4wfp4_mul_add,
)
from op_tests.op_benchmarks.triton.bench_fused_gemm_commons import (
make_mul_add_ab,
metric_to_scalar,
parse_fused_args,
run_fused_benchmark,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_afp4wfp4 import (
generate_gemm_afp4wfp4_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import torch
import triton
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.fusions.test_fused_mul_add import (
generate_fused_mul_add_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import sys
import torch
import triton
from aiter.ops.triton.fusions.fused_indexed_gated_residual import (
fused_indexed_gated_residual,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.fusions.test_fused_indexed_gated_residual import make


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
import triton.language as tl
from einops import rearrange
from aiter.ops.triton._triton_kernels.gated_delta_net.decode.fused_sigmoid_gating_recurrent import (
fused_sigmoid_gating_delta_rule_update,
)
from aiter.ops.triton.gated_delta_net.causal_conv1d_decode import (
causal_conv1d_update_split_qkv,
)
from aiter.ops.triton.gated_delta_net.fused_kda_decode import fused_kda_decode
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import sys
import torch
import triton
from aiter import dtypes
from aiter.ops.triton.fusions.fused_mla_prefill_qkv_fp8 import (
fused_mla_prefill_qkv_fp8,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.fusions.test_fused_mla_prefill_qkv_fp8 import (
generate_inputs,
torch_mla_prefill_qkv_fp8_ref,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
from aiter.ops.triton.fusions.fused_kv_cache import (
fused_qk_cat_and_cache_mla,
fused_qk_rope_cat_and_cache_mla,
)
from aiter.ops.triton.utils._triton import arch_info
from aiter.ops.triton.utils.types import e4m3_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
from aiter.ops.triton.rope.fused_qk_norm_rope_gate_fp8_quant import (
FP8_DTYPE,
MAX_SEQUENCES,
fused_qk_norm_rope_gate_fp8_quant,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import argparse
import functools
import torch
import triton
from aiter.ops.triton.fusions.fused_kv_cache import fused_qk_rope_reshape_and_cache
from aiter.ops.triton.utils._triton import arch_info
from aiter.ops.triton.utils.types import e4m3_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import argparse
import functools
import torch
import triton
from aiter.ops.triton.fusions.fused_reduce_qk_norm_rope_swa_write import (
fused_reduce_qk_norm_rope_swa_write,
)
from aiter.ops.triton.utils.types import str_to_torch_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.fusions.test_fused_reduce_qk_norm_rope_swa_write import (
_build_cos_sin,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import sys
import torch
import torch.nn.functional as F
import triton
from aiter.ops.triton.quant import dynamic_mxfp4_quant, fused_rms_gated_mxfp4_quant
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import triton
from aiter.ops.triton.normalization.fused_rmsnorm_add import fused_rmsnorm_add
from aiter.ops.triton.utils.types import str_to_torch_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.normalization.test_fused_rmsnorm_add import (
generate_fused_rmsnorm_add_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
from aiter.ops.triton.normalization.fused_rmsnorm_add_rmsnorm import (
fused_rmsnorm_add_rmsnorm,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.normalization.test_fused_rmsnorm_add_rmsnorm import (
generate_fused_rmsnorm_add_rmsnorm_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import sys
import torch
import triton
from aiter.ops.triton.fusions.fused_sigmoid_mul import fused_sigmoid_mul
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.fusions.test_fused_sigmoid_mul import (
generate_fused_sigmoid_mul_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from aiter.ops.triton.gemm.basic.gemm_a16w16 import _is_gluon_available, gemm_a16w16
from aiter.ops.triton.gemm.basic.gemm_a16w16_atomic import gemm_a16w16_atomic
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w16 import (
generate_gemm_a16w16_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from aiter.ops.triton.gemm.basic.gemm_a16w16_gated import gemm_a16w16_gated
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w16_gated import (
generate_gemm_a16w16_gated_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import sys
import matplotlib.pyplot as plt
import torch
import triton
from aiter.ops.triton.gemm.basic.gemm_a16w16_gated import gemm_a16w16_gated
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_shape_benchmark_object,
model_benchmark_shapes,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w16_gated import (
generate_gemm_a16w16_gated_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from aiter import logger
from aiter.ops.triton.gemm.basic.gemm_a16w4 import gemm_a16w4
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w4 import (
generate_gemm_a16w4_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from aiter.ops.triton.gemm.basic.gemm_a16w8_blockscale import (
gemm_a16w8_blockscale,
gemm_a16w8_blockscale_preshuffle,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w8_blockscale import (
generate_gemm_a16w8_blockscale_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import sys
from collections.abc import Callable
import triton
from aiter.ops.triton.gemm.basic.gemm_a8w8 import gemm_a8w8 as triton_gemm_a8w8
from aiter.ops.triton.gemm.basic.gemm_a8w8 import gemm_a8w8_preshuffle
from aiter.ops.triton.utils.types import str_to_torch_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8w8 import (
generate_gemm_a8w8_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
from collections.abc import Callable
import torch
import triton
from aiter.ops.triton.gemm.basic.gemm_a8w8_blockscale import (
gemm_a8w8_blockscale as triton_gemm_a8w8_blockscale,
)
from aiter.ops.triton.gemm.basic.gemm_a8w8_blockscale import (
gemm_a8w8_blockscale_preshuffle as triton_gemm_a8w8_blockscale_preshuffle,
)
from aiter.test_common import checkAllclose
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8w8_blockscale import (
generate_gemm_a8w8_blockscale_inputs,
run_torch,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from aiter.ops.triton.gemm.basic.gemm_a8w8_per_token_scale import (
gemm_a8w8_per_token_scale,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8w8_per_token_scale import (
generate_gemm_a8w8_per_token_scale_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import sys
import torch
import triton
from aiter.ops.triton.gemm.basic.gemm_a8wfp4 import (
gemm_a8wfp4,
)
from aiter.ops.triton.utils.types import get_fp8_dtypes
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8wfp4 import (
generate_gemm_a8wfp4_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from aiter.ops.triton.gemm.basic.gemm_afp4wfp4 import (
gemm_afp4wfp4 as triton_gemm_afp4wfp4,
)
from aiter.ops.triton.gemm.basic.gemm_afp4wfp4 import (
gemm_afp4wfp4_preshuffle,
)
from aiter.ops.triton.utils._triton import arch_info
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.gemm.basic.test_gemm_afp4wfp4 import (
generate_gemm_afp4wfp4_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import csv
import functools
import math
import warnings
from pathlib import Path
import torch
import triton
from aiter.ops.shuffle import shuffle_weight
from aiter.ops.triton.gemm.basic.gemm_afp8wfp8 import (
gemm_afp8wfp8,
gemm_afp8wfp8_preshuffle,
)
from aiter.ops.triton.utils._triton import arch_info
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import argparse
import functools
import logging
# PyTorch
import torch
# Triton
import triton
# AITER: Triton kernel wrappers
from aiter.ops.triton.gemm.grouped.gmm import (
gmm as triton_gmm,
)
from aiter.ops.triton.gemm.grouped.gmm import (
nptgmm as triton_nptgmm,
)
from aiter.ops.triton.gemm.grouped.gmm import (
ptgmm as triton_ptgmm,
)
# AITER: GMM defaults and utility functions
from aiter.ops.triton.utils.gmm_common import (
DTYPE,
DTYPE_STR,
GROUP_SIZES_DTYPE,
GROUP_SIZES_DTYPE_STR,
NUM_GROUP_SIZES,
RNG_SEED,
SUPPORTED_DTYPES_STR,
SUPPORTED_GROUP_SIZES_DTYPES_STR,
TRANS_LHS,
TRANS_RHS,
dtype_from_str,
gen_gmm_input,
gen_gmm_output,
gen_multiple_group_sizes,
gen_tgmm_bias_grad,
gen_tgmm_input,
gen_tgmm_output,
group_sizes_dtype_from_str,
str_from_dtype,
str_from_group_sizes_dtype,
)
# AITER: shared argparse helper
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
positive_int,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
from aiter.ops.triton.normalization.norm import (
layer_norm,
layernorm2d_fwd_with_add,
layernorm2d_fwd_with_add_dynamicquant,
layernorm2d_fwd_with_add_smoothquant,
layernorm2d_fwd_with_dynamicquant,
layernorm2d_fwd_with_smoothquant,
)
from aiter.ops.triton.utils.types import str_to_torch_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
from aiter.ops.triton.attention.mla import mla_decode_fwd, mla_prefill_fwd
from aiter.ops.triton.utils._triton import arch_info
from aiter.ops.triton.utils.types import e4m3_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.attention.test_mla import input_helper, shuffle_kv_buffer


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import argparse
import functools
import itertools
import sys
import torch
import triton
from aiter.ops.triton.attention.mla_decode_rope import (
decode_attention_fwd_grouped_rope,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_configs,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.attention.test_mla_decode_rope import (
input_helper,
ref_preprocess,
)
from op_tests.triton_tests.utils.rotary_embedding import DeepseekScalingRotaryEmbedding


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from aiter.ops.triton.normalization.rmsnorm import rms_norm
from aiter.ops.triton.quant.fused_mxfp4_quant import fused_rms_mxfp4_quant
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_configs,
print_vgpr,
)
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.normalization.test_rmsnorm import (
generate_rmsnorm_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
from aiter.ops.triton.rope.rope import (
RotateStyle,
rope_bwd,
rope_cached_bwd,
rope_cached_fwd,
rope_cached_fwd_inplace,
rope_cached_positions_bwd,
rope_cached_positions_fwd,
rope_cached_positions_fwd_inplace,
rope_cached_positions_offsets_bwd,
rope_cached_positions_offsets_fwd,
rope_cached_positions_offsets_fwd_inplace,
rope_cached_thd_positions_2c_bwd,
rope_cached_thd_positions_2c_fwd,
rope_cached_thd_positions_2c_fwd_inplace,
rope_cached_thd_positions_offsets_2c_bwd,
# rope_fwd_2d,
# rope_fwd_2d_inplace,
rope_cached_thd_positions_offsets_2c_fwd,
rope_cached_thd_positions_offsets_2c_fwd_inplace,
rope_fwd,
rope_fwd_inplace,
rope_thd_bwd,
rope_thd_fwd,
rope_thd_fwd_inplace,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_model_configs
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.triton_tests.rope.test_rope import generate_rope_inputs

@rahulbatra85
rahulbatra85 force-pushed the triton_gen_rot_buf_updates branch 2 times, most recently from cdc8ca0 to 5332c56 Compare November 2, 2026 17:48

@github-actions github-actions Bot left a comment

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.

Remaining comments which cannot be posted as a review comment to avoid GitHub Rate Limit

ruff

⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.fused.fused_gemm_a8w8_blockscale_mul_add import (
fused_gemm_a8w8_blockscale_mul_add,
)
from op_tests.op_benchmarks.triton.bench_fused_gemm_commons import (
make_mul_add_ab,
metric_to_scalar,
parse_fused_args,
run_fused_benchmark,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8w8_blockscale import (
generate_gemm_a8w8_blockscale_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.fused.fused_gemm_afp4wfp4_a16w16 import (
fused_gemm_afp4wfp4_a16w16,
)
from op_tests.op_benchmarks.triton.bench_fused_gemm_commons import (
metric_to_scalar,
parse_fused_args,
run_fused_benchmark,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w16 import (
generate_gemm_a16w16_inputs,
)
from op_tests.triton_tests.gemm.basic.test_gemm_afp4wfp4 import (
generate_gemm_afp4wfp4_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.fused.fused_gemm_afp4wfp4_mul_add import (
fused_gemm_afp4wfp4_mul_add,
)
from op_tests.op_benchmarks.triton.bench_fused_gemm_commons import (
make_mul_add_ab,
metric_to_scalar,
parse_fused_args,
run_fused_benchmark,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_afp4wfp4 import (
generate_gemm_afp4wfp4_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.triton_tests.fusions.test_fused_mul_add import (
generate_fused_mul_add_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import sys
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.fusions.fused_indexed_gated_residual import (
fused_indexed_gated_residual,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.triton_tests.fusions.test_fused_indexed_gated_residual import make


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
import triton.language as tl
from einops import rearrange
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton._triton_kernels.gated_delta_net.decode.fused_sigmoid_gating_recurrent import (
fused_sigmoid_gating_delta_rule_update,
)
from aiter.ops.triton.gated_delta_net.causal_conv1d_decode import (
causal_conv1d_update_split_qkv,
)
from aiter.ops.triton.gated_delta_net.fused_kda_decode import fused_kda_decode
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import sys
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter import dtypes
from aiter.ops.triton.fusions.fused_mla_prefill_qkv_fp8 import (
fused_mla_prefill_qkv_fp8,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.triton_tests.fusions.test_fused_mla_prefill_qkv_fp8 import (
generate_inputs,
torch_mla_prefill_qkv_fp8_ref,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.fusions.fused_kv_cache import (
fused_qk_cat_and_cache_mla,
fused_qk_rope_cat_and_cache_mla,
)
from aiter.ops.triton.utils._triton import arch_info
from aiter.ops.triton.utils.types import e4m3_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.fusions.fused_qk_concat import fused_qk_cat, fused_qk_rope_cat
from aiter.ops.triton.utils.types import str_to_torch_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
positive_int,
pow2_int,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.triton_tests.fusions.test_fused_qk_concat import (
generate_qk_inputs,
generate_rope_cached_freqs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.rope.fused_qk_norm_rope_gate_fp8_quant import (
FP8_DTYPE,
MAX_SEQUENCES,
fused_qk_norm_rope_gate_fp8_quant,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import argparse
import functools
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.fusions.fused_kv_cache import fused_qk_rope_reshape_and_cache
from aiter.ops.triton.utils._triton import arch_info
from aiter.ops.triton.utils.types import e4m3_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import argparse
import functools
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.fusions.fused_reduce_qk_norm_rope_swa_write import (
fused_reduce_qk_norm_rope_swa_write,
)
from aiter.ops.triton.utils.types import str_to_torch_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.triton_tests.fusions.test_fused_reduce_qk_norm_rope_swa_write import (
_build_cos_sin,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.normalization.fused_rmsnorm_add import fused_rmsnorm_add
from aiter.ops.triton.utils.types import str_to_torch_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.triton_tests.normalization.test_fused_rmsnorm_add import (
generate_fused_rmsnorm_add_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.normalization.fused_rmsnorm_add_rmsnorm import (
fused_rmsnorm_add_rmsnorm,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext
from op_tests.triton_tests.normalization.test_fused_rmsnorm_add_rmsnorm import (
generate_fused_rmsnorm_add_rmsnorm_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import sys
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.fusions.fused_sigmoid_mul import fused_sigmoid_mul
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)
from op_tests.triton_tests.fusions.test_fused_sigmoid_mul import (
generate_fused_sigmoid_mul_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.basic.gemm_a16w16 import _is_gluon_available, gemm_a16w16
from aiter.ops.triton.gemm.basic.gemm_a16w16_atomic import gemm_a16w16_atomic
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w16 import (
generate_gemm_a16w16_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.basic.gemm_a16w16_gated import gemm_a16w16_gated
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w16_gated import (
generate_gemm_a16w16_gated_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import sys
import matplotlib.pyplot as plt
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.basic.gemm_a16w16_gated import gemm_a16w16_gated
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_shape_benchmark_object,
model_benchmark_shapes,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w16_gated import (
generate_gemm_a16w16_gated_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter import logger
from aiter.ops.triton.gemm.basic.gemm_a16w4 import gemm_a16w4
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w4 import (
generate_gemm_a16w4_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.basic.gemm_a16w8_blockscale import (
gemm_a16w8_blockscale,
gemm_a16w8_blockscale_preshuffle,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a16w8_blockscale import (
generate_gemm_a16w8_blockscale_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import sys
from collections.abc import Callable
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.basic.gemm_a8w8 import gemm_a8w8 as triton_gemm_a8w8
from aiter.ops.triton.gemm.basic.gemm_a8w8 import gemm_a8w8_preshuffle
from aiter.ops.triton.utils.types import str_to_torch_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8w8 import (
generate_gemm_a8w8_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
from collections.abc import Callable
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.basic.gemm_a8w8_blockscale import (
gemm_a8w8_blockscale as triton_gemm_a8w8_blockscale,
)
from aiter.ops.triton.gemm.basic.gemm_a8w8_blockscale import (
gemm_a8w8_blockscale_preshuffle as triton_gemm_a8w8_blockscale_preshuffle,
)
from aiter.test_common import checkAllclose
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8w8_blockscale import (
generate_gemm_a8w8_blockscale_inputs,
run_torch,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.basic.gemm_a8w8_per_token_scale import (
gemm_a8w8_per_token_scale,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8w8_per_token_scale import (
generate_gemm_a8w8_per_token_scale_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import sys
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.basic.gemm_a8wfp4 import (
gemm_a8wfp4,
)
from aiter.ops.triton.utils.types import get_fp8_dtypes
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_a8wfp4 import (
generate_gemm_a8wfp4_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.gemm.basic.gemm_afp4wfp4 import (
gemm_afp4wfp4 as triton_gemm_afp4wfp4,
)
from aiter.ops.triton.gemm.basic.gemm_afp4wfp4 import (
gemm_afp4wfp4_preshuffle,
)
from aiter.ops.triton.utils._triton import arch_info
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)
from op_tests.triton_tests.gemm.basic.test_gemm_afp4wfp4 import (
generate_gemm_afp4wfp4_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import csv
import functools
import math
import warnings
from pathlib import Path
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.shuffle import shuffle_weight
from aiter.ops.triton.gemm.basic.gemm_afp8wfp8 import (
gemm_afp8wfp8,
gemm_afp8wfp8_preshuffle,
)
from aiter.ops.triton.utils._triton import arch_info
from op_tests.op_benchmarks.triton.utils.argparse import (
add_argparse_ff,
get_do_bench_graph_kwargs,
get_ff_args,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_model_benchmark_object,
get_shape_benchmark_object,
print_vgpr,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import argparse
import functools
import logging
# PyTorch
import torch
# Triton
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
# AITER: Triton kernel wrappers
from aiter.ops.triton.gemm.grouped.gmm import (
gmm as triton_gmm,
)
from aiter.ops.triton.gemm.grouped.gmm import (
nptgmm as triton_nptgmm,
)
from aiter.ops.triton.gemm.grouped.gmm import (
ptgmm as triton_ptgmm,
)
# AITER: GMM defaults and utility functions
from aiter.ops.triton.utils.gmm_common import (
DTYPE,
DTYPE_STR,
GROUP_SIZES_DTYPE,
GROUP_SIZES_DTYPE_STR,
NUM_GROUP_SIZES,
RNG_SEED,
SUPPORTED_DTYPES_STR,
SUPPORTED_GROUP_SIZES_DTYPES_STR,
TRANS_LHS,
TRANS_RHS,
dtype_from_str,
gen_gmm_input,
gen_gmm_output,
gen_multiple_group_sizes,
gen_tgmm_bias_grad,
gen_tgmm_input,
gen_tgmm_output,
group_sizes_dtype_from_str,
str_from_dtype,
str_from_group_sizes_dtype,
)
# AITER: shared argparse helper
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
positive_int,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.normalization.norm import (
layer_norm,
layernorm2d_fwd_with_add,
layernorm2d_fwd_with_add_dynamicquant,
layernorm2d_fwd_with_add_smoothquant,
layernorm2d_fwd_with_dynamicquant,
layernorm2d_fwd_with_smoothquant,
)
from aiter.ops.triton.utils.types import str_to_torch_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
print_vgpr,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.attention.mla import mla_decode_fwd, mla_prefill_fwd
from aiter.ops.triton.utils._triton import arch_info
from aiter.ops.triton.utils.types import e4m3_dtype
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_caller_name_no_ext
from op_tests.triton_tests.attention.test_mla import input_helper, shuffle_kv_buffer


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import argparse
import functools
import itertools
import sys
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.attention.mla_decode_rope import (
decode_attention_fwd_grouped_rope,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_configs,
print_vgpr,
)
from op_tests.triton_tests.attention.test_mla_decode_rope import (
input_helper,
ref_preprocess,
)
from op_tests.triton_tests.utils.rotary_embedding import DeepseekScalingRotaryEmbedding


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import math
import torch
import triton
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.normalization.rmsnorm import rms_norm
from aiter.ops.triton.quant.fused_mxfp4_quant import fused_rms_mxfp4_quant
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
override_parser_action,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import (
get_caller_name_no_ext,
get_model_configs,
print_vgpr,
)
from op_tests.triton_tests.normalization.test_rmsnorm import (
generate_rmsnorm_inputs,
)


⚠️ [ruff] <I001> reported by reviewdog 🐶
Import block is un-sorted or un-formatted

import functools
import torch
from op_tests.op_benchmarks.triton.utils.rotating_buffer_bench import (
RotatingBufferBench,
)
from aiter.ops.triton.rope.rope import (
RotateStyle,
rope_bwd,
rope_cached_bwd,
rope_cached_fwd,
rope_cached_fwd_inplace,
rope_cached_positions_bwd,
rope_cached_positions_fwd,
rope_cached_positions_fwd_inplace,
rope_cached_positions_offsets_bwd,
rope_cached_positions_offsets_fwd,
rope_cached_positions_offsets_fwd_inplace,
rope_cached_thd_positions_2c_bwd,
rope_cached_thd_positions_2c_fwd,
rope_cached_thd_positions_2c_fwd_inplace,
rope_cached_thd_positions_offsets_2c_bwd,
# rope_fwd_2d,
# rope_fwd_2d_inplace,
rope_cached_thd_positions_offsets_2c_fwd,
rope_cached_thd_positions_offsets_2c_fwd_inplace,
rope_fwd,
rope_fwd_inplace,
rope_thd_bwd,
rope_thd_fwd,
rope_thd_fwd_inplace,
)
from op_tests.op_benchmarks.triton.utils.argparse import (
get_do_bench_graph_kwargs,
get_parser,
get_rot_bufs_kwargs,
)
from op_tests.op_benchmarks.triton.utils.benchmark_utils import get_model_configs
from op_tests.triton_tests.rope.test_rope import generate_rope_inputs

@rahulbatra85
rahulbatra85 force-pushed the triton_gen_rot_buf_updates branch 2 times, most recently from b29eac5 to c98f944 Compare November 2, 2026 18:10
@rahulbatra85 rahulbatra85 changed the title [Triton/Gluon] : Add exclude list arg and allow list/dict/tuple of tensors in generate_rotating_buffers_pool [Triton/Gluon] : Change RotatingBuffersBench API and update any existing benchmarks Nov 2, 2026
@rahulbatra85
rahulbatra85 force-pushed the triton_gen_rot_buf_updates branch from c98f944 to fc148b7 Compare November 3, 2026 04:39

@Boss2002n Boss2002n left a comment •

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.

The new RotatingBufferBench class is a good cleanup.

Must fix before merge
Two benchmarks no longer run (benchmark_utils.py:378-380): bench_all_reduce.py and bench_all_reduce_rmsnorm_quant.py still import the deleted helpers and fail with ImportError. Fix: move them to RotatingBufferBench.
The default pool is one input set (rotating_buffer_bench.py:82-83): min_pool_len and max_pool_len default to 1, so benchmarks built without the CLI options stop rotating. Fix: default to 4 and 1000.
Settings leak between shapes (rotating_buffer_bench.py:175-184): per-call values stay on the object, so a later shape can get 64 entries (32 GiB) and pass the memory cap. Fix: keep them local.

Nice to fix
Restore del kwargs, v after sizing (rotating_buffer_bench.py:225-232).
Reject unknown keys in from_dict() (rotating_buffer_bench.py:117-126).
Update the --graph-* help, which names the deleted do_bench_aiter_triton (argparse.py:99-122).
Use logger instead of print, shorten the long docstrings, write gfx950 instead of MI355X, and keep main's names (_out_val_keys).

Merge order
Merge this PR first (nine PRs carry its API commit), after a rebase for the conflict with #6727.

	-Split out new rotating buffer pool generation and
do_bench_aiter_Triton into separate file and a class
	-Add exclude list arg and allow list/dict/tuple of tensors in generate_rotating_buffers_pool
@rahulbatra85
rahulbatra85 force-pushed the triton_gen_rot_buf_updates branch 2 times, most recently from 3c14021 to 719dfbb Compare November 3, 2026 19:38
@rahulbatra85
rahulbatra85 force-pushed the triton_gen_rot_buf_updates branch from 719dfbb to 88f4014 Compare November 3, 2026 19:54

@Boss2002n Boss2002n left a comment

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.

LGTM

@rahulbatra85
rahulbatra85 merged commit 076ba7f into main Nov 3, 2026
55 checks passed
@rahulbatra85
rahulbatra85 deleted the triton_gen_rot_buf_updates branch November 3, 2026 23:19
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants