---
id: pr-cutlass-2139
repo: NVIDIA/cutlass
pr: 2139
title: Blockwise and Groupwise GEMM for Blackwell and Improvements for Hopper
author: depaulmillz
date: '2025-02-26'
url: https://github.com/NVIDIA/cutlass/pull/2139
source_category: upstream-code
architectures:
- sm100
- sm90
tags:
- gemm
- grouped-gemm
- fp8
- block-scale
techniques:
- warp-specialization
- fine-grained-quantization
hardware_features:
- tcgen05
- tmem
- block-scale
- nvfp4
- fp8
kernel_types:
- gemm
- grouped-gemm
languages:
- cuda-cpp
- cute-dsl
captured_at: 2026-04-17
status: merged
merge_sha: ca4fdbea
inclusion_reason: manually curated, Blackwell kernel relevance
changed_paths:
- examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling.cu
- examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_groupwise_scaling.cu
- examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/hopper_fp8_commandline.hpp
- examples/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling.cu
- examples/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling/CMakeLists.txt
- examples/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling/hopper_fp8_commandline.hpp
- examples/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling/reference/host/gemm_with_groupwise_scaling.h
- examples/81_blackwell_gemm_blockwise/81_blackwell_gemm_blockwise.cu
- examples/81_blackwell_gemm_blockwise/81_blackwell_gemm_groupwise.cu
- examples/81_blackwell_gemm_blockwise/CMakeLists.txt
- examples/CMakeLists.txt
- include/cutlass/detail/sm100_blockwise_scale_layout.hpp
- include/cutlass/gemm/collective/builders/sm100_blockwise_umma_builder.inl
- include/cutlass/gemm/collective/builders/sm90_gmma_builder.inl
- include/cutlass/gemm/collective/collective_builder.hpp
- include/cutlass/gemm/collective/collective_mma.hpp
- include/cutlass/gemm/collective/fp8_accumulation.hpp
- include/cutlass/gemm/collective/sm100_mma_warpspecialized_blockwise_scaling.hpp
- include/cutlass/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input.hpp
- include/cutlass/gemm/collective/sm90_mma_array_tma_gmma_ss_warpspecialized.hpp
- include/cutlass/gemm/collective/sm90_mma_array_tma_gmma_ss_warpspecialized_fp8.hpp
- include/cutlass/gemm/collective/sm90_mma_array_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp
- include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp
- include/cutlass/gemm/dispatch_policy.hpp
- include/cutlass/gemm/kernel/gemm_universal.hpp
- include/cutlass/gemm/kernel/sm100_gemm_tma_warpspecialized_mma_transform.hpp
- include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp
- include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp
artifact_dir: evidence/pull-bundles/cutlass/gh-2139
head_ref_oid: null
---
# Blockwise and Groupwise GEMM for Blackwell and Improvements for Hopper

- Repository: `NVIDIA/cutlass`
- PR: [2139](https://github.com/NVIDIA/cutlass/pull/2139)
- Decision: `manually curated, Blackwell kernel relevance`
- Reason: manually curated, Blackwell kernel relevance

## Changed Paths

- `examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling.cu`
- `examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/67_hopper_fp8_warp_specialized_gemm_with_groupwise_scaling.cu`
- `examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/hopper_fp8_commandline.hpp`
- `examples/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling.cu`
- `examples/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling/CMakeLists.txt`
- `examples/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling/hopper_fp8_commandline.hpp`
- `examples/68_hopper_fp8_warp_specialized_grouped_gemm_with_blockwise_scaling/reference/host/gemm_with_groupwise_scaling.h`
- `examples/81_blackwell_gemm_blockwise/81_blackwell_gemm_blockwise.cu`
- `examples/81_blackwell_gemm_blockwise/81_blackwell_gemm_groupwise.cu`
- `examples/81_blackwell_gemm_blockwise/CMakeLists.txt`
- `examples/CMakeLists.txt`
- `include/cutlass/detail/sm100_blockwise_scale_layout.hpp`
- `include/cutlass/gemm/collective/builders/sm100_blockwise_umma_builder.inl`
- `include/cutlass/gemm/collective/builders/sm90_gmma_builder.inl`
- `include/cutlass/gemm/collective/collective_builder.hpp`
- `include/cutlass/gemm/collective/collective_mma.hpp`
- `include/cutlass/gemm/collective/fp8_accumulation.hpp`
- `include/cutlass/gemm/collective/sm100_mma_warpspecialized_blockwise_scaling.hpp`
- `include/cutlass/gemm/collective/sm90_mma_array_tma_gmma_rs_warpspecialized_mixed_input.hpp`
- `include/cutlass/gemm/collective/sm90_mma_array_tma_gmma_ss_warpspecialized.hpp`
- `include/cutlass/gemm/collective/sm90_mma_array_tma_gmma_ss_warpspecialized_fp8.hpp`
- `include/cutlass/gemm/collective/sm90_mma_array_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp`
- `include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp`
- `include/cutlass/gemm/dispatch_policy.hpp`
- `include/cutlass/gemm/kernel/gemm_universal.hpp`
- `include/cutlass/gemm/kernel/sm100_gemm_tma_warpspecialized_mma_transform.hpp`
- `include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp`
- `include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp`

## Evidence Bundle

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