Repository navigation
Benchmarks: Micro benchmark - Add Blackwell and Rubin CUTLASS kernels to gemm-flops - #883
Draft
Hongtao Zhang (polarG) wants to merge 27 commits into
Draft
Hongtao Zhang (polarG) wants to merge 27 commits into
Hongtao Zhang (polarG) wants to merge 27 commits into
Conversation
… comment Add four SM100 UMMA kernel name patterns to CUTLASS_LIBRARY_KERNELS so that the cutlass_profiler binary compiled for CUDA >= 12.9 (Blackwell GB200/GB300) actually contains SM100 kernels: cutlass3x_sm100_tensorop_gemm_tf32_* cutlass3x_sm100_tensorop_gemm_f16_* cutlass3x_sm100_tensorop_gemm_bf16_* cutlass3x_sm100_tensorop_gemm_s8_* Previously, all ten patterns in the filter were SM80-era legacy names (cutlass_simt_* / cutlass_tensorop_*). CUTLASS v4.1.0 generates SM100 kernels under the 3x naming scheme (cutlass3x_sm100_tensorop_gemm_*), so none of the old patterns matched, resulting in a profiler binary with zero SM100 kernels. When the SuperBench gemm-flops benchmark invoked the profiler on a GB300 node the profiler found no matching kernels and hung until the 1500 s timeout was reached (return code 124). The legacy SM80 patterns are preserved and are silently ignored when building for SM100-only targets; the new SM100 patterns are silently ignored when building for SM80-only targets, so the change is backward-compatible. Also fix the version comment which incorrectly referenced v3.9 instead of the v4.1.0 tag that is actually cloned.
…rsing for Blackwell
The gemm-flops benchmark on Blackwell (SM100/SM103, GB200/GB300) was timing
out after 1500 s (return code 124) because:
1. __kernel_map entries for compute capability 10.0 and 10.3 were copied from
the SM80 entry. SM80 used legacy SIMT and TensorOp names
(cutlass_simt_dgemm_*, cutlass_tensorop_h884gemm_*, etc.). CUTLASS v4.1.0
does NOT generate SIMT kernels for SM100; it generates 3x UMMA kernels
under a completely different naming scheme. Passing a non-existent kernel
pattern to cutlass_profiler causes it to loop indefinitely.
2. __parse_logline only contained SM80 name substrings, so even if SM100 output
had been produced it would not have been parsed.
Fix:
- Replace __kernel_map[10.0] and [10.3] with correct SM100 UMMA 3x patterns:
tf32_tc -> cutlass3x_sm100_tensorop_gemm_tf32_*
fp16_tc -> cutlass3x_sm100_tensorop_gemm_f16_*
bf16_tc -> cutlass3x_sm100_tensorop_gemm_bf16_*
int8_tc -> cutlass3x_sm100_tensorop_gemm_s8_*
fp64_tc, fp32, fp16, fp64 SIMT precisions are removed as they are not
supported on SM100 in CUTLASS v4.1.0.
- Add the four SM100 UMMA prefixes to __parse_logline so benchmark output
can be extracted from the cutlass_profiler CSV.
Kernel naming verified against CUTLASS v4.1.0 generator.py:
GenerateSM100() calls only TensorOp_*_UMMA_gemm functions.
3x kernels use prefix='3x', core_name()='gemm', yielding the
cutlass3x_sm{arch}_tensorop_gemm_{dtype}_* template.
…arsing Add test_flops_performance_cuda_sm100() covering all four precisions supported on Blackwell (SM100/SM103) in CUTLASS v4.1.0: - tf32_tc (cutlass3x_sm100_tensorop_gemm_tf32_*) - fp16_tc (cutlass3x_sm100_tensorop_gemm_f16_*) - bf16_tc (cutlass3x_sm100_tensorop_gemm_bf16_*) - int8_tc (cutlass3x_sm100_tensorop_gemm_s8_*) Each test case feeds a representative cutlass_profiler CSV output line (using a realistic 3x UMMA kernel name) through the benchmark's run() path and asserts that the correct GFLOPS/GIOPS metric is extracted. The test mirrors the structure of the existing SM80 test.
…kwell)
During a SuperBench gemm-flops run on a GB300 (SM100/SM103) node, the
cutlass_profiler hung for the full 1500 s timeout and returned exit code 124
(SIGTERM). Investigation revealed two root causes:
1. CUTLASS_LIBRARY_KERNELS in the Makefile only listed SM80-era kernel name
patterns, so zero SM100 kernels were compiled into the profiler binary.
2. The Python __kernel_map for SM10.0/SM10.3 referenced non-existent SIMT
kernel names (fixed in a separate commit).
While cause (2) is fixed in the Python layer, upgrading CUTLASS also brings
in upstream profiler reliability fixes that directly address the hang-on-
no-match behaviour observed in v4.1.0.
**v4.2.0 (2025-09-15):**
- SM103 (B300/GB300) blockscaled FP4 (Ultra) GEMM kernel support — enables
future SuperBench benchmarks targeting the SM103 FP4 path.
- Heuristics-based kernel autotuning via nvidia-matmul-heuristics — allows
the profiler to find the best kernel variant automatically.
- CUTLASS_LIBRARY_INSTANTIATION_LEVEL flag for SM100/103: lets the build
instantiate all possible kernel combinations without listing each pattern
explicitly — more robust against future precision additions.
- Multiple profiler bug fixes: timeout/no-output issues on SM100 resolved.
- SM100 MoE grouped GEMM, mixed-input GEMM, cpasync, and FP4 GEMV kernels.
- SM121 (DGX Spark) GPU support.
**v4.3.0–v4.3.1 (2025-11-21 / 2025-11-26):**
- SM103 support added to the CuTe Python DSL.
- SM100 convolution stream-K kernel.
- Additional profiler hang and no-output fixes for Blackwell SM100.
- Ragged contiguous grouped GEMM for SM100 MoE workloads.
- SM100 FMHA backward + MLA (multi-latent attention) kernels.
**v4.3.2–v4.3.5 (2025-12-05 – 2026-01-09):**
- Bug fixes and stability improvements; copyright updated to 2026.
- Optimal code generation with CUDA toolkit 13.0U1.
- Kernel name patterns (cutlass3x_sm100_tensorop_gemm_*) are unchanged
between v4.1.0 and v4.3.5; no changes to CUTLASS_LIBRARY_KERNELS or the
Python __kernel_map / __parse_logline are needed.
- The CUDA >= 12.9 version threshold is unchanged (v4.3.5 still supports
CUDA 12.9; optimal code generation targets CUDA 13.0U1).
- SM80-era and SM90 build paths are unaffected (still use their existing
CUTLASS versions).
CUTLASS_LIBRARY_KERNELS was a single flat list shared across all docker
image build paths (SM70/80, SM90, SM100/103). This has two issues:
1. Intent is opaque: SM80 legacy names (cutlass_simt_*, cutlass_tensorop_*)
sit alongside SM100 3x UMMA names (cutlass3x_sm100_*) with no indication
of which applies to which GPU family.
2. Every cmake invocation carries all patterns regardless of the target
architecture. Unmatched patterns are silently ignored by CUTLASS, so
correctness is not affected, but the build definition misleads readers
into thinking all kernels are built for all images.
Each build path now declares only the kernel patterns relevant to its
target GPU architecture, mirroring the existing per-arch ARCHS variable.
The cmake call uses $(CUTLASS_KERNELS) instead of the hardcoded list.
CUDA >= 12.9 (SM100/103, v4.3.5):
cutlass3x_sm100_tensorop_gemm_{tf32,f16,bf16,s8}_*
SM103 standard GEMM kernels use SM100 UMMA naming (same cutlass3x_sm100_*
prefix) because SM103 inherits SM100's UMMA ISA.
CUDA >= 12.8 (SM90+SM100, v3.9.2):
cutlass3x_sm90_tensorop_gemm_{tf32,f16,bf16,s8}_* (Hopper WGMMA)
cutlass3x_sm100_tensorop_gemm_{tf32,f16,bf16,s8}_* (Blackwell UMMA)
This also fixes a pre-existing gap: the old flat list had zero SM90 3x
WGMMA patterns, so Hopper kernels were never compiled into the profiler
on CUDA 12.8 images.
CUDA >= 11.8 (SM70–SM90, existing cutlass):
cutlass_simt_{dgemm,sgemm,hgemm}_* (SIMT legacy)
cutlass_tensorop_*gemm_* (SM80 TensorOp)
Unchanged from the previous flat list.
else (SM70–SM86 subset):
cutlass_simt_* and Volta TensorOp (h884) only — no SM80 A100-specific
kernels since those GPUs pre-date A100.
$(eval VAR := ...) inside a Make recipe must be a single unbroken line.
Backslash-newline inside $(eval ...) makes Make insert a space before each
continuation, which corrupts the comma-separated kernel list sent to cmake.
All CUTLASS_KERNELS assignments are therefore kept on one line.
The cmake argument is quoted as "-DCUTLASS_LIBRARY_KERNELS=$(CUTLASS_KERNELS)"
to handle any shell-level interpretation of special characters in the list.
Register fp8_tc as a supported precision in GemmFlopsBenchmark. - Add 'fp8_tc' to _support_precisions list - Add 'fp8_tc': 'fp8_tc_flops' to _metric_map FP8 (e4m3) is the highest-throughput floating-point format natively supported by both Hopper (SM90) and Blackwell (SM100/SM103): H100/H200: ~3.9 PFLOPS FP8 (CUTLASS WGMMA, 3x naming) GB200: ~4.5 PFLOPS FP8 (CUTLASS UMMA, 3x naming) GB300: ~7.0 PFLOPS FP8 (CUTLASS UMMA, 3x naming) Not adding fp8_tc to _support_precisions alone is sufficient since _support_precisions is overridden at runtime by __kernel_map[capability] keys; older GPUs (SM80 and below) will not see fp8_tc because there is no __kernel_map entry for them.
Add fp8_tc (FP8 e4m3) precision to the CUDA gemm-flops benchmark kernel map and CSV output parser for Hopper and Blackwell GPUs. Kernel map additions: - SM90 (H100/H200): 'cutlass3x_sm90_tensorop_gemm_e4m3_*' Uses CUTLASS WGMMA 3x kernels from v3.9.2+. - SM100 (GB200): 'cutlass3x_sm100_tensorop_gemm_e4m3_*' Uses CUTLASS UMMA 3x kernels from v4.3.5. - SM103 (GB300): inherits SM100 map via dict() copy (no change needed). FP8 is NOT added to SM80 or older maps since CUTLASS does not generate e4m3 SIMT/TensorOp kernels for those architectures. The runtime _support_precisions override (line 101: list(__kernel_map[cap].keys())) means fp8_tc is silently excluded on GPUs whose map has no fp8_tc entry, requiring no config-level exclusion by the user. Parse pattern additions in __parse_logline: - 'gemm,cutlass3x_sm90_tensorop_gemm_e4m3_' for SM90 FP8 output rows - 'gemm,cutlass3x_sm100_tensorop_gemm_e4m3_' for SM100 FP8 output rows
Extend per-architecture CUTLASS_KERNELS with e4m3 FP8 patterns so the cutlass_profiler binary contains FP8 benchmarking support. Changes: - CUDA >= 12.9 (SM100/103, CUTLASS v4.3.5): + cutlass3x_sm100_tensorop_gemm_e4m3_* - CUDA >= 12.8 (SM90/100, CUTLASS v3.9.2): + cutlass3x_sm90_tensorop_gemm_e4m3_* (Hopper WGMMA FP8) + cutlass3x_sm100_tensorop_gemm_e4m3_* (Blackwell UMMA FP8) FP8 e4m3 delivers the highest floating-point throughput before FP4: H100/H200 (SM90): ~3.9 PFLOPS GB200 (SM100): ~4.5 PFLOPS GB300 (SM103): ~7.0 PFLOPS Older CUDA paths (< 12.8, targeting SM70-SM86) are unchanged since CUTLASS SIMT/TensorOp kernels for those architectures do not include FP8 support.
Add test_flops_performance_cuda_fp8() to verify that the CSV output parser correctly extracts GFLOPs from FP8 e4m3 kernel result lines for both Hopper (SM90) and Blackwell (SM100) architectures. Test cases cover: - SM90 WGMMA kernel: cutlass3x_sm90_tensorop_gemm_e4m3_e4m3_f32_* Expected max GFLOPs: 2317200 (from two kernel variants) - SM100 UMMA kernel: cutlass3x_sm100_tensorop_gemm_e4m3_e4m3_f32_* Expected max GFLOPs: 5425400 (from two kernel variants) The test reuses the existing _process_raw_result() path and verifies that 'fp8_tc_flops' is correctly populated in benchmark.result.
The per-architecture refactoring (c8f21fe) accidentally dropped SM80 legacy kernel patterns from the CUDA >= 12.8 build path, which would cause the same timeout hang on Hopper H100/H200 that we fixed for Blackwell. Root cause: __kernel_map[9.0] maps non-FP8 precisions (fp64, fp32, fp16, fp64_tc, tf32_tc, bf16_tc, fp16_tc, int8_tc) to SM80-era legacy names (cutlass_simt_*, cutlass_tensorop_*). These kernels are backward-compatible on SM90 Hopper, but the Makefile was only building cutlass3x_sm90_* patterns, so cutlass_profiler would find zero matching kernels for non-FP8 precisions and hang. Changes: - Restored SM80 legacy patterns (SIMT + TensorOp) in CUTLASS_KERNELS - Restored ARCHS '80;90;100' (was '90;100') to compile SM80 kernels - Kept SM90 FP8 (cutlass3x_sm90_tensorop_gemm_e4m3_*) and all SM100 3x UMMA patterns alongside the legacy ones This matches the original main branch ARCHS='80;90;100' while adding the new 3x patterns for FP8 and SM100 support.
The SM100 and FP8 tests used Python name-mangled attribute _CudaGemmFlopsBenchmark__precision_need_to_run which creates a separate attribute from the _precision_need_to_run that _process_raw_result() actually reads. Without calling _preprocess() (which parses args, initializes _result, and sets the real attribute), both tests would fail with IndexError or AttributeError at runtime. Fixes: - Initialize args via add_parser_arguments() + parse_args() so _args.log_raw_data is available for _process_raw_result() - Create BenchmarkResult so _result.add_raw_data() works - Set benchmark._precision_need_to_run (the correct base-class attribute) instead of the name-mangled variant Note: the pre-existing SM80 test (test_flops_performance_cuda) has the same name-mangled assignment at line 56, but it is harmless there because _preprocess() is called first which correctly populates _precision_need_to_run via the --precision CLI argument.
…or to sub-makefile Three bugs in the CUDA>=12.9 CUTLASS build caused generated_kernels.txt to never be created, failing the cmake install step: 1. Bare arch tokens "100;103" do not trigger Blackwell code generation. CUTLASS v4.3.5's generator.py requires suffixed tokens (100a, 103a, etc.) to enter the GenerateSM100() path. Changed to "100a;103a". 2. TF32 kernel filter pattern "gemm_tf32_*" does not match actual kernel procedural names. In CUTLASS v4.3.5, the intermediate type precedes "gemm" in core_name(), producing "tf32gemm" not "gemm_tf32". Changed to "tf32gemm*". 3. CUTLASS gates int8 GEMM generation when 103a is in the arch list (arch_family_cc pruning in GenerateSM100()). SM103 hardware fully supports tcgen05.mma.kind::i8, and generated int8 kernels use CC range [100, 110] via ThorSMRenumbering, so runtime dispatches them on both SM100 and SM103. The gate is a dead-kernel pruning optimization for family-binary builds, not a hardware limitation. A sed patch clears arch_family_cc to restore int8 GEMM in the combined 100a+103a build. Also fixed the same TF32 naming issue in the CUDA>=12.8 block. Refactored the entire cuda_cutlass target into a dedicated cutlass.mk sub-makefile to isolate the growing per-architecture complexity: - Reusable kernel list variables (CUTLASS_KERNELS_LEGACY, _VOLTA, _SM100) eliminate duplicated 200+ char strings across CUDA version tiers. - Each tier now sets 4 variables (TAG, ARCHS, KERNS, PATCHES) and shares a single clone -> patch -> cmake -> build -> install pipeline. - _strip_spaces helper handles Make backslash-continuation space injection. - Adding new GPU architectures or CUDA tiers is now a 4-line addition. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
On ARM64 (Grace CPU), the per-family shared library libcutlass_gemm_sm100_f16_gemm_e4m3.so overflows R_AARCH64_PREL32 relocations in .eh_frame when thousands of FP8 kernel objects are linked into a single SO. The 32-bit relative offset cannot span the large .text section produced by the FP8 e4m3 kernel family. Fix by: 1. Building static libs only (-DCUTLASS_BUILD_SHARED_LIBS=OFF) to avoid the SO link step entirely. SuperBench only uses the cutlass_profiler binary which links statically. 2. Suppressing asynchronous unwind tables in generated kernel TUs (-fno-asynchronous-unwind-tables) to reduce .eh_frame size as a defense-in-depth measure. This is safe because generated kernel code does not throw C++ exceptions; the core library objects (handle.cu, util.cu) that do throw retain synchronous unwind tables which are sufficient for exception propagation. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
- Narrow SM100 kernel filter from 5 broad wildcards (~10K objects) to 10 exact dtype patterns (~3,900 objects) to stay under aarch64 R_AARCH64_PREL32 relocation limit. Remove void-output, mixed-output, and e5m2/e3m2 input variants not needed for benchmarking. - Switch from static-only to shared per-family libraries, eliminating the need for sed alias patch and objcopy workarounds. - Fix TF32 kernel name mismatch: gemm_tf32_* -> tf32gemm_* to match actual CUTLASS v4.3.5 kernel names. - Fix FP8 kernel filter overlap: gemm_e4m3_* -> gemm_e4m3_e4m3_* to prevent matching FP4/FP6 kernels and reporting inflated FLOPS. - Add FP4 (e4m3 x e2m1) precision support to gemm-flops benchmark for Blackwell GB200/GB300. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Narrow both build-time and runtime kernel filters to reduce profiler execution time from ~7 hours to ~58 minutes while reliably measuring peak FLOPS across all 6 precisions on GB200/GB300. Build filter (cutlass.mk - CUTLASS_LIBRARY_KERNELS): - Drop redundant sub-groups that don't contribute to peak measurement: - f16_f16_f16_f16_f16 (FP16 accumulator, same throughput as FP32-accum) - f16_f16_f32_f32_f32 (FP32 output, redundant with FP16-output variant) - bf16_bf16_f32_f32_f32 (FP32 output, same as above for BF16) - s8_s8_s32_s32_s32 (INT32 output, canonical INT8 uses s8 output) - Reduces build from 10 patterns (~3,900 objects) to 6 patterns (~2,200) - Each per-family shared library sub-group stays well under the aarch64 R_AARCH64_PREL32 relocation limit (~1000 objects max) Runtime filter (cuda_gemm_flops_performance.py - --kernels): - Select only 256x256 tiles with stream_k scheduling and 2sm cooperative mode — the competitive subset for peak FLOPS at 16384x16384x16384 - Based on actual GB300 TF32 profiling data: the top 10 kernels are all 256x256 tiles, 2x2x1 cluster, stream_k, 2sm; all 8 layouts within 3% - ~56 kernels per precision (TF32/FP16/BF16/INT8/FP8) + 49 FP4 (all) - Total: ~329 kernels at ~10.5 sec each = ~58 minutes - Validated: top-performing TF32 kernel (1.02 PFLOPS) is included Pattern format uses cutlass_profiler's glob matching (split on '*', substring search in order): *tf32gemm*256x256x*stream_k_2sm matches both '...stream_k_2sm' and '...stream_k_2sm_epi_tma' FP4 (e4m3 x e2m1) keeps the broad wildcard since it only has ~49 kernels total — no runtime savings from filtering. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
INT8 and FP8 256x256 stream_k_2sm kernels (without epi_tma) fail at runtime on GB200/GB300, producing zero results and return code 32. This happens because INT8/FP8 use 128-byte K-tiles (vs 32 for TF32) which changes SMEM pressure and supported tile configurations. Fix: broaden INT8/FP8 runtime patterns from *256x256x*stream_k_2sm to *256x*stream_k_2sm* which matches: - 128x256, 256x128, 256x256 tiles (competitive for large GEMMs) - Both stream_k_2sm and stream_k_2sm_epi_tma scheduling - ~168 kernels each (failed kernels return in ~2 sec, not 10.5) TF32/BF16/FP16 keep 256x256-only (56 each) since their 32/64-byte K-tiles work well with 256x256 tiles. Also add trailing * to all patterns so stream_k_2sm matches both '_2sm' and '_2sm_epi_tma' variants consistently. Validated: test1.log shows TF32 (1.03 PFLOPS), BF16 (1.96 PFLOPS), FP16 (1.99 PFLOPS) all succeeded; INT8 failed with the old pattern. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
CUTLASS shared per-family libraries (libcutlass_gemm_sm100_*.so) are installed to /opt/superbench/lib by cmake -DCMAKE_INSTALL_LIBDIR, but the container's LD_LIBRARY_PATH did not include this path. This caused cutlass_profiler to fail at runtime with: libcutlass.so: cannot open shared object file: No such file or directory INT8, FP8, and FP4 precisions silently failed (return code 32) while TF32/BF16/FP16 appeared to work from a fallback static-linked subset. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
SM103a (GB300 Blackwell Ultra) does NOT support INT8 UMMA — the tcgen05.mma.kind::i8 instruction was removed in favor of doubled FP4/NVFP4 throughput. CuTe's config.hpp only defines CUTE_ARCH_TCGEN05_S8_MMA_ENABLED for SM100a/SM101a; SM103a is deliberately excluded. INT8 kernels compiled for SM103a contain a runtime trap (CUTE_INVALID_CONTROL_PATH) instead of real PTX. Previous sed patch removed the arch_family_cc gate entirely, which forced INT8 generation for all archs but produced broken SM103a cubins that fail at runtime with: 'Attempting to use SM100_MMA_S8_SS without CUTE_ARCH_TCGEN05_S8_MMA_ENABLED' Fix (two parts): 1. cutlass.mk: Change sed patch to remove only '103a' from the gate (not the entire gate). This preserves INT8 generation for SM100a (GB200) while the gate still blocks on '100f' and '101f' family binaries as CUTLASS intended. 2. cuda_gemm_flops_performance.py: Exclude 'int8_tc' from the SM10.3 kernel map so the profiler never attempts INT8 on GB300. This is the primary safety net — the SM103a cubin trap is a backup. Result: INT8 works on GB200 (SM100a), cleanly skipped on GB300 (SM103a). Verified by GPT-5.4 and Claude Opus 4.7 code review. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
The previous fp4_tc metric used cutlass3x_sm100_tensorop_gemm_e4m3_e2m1_* which is FP8(e4m3) x FP4(e2m1) mixed-precision GEMM. This kernel runs on the FP8 datapath (tcgen05.mma.kind=f8f6f4) at FP8 throughput regardless of operand width — narrowing one input to FP4 saves bandwidth but does NOT double the math throughput. Result: measured ~3.8 PFLOPS vs NVIDIA NVFP4 spec of 15 PFLOPS (25%). Both GPT-5.4 and Claude Opus 4.7 confirmed this was a wrong-kernel issue. Fix: switch to true NVFP4 via CUTLASS BlockScaledGemm operation kind. SM103a (GB300 Blackwell Ultra) has dedicated FP4 Ultra kernels: cutlass3x_sm103_bstensorop_gemm_ue4m3xe2m1_*_ultra_2sm_* - ue4m3 = FP8 scale factor (NVFP4 spec) - vs16 = 16-element scale block (NVFP4 spec) - bstensorop = block-scaled tensor op - ultra_2sm = SM103-specific FP4 Ultra dual-CTA path Changes: 1. cutlass.mk: add bstensorop NVFP4 kernel patterns to build filter 2. cuda_gemm_flops_performance.py: - Replace 'fp4_tc' with 'nvfp4_tc' in SM10.3 kernel map - Use --operation=block_scaled_gemm (instead of gemm) for nvfp4_tc - Add parse_logline entry for block_scaled_gemm OperationKind 3. gemm_flops_performance_base.py: replace fp4_tc with nvfp4_tc precision 4. tests: replace test_flops_performance_cuda_fp4 with cuda_nvfp4 The mixed FP4 measurement is dropped entirely. cuBLASLt's fp4e2m1 benchmark already covers the production FP4 inference path. CLI verified: --operation=block_scaled_gemm works in cutlass_profiler. Kernel naming verified by inspecting CUTLASS v4.3.5 generator output. Validated by GPT-5.4 and Claude Opus 4.7 code review. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
# Conflicts: # superbench/benchmarks/micro_benchmarks/cuda_gemm_flops_performance.py # tests/benchmarks/micro_benchmarks/test_cuda_gemm_flops_performance.py # third_party/Makefile
…in (SM107a) install
CUTLASS v4.8.0 is the first release with SM107. When any 107 arch is in the
build, its generator retargets the SM100 generators to SM107, so the SM100
kernels disappear from that build, and cutlass_profiler only runs kernels from
its own libcutlass.so. The CUDA 13.4 tier now builds one v4.8.0 clone twice:
- "100a;103a" -> $SB_MICRO_PATH/{bin,lib} (GB200/GB300, same kernel names as v4.3.5)
- "107a" -> $SB_MICRO_PATH/cutlass_sm107/{bin,lib} (Rubin)
The v4.8.0 INT8 gate is ['100f', '101f', '103a', '107f']; the patch drops
'103a' as before and a grep fails the build if it no longer applies.
gemm-flops on CC 10.7 no longer aliases the Ampere (SM80) kernels:
- TF32/FP16/BF16: SM100 generators retargeted to Sm107 (stream-K 2SM 256x256)
- FP8: native SM107 runtime-typed f8 kernels with B-reuse schedules, launched
with --runtime_input_datatype_a/b=e4m3
- NVFP4: native SM107 block-scaled e2m1/ue4m3 vs16 kernels
- INT8 is not offered: SM107a has no tcgen05 kind::i8 MMA (the generated
kernel compiles without any tensor-core MMA)
On CC 10.7 the profiler is taken from cutlass_sm107/bin and its lib folder is
put first in LD_LIBRARY_PATH, since LD_LIBRARY_PATH outranks the profiler's
RUNPATH and both installs ship a libcutlass.so. An explicit --bin_dir still wins.
- CUDA < 11.8 (cuda11.1.1 image): build CUTLASS v2.4.0, the commit pinned as the third_party/cutlass submodule, with the full legacy kernel list, as main did. v3.9.2 needs CUDA 11.4+ and CMake 3.19+, and the Volta-only list dropped every A100 tensor-core kernel. - H100 FP8: build the e4m3 x e4m3 -> f32 WGMMA kernels (66) for "90a" in the CUDA 12.0-12.7 and 12.8 builds. Plain "90" compiles them without any WGMMA instruction. CUTLASS v3.x names put the MMA shape before "gemm" (s64x256x32gemm), so the build filter, the 9.0 fp8_tc pattern and the parse prefix accept both v3.x and v4.x names. - CUDA 12.8: use v4.3.5 with "80;90a;100a". v3.9.2 SM100 names (s128x256x8tf32gemm) matched neither the SM100 filters nor the GB200 patterns, and plain "100" generated no SM100 kernels. CUDA 12.9, 13.0 and 13.4 build commands are unchanged.
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Precision availability mismatches, explicit Rubin library selection, and architecture-dependent test failures remain unresolved.
4 open findings
What changed in this PR
Extends gemm-flops with Blackwell and Rubin CUTLASS kernels and low-precision benchmarking.
Changes:
- Adds architecture-specific kernel selection, FP8 and NVFP4 metrics.
- Centralizes CUDA-version-specific CUTLASS builds, with a separate Rubin installation.
- Updates library paths and expands kernel-selection, command-generation and parsing tests.
| File | Description |
|---|---|
| third_party/Makefile | Delegates CUTLASS builds to the new configuration. |
| third_party/cutlass.mk | Defines version-specific builds and kernel filters. |
| tests/benchmarks/micro_benchmarks/test_cuda_gemm_flops_performance.py | Adds architecture and precision tests. |
| superbench/benchmarks/micro_benchmarks/gemm_flops_performance_base.py | Registers FP8 and NVFP4 precisions and metrics. |
| superbench/benchmarks/micro_benchmarks/cuda_gemm_flops_performance.py | Adds kernel maps and Rubin profiler selection. |
| dockerfile/cuda13.4.dockerfile | Adds library lookup path and documents separate installations. |
| dockerfile/cuda13.0.dockerfile | Adds SuperBench's library lookup path. |
🧠 Review effort: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Comment on lines
+64
to
+65
| self.__kernel_map[10.0] = { | ||
| 'tf32_tc': 'cutlass3x_sm100_tensorop_tf32gemm*256x256x*stream_k_2sm*', |
| self.__kernel_map[10.3] = {k: self.__kernel_map[8.0][k] for k in self.__kernel_map[8.0] if 'int4_tc' not in k} | ||
| self.__kernel_map[10.7] = {k: self.__kernel_map[8.0][k] for k in self.__kernel_map[8.0] if 'int4_tc' not in k} | ||
| # FP8 (e4m3) on Hopper SM90a: CUTLASS 3x WGMMA kernels; v3.x names put the MMA shape before "gemm". | ||
| self.__kernel_map[9.0]['fp8_tc'] = 'cutlass3x_sm90_tensorop_*gemm_e4m3_e4m3_f32_f32_f32_*' |
Comment on lines
+169
to
+175
| subdir = self.__cutlass_subdir.get(self.__capability) | ||
| if subdir and self._args.bin_dir is None and os.getenv('SB_MICRO_PATH'): | ||
| prefix = os.path.join(os.getenv('SB_MICRO_PATH'), subdir) | ||
| self._args.bin_dir = os.path.join(prefix, 'bin') | ||
| # LD_LIBRARY_PATH outranks RUNPATH, so put this build's libcutlass*.so ahead of SB_MICRO_PATH/lib. | ||
| os.environ['LD_LIBRARY_PATH'] = os.path.join(prefix, 'lib') + os.pathsep + os.getenv('LD_LIBRARY_PATH', '') | ||
| return super()._set_binary_path() |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## main #883 +/- ##
==========================================
+ Coverage 86.31% 86.35% +0.04%
==========================================
Files 107 107
Lines 8256 8273 +17
==========================================
+ Hits 7126 7144 +18
+ Misses 1130 1129 -1
Flags with carried forward coverage won't be shown. Click here to find out more. ☔ View full report in Codecov by Harness. 🚀 New features to boost your workflow:
|
CUTLASS v4.3.5 only runs GenerateSM103_TensorOp_fp4_ultra_UMMA_gemm_with_block_scaled for CUDA 13.0+, so the cuda12.9 image had no kernels for GB300 nvfp4_tc. CUTLASS enables SM103 from CUDA 12.9 (arch/config.h) and the Ultra MMA has a CUDA 12.9 PTX form; all six Ultra schedule variants compile with CUDA 12.9.41 to the same UTCOMMA.2CTA.BLOCK16 instructions as with CUDA 13.4. The 12.9 tier now lowers only that generator's gate to 12.9 (sed limited to the function) and a grep fails the build if the patch no longer applies. With CUDA 12.9 the generated kernel list is identical to CUDA 13.0 (3,832 kernels, GB300 nvfp4_tc selects 72). Other tiers are unchanged.
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
NVFP4 output parsing rejects actual CUTLASS rows, and default CPU verification adds excessive work before throughput measurements.
5 open findings
1 resolved since last review
🧠 Review effort: Balanced
Comment on lines
+108
to
+111
| 'block_scaled_gemm,cutlass3x_sm103_bstensorop_gemm_ue4m3xe2m1_', | ||
| 'block_scaled_gemm,cutlass3x_sm103_bstensorop_gemm_ue4m3xf4_', 'gemm,cutlass3x_sm107_tensorop_tf32gemm_', | ||
| 'gemm,cutlass3x_sm107_tensorop_gemm_f16_', 'gemm,cutlass3x_sm107_tensorop_gemm_bf16_', | ||
| 'gemm,cutlass3x_sm107_tensorop_gemm_f8_f8_', 'block_scaled_gemm,cutlass3x_sm107_bstensorop_gemm_ue4m3xe2m1_' |
| # in cutlass_profiler — they have separate kernel registries from | ||
| # the standard gemm operation. | ||
| operation = 'block_scaled_gemm' if p in ('nvfp4_tc', 'mxfp4_tc') else 'gemm' | ||
| command += (' --operation=' + operation) |
…m-flops on Rubin New gemm-flops precisions for CC 10.7, using CUTLASS v4.8.0's native SM107 block-scaled kernels (B-reuse schedules, N=256 2SM tiles): - nvfp4_ue5m3_tc: NVFP4 (e2m1) with the new UE5M3 scale type, 16-element blocks - mxfp8_tc: MXFP8 (runtime-typed f8 run as e4m3, UE8M0 scales, 32-element blocks) - mxfp4_tc: MXFP4 (e2m1, UE8M0 scales, 32-element blocks) The profiler operation now follows the kernel family (bstensorop -> block_scaled_gemm), and every runtime-typed FP8 family (gemm_f8_f8, ue8m0xf8) gets --runtime_input_datatype_a/b=e4m3. SM107 block-scaled rows are parsed by one prefix. The Rubin build adds the native 2SM kernels of these formats (672 kernels, 3,360 total; largest library group 689 files). Docs list the FP8/FP4 metrics.
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Block-scaled result selectors reject actual CUTLASS CSV output, preventing NVFP4 and MXFP metrics from being recorded.
6 open findings
OperationKind::kGemmin… · New
🧠 Review effort: Balanced
Comment on lines
+115
to
+118
| 'block_scaled_gemm,cutlass3x_sm103_bstensorop_gemm_ue4m3xe2m1_', | ||
| 'block_scaled_gemm,cutlass3x_sm103_bstensorop_gemm_ue4m3xf4_', 'gemm,cutlass3x_sm107_tensorop_tf32gemm_', | ||
| 'gemm,cutlass3x_sm107_tensorop_gemm_f16_', 'gemm,cutlass3x_sm107_tensorop_gemm_bf16_', | ||
| 'gemm,cutlass3x_sm107_tensorop_gemm_f8_f8_', 'block_scaled_gemm,cutlass3x_sm107_bstensorop_' |
This branch has not been deployed
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
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.
Description
Adds Blackwell (GB200 SM100a, GB300 SM103a) and Rubin (SM107a) support to
gemm-flops, with FP8 and NVFP4 precisions (plus NVFP4 with UE5M3 scales, MXFP8 and MXFP4 on Rubin), and moves the CUTLASS build intothird_party/cutlass.mkwith one build per CUDA version.Major Revision
third_party/cutlass.mk(included bythird_party/Makefile) picks the CUTLASS tag, archs and kernel filter per CUDA version:On CUDA 13.4 the Rubin build is installed to
$SB_MICRO_PATH/cutlass_sm107/{bin,lib}. With any 107 arch, CUTLASS v4.8.0 retargets the SM100 generators to SM107, andcutlass_profileronly runs kernels from its ownlibcutlass.so, so Blackwell and Rubin need separate builds.On CUDA 12.9, CUTLASS v4.3.5 only generates the GB300 NVFP4 Ultra kernels for CUDA 13.0+. CUTLASS enables SM103 from CUDA 12.9, so
cutlass.mklowers only that generator's gate to 12.9; a grep stops the build if the patch no longer applies.gemm-flops:fp8_tcandnvfp4_tc, and on Rubinnvfp4_ue5m3_tc(NVFP4 with the new UE5M3 scale type),mxfp8_tcandmxfp4_tc.cutlass_profilerruns fromcutlass_sm107/bin, with itslibfolder first inLD_LIBRARY_PATH.--runtime_input_datatype_a/b=e4m3; block-scaled kernels run with--operation=block_scaled_gemm.Minor Revision
/opt/superbench/libtoLD_LIBRARY_PATH.Validation
make -nwere applied to a pristine copy of that CUTLASS tag's kernel generator, which was then run offline with the exact filters. Every precision of every GPU an image serves selects kernels, e.g. GB200 TF32/BF16/FP16 56 and INT8/FP8 168, GB300 NVFP4 72 (CUDA 12.9, 13.0 and 13.4), Rubin TF32/BF16/FP16 56, FP8 112, NVFP4 28, NVFP4-UE5M3 28, MXFP8 112 and MXFP4 28, H100 FP8 66. The one exception is H100 FP8 with CUDA 11.8, which has no sm_90a; no image uses CUDA 11.8.UTCOMMA.2CTA.BLOCK16instructions as with CUDA 13.4.