Conversation
* feat(opus): enable gfx1250 support for OPUS tests
Add gfx1250 (GFX12) support across OPUS header and device tests:
- opus.hpp: Add __GFX12__ guard for bf16 conversion, buffer_default_config,
async_load fallback (no vmem-to-lds-load-insts), and gfx12 waitcnt
(s_wait_loadcnt/s_wait_dscnt/s_wait_expcnt)
- opus.hpp: Add gfx1250 fp4 conversion using pk8 builtins
(cvt_scalef32_pk8_fp4_f32 / cvt_scale_pk8_f32_fp4) which convert
8 fp4 values at once, unlike gfx950's pk=2 style
- test_opus_device.py: Add gfx1250 to FP8/FP4 supported arch sets,
fix fp8/bf8 dtype detection for gfx1250 (e4m3fn/e5m2 like gfx950)
- test_mfma_*.cu: Move #endif guard after template instantiations to
fix compilation on non-gfx942/gfx950 targets
Verified all tests pass on both gfx1250 and gfx942.
* feat(opus): replace inline asm with LLVM IR intrinsics, add arch guards for mfma/wmma
- Replace gfx1250 waitcnt inline asm with __asm("llvm.amdgcn.s.wait.*") LLVM IR
intrinsic bindings, exposing native s_wait_* instructions with number<cnt> API
- Guard struct mfma, DISPATCH_MFMA_ macros, mfma type aliases, mfma_adaptor,
and make_mfma with #if defined(__GFX9__) || !defined(__HIP_DEVICE_COMPILE__)
- Guard struct wmma, DISPATCH_WMMA_ macros, wmma type aliases, wmma_adaptor,
and make_wmma with #if defined(__gfx1250__) || !defined(__HIP_DEVICE_COMPILE__)
- make_tiled_mma auto-selects wmma_adaptor/make_wmma on gfx1250 vs
mfma_adaptor/make_mfma on GFX9 via conditional default template params
- Add gfx1250 async_load via global_load_async_to_lds builtins with compact
GPTR_/LPTR_ macros
- Add wmma f16/f32/f8 device tests for gfx1250
* style: format test_opus_device.py with black
…and improve throughput (ROCm#2414) Three optimizations for the MXFP4 blockscale sort kernel used in MoE inference: 1. Remove unnecessary tl.constexpr annotations (token_num, M_i, strides) - Only BLOCK_SIZE_M, BLOCK_SIZE_N, TOPK remain as constexpr - Reduces compiled kernel variants from O(token_num * stride) to 2 (TOPK=1 and TOPK=k) - Strides annotated with tl.int64 for type stability without recompilation 2. Reduce sorted_ids loads from 4x to 2x per program - Restructure loop: outer m_idx loads sorted_ids once, inner n_idx reuses row addresses - 5-10% speedup for decode (128-1024 tokens) 3. Add fused-N kernel variant for large token counts (>2048) - 1D grid: each program handles ALL N-tiles, loading sorted_ids once - Dispatch threshold at token_num=2048: 2D-grid below, fused-N above - 30-52% speedup for prefill (4096-32768 tokens) 4. Remove unused M_i parameter from _fused_dynamic_mxfp4_quant_moe_sort_kernel and dispatch smaller BLOCK_SIZE_Mx=32 for token_num<=32 to reduce register pressure in the quantization phase (2-7% decode improvement) Benchmark (DeepSeek-R1 config: E=256, topk=8, dim=7168, stage1): token=1: 4.64 -> 4.92 us (within noise) token=128: 6.78 -> 6.43 us (+5%) token=1024: 10.18 -> 9.57 us (+6%) token=4096: 23.77 -> 16.50 us (+31%) token=32768: 197 -> 94.6 us (+52%)
Signed-off-by: ganyi <ygan@amd.com>
* refactor(pa): use ctypes binding for pa_fwd and pa_ps_fwd * update * use const HipDeviceGuard device_guard(Q->device_id); --------- Co-authored-by: amd-ruitang3 <rui.tang2@amd.com> Co-authored-by: amd-ruitang3 <145657428+amd-ruitang3@users.noreply.github.com>
* refactor asm kl bind * update * refactor_topk_softmax_asm_bind * fix lint * Update rocm_ops.hpp --------- Co-authored-by: amd-ruitang3 <Rui.Tang2@amd.com>
* [FEAT] exclude torch.h in a8w8 cu files * format * format * [FEAT] exclude torch.h in a4w4, blockscale_bpreshuffle, mi350 cu files Extend the ctypes FFI refactoring from asm_gemm_a8w8 to the remaining 3 ASM GEMM kernels: - asm_gemm_a4w4.cu - asm_a8w8_blockscale_bpreshuffle.cu - asm_mi350_a8w8_blockscale.cu Changes per kernel: - .cu: Replace torch/pybind11 deps with aiter_hip_common.h + extern "C" - .py: Switch to ffi_type="ctypes" with wrapper functions for type conversion - .json: Remove pybind .cu from srcs (no longer needed with ctypes FFI) Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * update * remove useless .h * use HipDeviceGuard --------- Co-authored-by: junxiaguo <junxiaguo@amd.com> Co-authored-by: Chuanbo Wang <Chuanbo.Wang@amd.com> Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: amd-ruitang3 <rui.tang2@amd.com>
* replace ck_tile type covert by opus cast * [OPUS] Add finfo class for float-valued type properties (eps/max/min/tiny/bits) Supports fp32, fp16, bf16, fp8, bf8, fp4, e8m0 with gfx950/gfx942 specializations. Verified bitwise against torch.finfo on both MI355 (gfx950) and MI308 (gfx942). * [OPUS] Use explicit opus:: namespace in test_finfo.cu * update cache_kernels, quant_kernels, rmsnorm_quant_kernels * fix for fp4 * remove quant_common.cuh --------- Co-authored-by: carlushuang <carlus.huang@amd.com> Co-authored-by: chenjun <junchen2@amd.com>
Made-with: Cursor
* update kernels & interface * rebase kernel to latest * update * update * update
…fix nhead=128 reduce mgc (ROCm#2319) * fix mla nps mode nhead=128 split and reduce error * delete unuseless code * fix num_kv_splits = 0 error
* Migrate MoE ASM kernels from pybind to C ABI + ctypes Convert fmoe, fmoe_int8_g1u0, fmoe_g1u1, fmoe_g1u1_tkw1, fmoe_int8_g1u0_a16, fmoe_g1u1_a16, fmoe_fp8_blockscale_g1u1, and moe_stage1_g1u1 from torch::Tensor& (pybind11) to AiterTensor* + hipStream_t (C ABI called via ctypes). - asm_fmoe.cu: Remove torch/ATen includes, use AiterTensor*, AITER_DTYPE_*, AITER_CHECK; template <int I_elemSize, int O_elemSize> - asm_moe_2stage.cu: Same conversion for moe_stage1_g1u1 - moe_op.h: Remove fmoe pybind declarations (now extern "C") - rocm_ops.hpp: Remove fmoe entries from MOE_OP_PYBIND macro - moe_op.py: Use ffi_type="ctypes" with new module_moe_fmoe_asm - optCompilerConfig.json: Split ctypes sources into module_moe_fmoe_asm Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * fix: add missing #include <memory> for std::unique_ptr/make_unique After pip install -e . refreshed aiter_meta from csrc/, the indirect include chain changed and <memory> was no longer transitively included. * fix: address Copilot review comments - kernelName: str -> Optional[str] for correct ctypes c_char_p conversion - Remove extra activation param from fmoe_int8_g1u0_a16 (C ABI has no such param) - Fix typo "supput" -> "support" in asm_fmoe.cu * fix: add HipDeviceGuard to all C ABI MoE kernel functions Address review comment from amd-ruitang3: the pybind->ctypes migration removed device_guard. Now that PR ROCm#2377 has merged, use the new HipDeviceGuard in all 8 extern "C" fmoe/moe_stage functions. * fix: restore activation parameter in fmoe_int8_g1u0_a16 C ABI The activation parameter was dropped during pybind-to-ctypes migration. The original implementation uses it to select between silu/gelu config maps. Restore it in C ABI signature, Python ctypes declaration, and call site. * fix: remove stale topk_softmax_asm pybind declaration from moe_op.h This declaration was erroneously re-added during rebase conflict resolution. PR ROCm#2327 already removed it when migrating to ctypes. --------- Co-authored-by: root <root@hjbog-srdc-39.amd.com> Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
ROCm#2274) * Add Attention support to bench_models.py * Add MHA layout CLI arg * Add support for batched_gemm_a16wfp4 * Refactor TP logic and _get_handler * Remove unified attention from this branch
* update decode_update_mla_metadata_v1 natively_supported logic * edit get_mla_metadata_v1_2_device params.num_heads = num_heads;
--------- Co-authored-by: lalala-sh <Jiaxing.Wen@amd.com> Co-authored-by: Felix Li <felix.li@amd.com> Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
* mdf_asm_kl_bind * mdf topk * split “module_asm_communication” out from "module_custom_all_reduce" * update * update * update * remove mdf of communication * update * update
* [setup][flydsl] Add flydsl into aiter install requires flydsl was only declared in requirements.txt and pyproject.toml build-system requires, but missing from setup.py install_requires. This caused flydsl to be absent in Docker images because `pip install -e . --no-build-isolation` skips build deps and only installs install_requires Signed-off-by: zejunchen-zejun <zejun.chen@amd.com> * force flydsl version >=0.1.1 Signed-off-by: zejunchen-zejun <zejun.chen@amd.com> * force flydsl version according to requirements.txt Signed-off-by: zejunchen-zejun <zejun.chen@amd.com> --------- Signed-off-by: zejunchen-zejun <zejun.chen@amd.com>
* Improve config selection for RDNA gpus * Add comments for config flow logic
* Fix `test_moe_routing_sigmoid_top1_fused.py` The kernel's tuning config JSON file contained trailing commas after the last entry, which is invalid JSON. * Fix `test_gather_kv_b_proj.py` Use `gfx950` FP8 type instead of `gfx942` one (`torch.float8_e4m3fn` instead of `torch.float8_e4m3fnuz`). * Fix `test_causal_conv1d.py` Relax `bf16` absolute tolerance from 5e-2 to 6e-2 for marginal precision differences on `gfx950`.
* [HIP] Optimized fused split GDR decode
…OCm#2436) Co-authored-by: RoyWang <RoyWang@amd.com>
…Cm#2444) pull_request events from forks only get read-only GITHUB_TOKEN, causing createComment to fail with permission errors. pull_request_target runs in the base repo context with write access. This is safe since the workflow only posts a comment and never checks out or executes PR code.
…OCm#2448) Update recommended tool versions to match current CI: - black: 25.1.0 -> 26.3.0 - ruff: 0.11.11 -> 0.15.7 - clang-format: require version 18/19/20 Update pre-commit hook to prefer versioned clang-format (20 > 19 > 18) with fallback to unversioned clang-format. Signed-off-by: Xin Huang <Xin.Huang@amd.com>
* CI: Switch Triton tests from MI325 to MI355 runners * Update triton-test.yaml * CI: Add ci:triton-355 label to opt into MI355 Triton runners on PRs * Revert "CI: Add ci:triton-355 label to opt into MI355 Triton runners on PRs" This reverts commit 99b42f8. * CI: Default Triton to MI325; opt into MI355 via ci:triton-355 label * Fix: Restore MI325 runner after merge with main The merge of main (which included ROCm#2380) left runs-on referencing matrix.runner, but the matrix was simplified to shard-only. This caused MI325 shard jobs to silently not schedule. * Fix: Add 'labeled' event type to trigger MI355 Triton tests The triton-mi355 job checks for the ci:triton-355 label, but the workflow only triggers on opened/synchronize/reopened/ready_for_review. Adding a label after PR creation does not re-trigger the workflow, so the MI355 jobs are always skipped. Add 'labeled' to pull_request types so adding ci:triton-355 triggers the workflow and the MI355 jobs can run. * Fix: MI355 Triton job reuses pre-built wheel instead of compiling from source Add build-triton dependency, download triton-wheel artifact, and pass TRITON_WHEEL_DIR to build_aiter_triton.sh — matching the MI325 job behavior. Previously each MI355 shard compiled Triton from source, wasting ~15 min per shard.
* fix * add e=256 k=8 tuned config * updalte flydsl --------- Co-authored-by: coderfeli <felix.li@amd.com>
* Improve RDNA config selection for FA * Formatting --------- Co-authored-by: Saeid Rostami <srostami@amd.com>
…eam (ROCm#2564) * fix(hip): launch FMHA Philox, sampling, and MM kernels on current stream - FMHA: ParsePhiloxCudaState hipLaunchKernelGGL now uses the same stream as mha_fwd/mha_bwd (at::hip::getCurrentHIPStream) in asm and CK py_itfs paths. - sample_kernels: use <<<grid, block, 0, stream>>> for greedy/random/mixed/ exponential sampling ops that previously omitted the stream argument. - custom_kernels: MMGPUKernel passes stream into matrixMultiplyShared launch. - moe_cktile2stages: remove no-op getCurrentHIPStream() calls. Related to stream-ordering issues like ROCm#2520 when callers use non-default streams (e.g. overlap scheduling). Made-with: Cursor * Resolve Copilit review comment --------- Co-authored-by: wunhuang <wunhuang@amd.com>
* Fix nondeterministic RNG in test_fused_mxfp4_quant Tests in test_fused_mxfp4_quant.py were failing in CI, especially when executed as part of shard 3. The failures were not reproducible when running the test line in isolation. Thanks to Bruno for providing the command line. Root cause: The random seed was previously set to be at the top-level part of the module just after imports via torch.manual_seed(). This caused test behaviour to depend on the global RNG state, which is affected by previously executed tests in the same shard (which makes sense why it worked in isolation, but not in the shard). As a result, the test outcomes were order-dependent and non-deterministic. Fix: - Removed torch.manual_seed() from top-level part of module - Added this deterministic seeding behaviour to the test case that was being impacted by this to ensure order-independent behaviour Validation: - Reproduced failure using CI shard 3 command locally - Verified the failures occuring in op_tests/triton_tests/quant/test_fused_mxfp4_quant.py::test_fused_rms_quant - After fix: - All tests pass in shard 3 with TRITON_HIP_USE_ASYNC_COPY=0 - Test_fused_rms_quant also passes with ASYNC_COPY enabled (in command line run with shard 3 and isolation) - Tests pass consistently in isolation and repeated runs Additional Notes: - Remaining failures with TRITON_HIP_USE_ASYNC_COPY=1 are affected (MoE + GEMM known issues with ASYNC enabled). This is unrelated to the current task and can be addressed separately * Set RNG seed before all test cases to make everything deterministic. Moved the seeds after skip condition, and used black to format file
…device && rm hip_compat.h (ROCm#2525) * add WARP_SIZE define for host and device && rm hip_compat.h * update * update reduce * update quant kernels * add permlanex16 to hip_reduce * rm WARP_SIZE define in topk_softmax_kernels_group.cu
The 1-stage fused allreduce+RMSNorm kernel produces numerically different residual outputs compared to the unfused (allreduce -> bf16 -> residual add) path. The divergence is small per element (1-4 ULPs in bf16) but compounds across transformer layers during decode, causing measurable accuracy regression (e.g. -2.6pp on GSM8K for a 60-layer MoE model at TP=4). Root cause: the 1-stage kernel accumulates in f32 and adds the residual before downcasting to bf16, skipping the intermediate bf16 rounding that the unfused path naturally performs. This extra f32 precision shifts ~25% of output elements by 1+ ULPs. Fix: insert a register-level bf16 round-trip (downcast+upcast) after the f32 allreduce accumulation and before the residual addition, so the fused kernel matches the unfused path bit-for-bit. No memory traffic added; no measurable impact on kernel latency. Made-with: Cursor
Co-authored-by: junxiaguo <junxiaguo@amd.com>
* OPUS: add gfx950 smem transpose load path Add smem tr_load/tr_load_if APIs and wire _tr_load to gfx950 ds_read_tr* builtins with scalar/vec dispatch, including clang>=20 u16 support and simplified diagnostics. * tr_load example layout and unit test
* replace ck_tile api with opus api in some hip kernels(topk_softmax, moe_fused_gate. sample) * update * rm ck_tile in topk_softmax_kernels_group.cu --------- Co-authored-by: Xin Huang <Xin.Huang@amd.com>
…m#2555) * Fix some benchmark scripts so that they generate the output CSVs Affects the following Triton-based benchmarks: * bench_moe_gemm_a4w4.py * bench_moe_gemm_a8w4.py * bench_moe_gemm_a8w8.py * bench_moe_gemm_a8w8_blockscale.py * bench_moe_gemm_int8_smoothquant.py * Reformat some MoE GEMM benchmarks with Black * Change comments to proper type annotations
* adding sliding window for sink attn * format * split sink attention tests and guard fused backward * prune sliding window blocks in Triton kernels Skip tiles that cannot overlap the active sliding window in the forward and one-kernel backward paths so local attention avoids paying full-context compute on long sequences. Made-with: Cursor * clarify sink test skip comments Make the sink-specific skips explicitly reference the existing baseline MHA backward limitations so they are not mistaken for sink-only issues. Made-with: Cursor
Reduce UTs by removing unnecessary tests. This should be a reduction of ~88%. Mainly done by Reduce number of shapes and keep the relevant ones Have another set of smaller shapes to use for different layouts, output tensor arg and float16.
…ld (ROCm#2548) (ROCm#2603) * fix: split asm_topksoftmax into separate module to fix ctypes JIT build When topk_softmax_asm (ffi_type="ctypes") triggers JIT compilation, torch_exclude=True is forced. But module_moe_asm also contains pybind .cu files that depend on torch, causing undefined symbol errors. Split asm_topksoftmax.cu into its own module_moe_topksoftmax_asm so it compiles torch-free independently. No .cu files modified. Fixes ROCm#2548 * style: black formatting for compile_ops decorator line * fix: remove unnecessary ck_tile include from module_moe_topksoftmax_asm asm_topksoftmax.cu does not use ck_tile headers. --------- Co-authored-by: root <root@hjbog-srdc-39.amd.com>
Rope had 138757 UTs. We don't need that many. Reduced to 1987.
* replace ck_tile api by opus in activation * change warp size 64 to WARP_SIZE * fix include * fix warp size and rm check * fix format
* add fused_qk_norm_group_quant kernel * Optimize fused_qk_rmsnorm_group_quant kernel and add fp4x2 test support Kernel optimizations (fused_qk_rmsnorm_group_quant.cu): - Add row_active guard to skip OOB threads in load/store/compute paths - Fuse x2 (K) processing into same block as x1 when grid_y==1, halving block count for large token sizes with second input - Adaptive grid_y strategy: grid_y=2 for small tokens (m<=1024) with n2>0 to leverage CU parallelism; grid_y=1 for large tokens to reduce launch overhead - Architecture-specific dispatch for gfx950 fp8 in <=2048 bucket with multiple BlockSize/thread_data_size configs (128x16, 64x32, 256x8, 128x8, 256x16, 64x16, 128x32) and env var override (AITER_FUSED_QK_RMS_2048_CFG) Test improvements (test_fused_qk_rmsnorm_group_quant_hip.py): - Add fp4x2 quantization support: reference impl, Triton mxfp4 baseline, HIP fp4x2 path with e8m0 scale handling - Auto-detect fp4x2 capability on gfx950/gfx1250 and include in default test matrix - Expand default test matrix: add token=16384, residual=[0,1] - Add --quant_out_dtype CLI arg, remove redundant --quant_type - Add gfx and quant_type columns to summary output Made-with: Cursor * opt perf * rename test and fix lint * opt multithread_reduce * update review comments
* CK mha bwd: add sink attention score gradient support * test: add varlen sink bwd tests to test_mha_sink_bwd * Update op_tests/test_mha_sink_bwd.py Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * style: apply black formatting to test_mha_sink_bwd * test: move sink bwd tests into test_mha.py and test_mha_varlen.py * style: apply black formatting to sink bwd tests in test_mha and test_mha_varlen * fix: adapt mha bwd to updated CK fmha_bwd API and zero dq_accum Three fixes required after the CK submodule was updated to the sink_bwd_cherry_pick branch: 1. fmha_bwd_traits no longer carries seqlen/batch/nhead fields. Remove the now-stale seqlen_q, seqlen_k, batch, max_seqlen_*, nhead_q, nhead_k arguments from the traits initializer lists in mha_bwd.cu, mha_bwd_kernels.cu, and mha_varlen_bwd_kernels.cu. 2. nhead_stride_dq_acc / batch_stride_dq_acc are int64_t in mha_bwd_args but ck_tile::index_t (int) in fmha_bwd_args. Add explicit static_cast<ck_tile::index_t> to silence the narrowing-conversion errors. 3. fmha_bwd_launcher was removed from the new CK API. Replace launcher.dq_acc_splits with the equivalent expression ceil(seqlen_k / 16) for deterministic mode and 1 otherwise, matching the logic documented in fmha_bwd_runner.hpp. Replace launcher.needs_zero_dq_acc with unconditional torch::zeros: the dq_dk_dv kernel always writes dq_acc via atomicAdd (even in non-deterministic mode), so an uninitialized accumulator silently corrupts dQ for hdim >= 128 where the convert_dq kernel is active. All 22 sink-bwd tests pass after this change. * update ck to ROCm/rocm-libraries#5504 * Revert "update ck to ROCm/rocm-libraries#5504" This reverts commit 7481fd6. * update ck commit Signed-off-by: Linjun-AMD <Jun.Lin@amd.com> * update bwd args Signed-off-by: Linjun-AMD <Jun.Lin@amd.com> * [CK] update mha bwd traits args and fix sink_ptr comments * [CK] fix mha_bwd_args initializer in benchmark_mha_bwd.cpp for sink_ptr/d_sink_ptr --------- Signed-off-by: Linjun-AMD <Jun.Lin@amd.com> Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
* Fix group topk dispatch for glm5 * update grouped_topk not compute topk group when group=1
…topk, cache, sample) (ROCm#2599) * update topk_softmax * update hip group topk * rm warpsize in sample_kernels.cu * update cache.cu * update * update2
* mi350 mla ps mode support nhead8 mtp4 * upload lse co * add return lse test * fix kPackedQoLenPerWg = 16 only when (num_heads == 8) && (max_seqlen_qo == 4) && q_is_fp8 && kv_is_fp8) * fix the err * up the perf * uplift perf to 545 TFLOPS * rename the kernel name
* [FlyDSL] Upgrade MOE kernels: split-K, fuse quant, flyc.compile fast dispatch - Refactor stage1/stage2 to use flyc.compile for fast kernel dispatch (~5us) - Add split-K support with fused silu_and_mul + mxfp4 quant + scale-sort - Add persistent round-robin mode for stage2 - L2 cache optimization for preshuffle pipeline - Update tuner for new kernel configurations Made-with: Cursor * code clean * fix typo * update config * update * code clean * atach link * clean dead code * format * moe ut collect ds as default * Update kimik2_fp4_tuned_fmoe.csv * Delete aiter/configs/model_configs/kimi-2_fp4_tuned_fmoe.csv drop duplicated configs
Add FlyDSL-based implementation of fused_qk_rope_reshape_and_cache as a drop-in replacement for the Triton version. The FlyDSL kernel fuses NeoX-style RoPE rotation with paged KV cache writes in two kernel launches (Q RoPE + K RoPE + KV cache write). Performance: 1.50x average speedup over Triton across 12 LLM models on MI355X. Accuracy: bit-identical to Triton for bf16, within 4e-3 for f16. E2E validation: Llama-3.1-8B-Instruct produces identical outputs with both backends. Changes: - aiter/ops/flydsl/rope_kernels.py: wrapper matching fused_qk_rope_reshape_and_cache interface, with automatic Triton fallback for unsupported features (GPT-J style, offsets, KV scaling, zeros output) - aiter/ops/flydsl/__init__.py: export flydsl_fused_qk_rope_reshape_and_cache - op_tests/flydsl_tests/test_flydsl_rope.py: unit tests covering int32/int64 positions, flash/non-flash layouts, bf16/f16, head_dim=64/128, multi-model sweep Requires FlyDSL >= 0.1.1 with fused_rope_cache_kernel supporting pos_int64=True. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
![github-actions[bot]](https://avatars.githubusercontent.com/in/15368?s=60&v=4){kind=small}
| if _FLYDSL_ROOT not in sys.path: | ||
| sys.path.insert(0, _FLYDSL_ROOT) | ||
|
|
||
| from kernels.fused_rope_cache_kernel import build_fused_rope_cache_module |
Comment on lines
+20
to
+21
| import time | ||
|
|
Owner
Author
|
|
Owner
Author
E2E GSM8K Evaluation CompleteRan GSM8K 3-shot evaluation on Llama-3.1-8B-Instruct with vLLM 0.18 on MI355X:
Result: Identical. Both backends produce the same correct/incorrect distribution on all 100 GSM8K questions. Note: The low raw accuracy is due to 3-shot prompting without chat template — Llama-3.1-8B-Instruct expects the chat format for proper reasoning. The key metric is that both backends match exactly, confirming FlyDSL RoPE is a drop-in replacement with no accuracy impact. Full Validation Summary
PR is ready for review. |
Owner
Author
E2E GSM8K v2 — Proper Evaluation ResultsFixed evaluation pipeline (answer extraction + stop tokens). Re-ran GSM8K 3-shot on Llama-3.1-8B-Instruct with vLLM 0.18 on MI355X:
Both backends well above the ATOM CI threshold. The 2pp difference is within normal variance for 100-sample evaluation. This confirms FlyDSL RoPE is a safe drop-in replacement with no model accuracy impact. |
Owner
Author
|
Closed: branch base was stale. Recreating with clean base. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
fused_qk_rope_reshape_and_cache(Triton)Changes
aiter/ops/flydsl/rope_kernels.py— wrapper with Triton fallback for unsupported featuresaiter/ops/flydsl/__init__.py— export new functionop_tests/flydsl_tests/test_flydsl_rope.py— comprehensive unit testsTest plan
🤖 Generated with Claude Code