---
id: pr-tilelang-1676
repo: tile-ai/tilelang
pr: 1676
title: '[Feature] Atomic Reduction Operations and Vectorization Enhancement'
author: LeiWang1999
date: '2026-01-15'
url: https://github.com/tile-ai/tilelang/pull/1676
source_category: upstream-code
architectures:
- sm100
tags:
- fused-op
techniques:
- kernel-fusion
- vectorized-loads
hardware_features: []
kernel_types:
- reduction
languages: []
captured_at: '2026-05-16'
status: merged
merge_sha: b27fb9286a0f05f3adffd6911bb71bf54f42c4e8
inclusion_reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.
changed_paths:
- src/op/atomic_add.cc
- src/op/atomic_add.h
- src/op/atomic_reduce.cc
- src/op/atomic_reduce.h
- src/op/builtin.cc
- src/op/builtin.h
- src/op/copy.cc
- src/op/fill.cc
- src/op/parallel.cc
- src/target/codegen_cuda.cc
- src/target/codegen_hip.cc
- src/target/utils.cc
- src/target/utils.h
- src/transform/atomicadd_vectorize.cc
- src/transform/atomicadd_vectorize.h
- src/transform/legalize_safe_memory_access.cc
- src/transform/loop_partition.cc
- src/transform/loop_partition.h
- src/transform/loop_vectorize.cc
- src/transform/loop_vectorize.h
- src/transform/lower_tile_op.cc
- testing/python/language/test_tilelang_language_atomic.py
- testing/python/transform/test_tilelang_transform_legalize_safe_memory_access.py
- tilelang/language/atomic.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1676
head_ref_oid: ec37e0dd2d2eabeb5f82a588f47c2010f5520520
---
# [Feature] Atomic Reduction Operations and Vectorization Enhancement

- Repository: `tile-ai/tilelang`
- PR: [1676](https://github.com/tile-ai/tilelang/pull/1676)
- Decision: `Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.`
- Reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.

## Changed Paths

- `src/op/atomic_add.cc`
- `src/op/atomic_add.h`
- `src/op/atomic_reduce.cc`
- `src/op/atomic_reduce.h`
- `src/op/builtin.cc`
- `src/op/builtin.h`
- `src/op/copy.cc`
- `src/op/fill.cc`
- `src/op/parallel.cc`
- `src/target/codegen_cuda.cc`
- `src/target/codegen_hip.cc`
- `src/target/utils.cc`
- `src/target/utils.h`
- `src/transform/atomicadd_vectorize.cc`
- `src/transform/atomicadd_vectorize.h`
- `src/transform/legalize_safe_memory_access.cc`
- `src/transform/loop_partition.cc`
- `src/transform/loop_partition.h`
- `src/transform/loop_vectorize.cc`
- `src/transform/loop_vectorize.h`
- `src/transform/lower_tile_op.cc`
- `testing/python/language/test_tilelang_language_atomic.py`
- `testing/python/transform/test_tilelang_transform_legalize_safe_memory_access.py`
- `tilelang/language/atomic.py`

## Evidence Bundle

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