---
id: pr-tilelang-1762
repo: tile-ai/tilelang
pr: 1762
title: '[Feature] Hierarchical reduction and warp reduction intrinsics support'
author: tzj-fxz
date: '2026-02-13'
url: https://github.com/tile-ai/tilelang/pull/1762
source_category: upstream-code
architectures:
- sm90
tags:
- WGMMA
- arch_pipeline
- benchmark_test
- compiler_runtime
- cuda backend
- intrinsic
- memory_primitives
techniques:
- pipeline-stages
- warp-specialization
hardware_features:
- wgmma
kernel_types:
- compiler
- memory
- reduction
- runtime
languages:
- cuda-cpp
- cute-dsl
- ptx
captured_at: '2026-05-16'
status: merged
merge_sha: 48a8f4a6ce45d5b58046279bd115364eb884158c
inclusion_reason: 'Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.'
changed_paths:
- src/op/finalize_reducer.cc
- src/op/reduce.cc
- src/tl_templates/cuda/reduce.h
- testing/python/language/test_tilelang_language_reduce.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1762
head_ref_oid: 1700eb1009743f183237799d2cbd7525e2ec9e5d
---
# [Feature] Hierarchical reduction and warp reduction intrinsics support

- Repository: `tile-ai/tilelang`
- PR: [1762](https://github.com/tile-ai/tilelang/pull/1762)
- Decision: `Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.`
- Reason: Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.

## Changed Paths

- `src/op/finalize_reducer.cc`
- `src/op/reduce.cc`
- `src/tl_templates/cuda/reduce.h`
- `testing/python/language/test_tilelang_language_reduce.py`

## Evidence Bundle

`evidence/pull-bundles/tilelang/gh-1762`
