---
id: pr-tilelang-1858
repo: tile-ai/tilelang
pr: 1858
title: '[Feature] Add TIR builtins for warp-level vote and block-level predicate sync'
author: sepcnt
date: '2026-04-13'
url: https://github.com/tile-ai/tilelang/pull/1858
source_category: upstream-code
architectures:
- sm100
tags:
- arch_pipeline
- benchmark_test
- compiler_runtime
- cuda backend
- intrinsic
techniques:
- pipeline-stages
- warp-specialization
hardware_features: []
kernel_types:
- compiler
- runtime
languages:
- cuda-cpp
captured_at: '2026-05-16'
status: merged
merge_sha: 19236b44c07104f7453013081f6bc01ad3ca5aa5
inclusion_reason: 'Extract the architecture assumption first: SM target, TMA/WGMMA/tcgen path, cluster use, and pipeline stages.'
changed_paths:
- docs/programming_guides/instructions.md
- src/op/builtin.cc
- src/op/builtin.h
- src/target/codegen_cuda.cc
- src/target/codegen_hip.cc
- testing/python/language/test_tilelang_language_warp_vote.py
- tilelang/language/__init__.py
- tilelang/language/builtin.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1858
head_ref_oid: 220c0e9e32e499b6f52a6ab606d5b1acbebf767c
---
# [Feature] Add TIR builtins for warp-level vote and block-level predicate sync

- Repository: `tile-ai/tilelang`
- PR: [1858](https://github.com/tile-ai/tilelang/pull/1858)
- 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

- `docs/programming_guides/instructions.md`
- `src/op/builtin.cc`
- `src/op/builtin.h`
- `src/target/codegen_cuda.cc`
- `src/target/codegen_hip.cc`
- `testing/python/language/test_tilelang_language_warp_vote.py`
- `tilelang/language/__init__.py`
- `tilelang/language/builtin.py`

## Evidence Bundle

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