---
id: pr-cutlass-1932
repo: NVIDIA/cutlass
pr: 1932
title: Blockwise Scaling for FP8
author: manishucsd
date: '2025-01-09'
url: https://github.com/NVIDIA/cutlass/pull/1932
source_category: upstream-code
architectures:
- sm90
tags:
- arch_pipeline
- block scaled
- collective builder
- epilogue
- fp8
- gemm_quant
- hopper
- wgmma
techniques:
- async-copy
- epilogue-fusion
- fine-grained-quantization
- pipeline-stages
- warp-specialization
hardware_features:
- fp8
- hopper
- tma
- wgmma
kernel_types:
- gemm
- runtime
languages:
- cuda-cpp
- cute-dsl
captured_at: '2026-05-16'
status: merged
merge_sha: ef5620dd1d5e7c0af33083fae17b43c34f1d8950
inclusion_reason: Inspect scale layout, accumulator type, tile/schedule choice, epilogue fusion, and partial-tile guards before deriving a candidate.
changed_paths:
- examples/54_hopper_fp8_warp_specialized_gemm/54_hopper_fp8_warp_specialized_gemm.cu
- examples/65_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/65_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling.cu
- examples/65_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/CMakeLists.txt
- examples/65_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/hopper_fp8_commandline.hpp
- examples/65_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/reference/host/gemm_with_blockwise_scaling.h
- examples/CMakeLists.txt
- include/cutlass/arch/barrier.h
- include/cutlass/gemm/collective/builders/sm90_gmma_builder.inl
- include/cutlass/gemm/collective/collective_builder_decl.hpp
- include/cutlass/gemm/collective/collective_mma.hpp
- include/cutlass/gemm/collective/fp8_accumulation.hpp
- include/cutlass/gemm/collective/sm90_mma_tma_gmma_rs_warpspecialized.hpp
- include/cutlass/gemm/collective/sm90_mma_tma_gmma_rs_warpspecialized_mixed_input.hpp
- include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized.hpp
- include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8.hpp
- include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp
- include/cutlass/gemm/collective/sm90_sparse_mma_tma_gmma_ss_warpspecialized.hpp
- include/cutlass/gemm/dispatch_policy.hpp
- include/cutlass/gemm/gemm.h
- include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp
- include/cutlass/pipeline/sm90_pipeline.hpp
artifact_dir: evidence/pull-bundles/cutlass/gh-1932
head_ref_oid: b2695263dc9f63ae5a2b5f2174a4036d84572096
---
# Blockwise Scaling for FP8

- Repository: `NVIDIA/cutlass`
- PR: [1932](https://github.com/NVIDIA/cutlass/pull/1932)
- Decision: `Inspect scale layout, accumulator type, tile/schedule choice, epilogue fusion, and partial-tile guards before deriving a candidate.`
- Reason: Inspect scale layout, accumulator type, tile/schedule choice, epilogue fusion, and partial-tile guards before deriving a candidate.

## Changed Paths

- `examples/54_hopper_fp8_warp_specialized_gemm/54_hopper_fp8_warp_specialized_gemm.cu`
- `examples/65_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/65_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling.cu`
- `examples/65_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/CMakeLists.txt`
- `examples/65_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/hopper_fp8_commandline.hpp`
- `examples/65_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/reference/host/gemm_with_blockwise_scaling.h`
- `examples/CMakeLists.txt`
- `include/cutlass/arch/barrier.h`
- `include/cutlass/gemm/collective/builders/sm90_gmma_builder.inl`
- `include/cutlass/gemm/collective/collective_builder_decl.hpp`
- `include/cutlass/gemm/collective/collective_mma.hpp`
- `include/cutlass/gemm/collective/fp8_accumulation.hpp`
- `include/cutlass/gemm/collective/sm90_mma_tma_gmma_rs_warpspecialized.hpp`
- `include/cutlass/gemm/collective/sm90_mma_tma_gmma_rs_warpspecialized_mixed_input.hpp`
- `include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized.hpp`
- `include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8.hpp`
- `include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp`
- `include/cutlass/gemm/collective/sm90_sparse_mma_tma_gmma_ss_warpspecialized.hpp`
- `include/cutlass/gemm/dispatch_policy.hpp`
- `include/cutlass/gemm/gemm.h`
- `include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp`
- `include/cutlass/pipeline/sm90_pipeline.hpp`

## Evidence Bundle

`evidence/pull-bundles/cutlass/gh-1932`
