---
id: pr-tilelang-1717
repo: tile-ai/tilelang
pr: 1717
title: '[Enhancement] Add explicit global memory load/store intrinsics (ldg/stg 32/64/128)'
author: LeiWang1999
date: '2026-01-26'
url: https://github.com/tile-ai/tilelang/pull/1717
source_category: upstream-code
architectures:
- sm100
tags:
- arch_pipeline
- benchmark_test
- compiler_runtime
- copy lowering
- memory_primitives
- scheduler_autotune
techniques:
- autotuning
- pipeline-stages
- tile-scheduling
- vectorized-loads
hardware_features:
- blackwell
kernel_types:
- compiler
- memory
- runtime
- scheduler
languages:
- cuda-cpp
captured_at: '2026-05-16'
status: merged
merge_sha: 8bd0293740928d067d2690d83ca1d0752f263686
inclusion_reason: 'Treat compiler/runtime integration as part of the kernel: launch wrapper, cache/JIT behavior, generated code, and build flags need tests.'
changed_paths:
- 3rdparty/tvm
- src/op/builtin.cc
- src/op/builtin.h
- src/target/codegen_cuda.cc
- src/tl_templates/cuda/copy.h
- src/tl_templates/cuda/copy_sm100.h
- src/transform/lower_ldg_stg.cc
- testing/python/language/test_tilelang_language_ldg.py
- testing/python/language/test_tilelang_language_stg.py
- testing/python/transform/test_tilelang_transform_lower_ldgstg.py
- tilelang/engine/phase.py
- tilelang/jit/__init__.py
- tilelang/language/__init__.py
- tilelang/language/builtin.py
- tilelang/language/utils.py
- tilelang/transform/__init__.py
- tilelang/transform/pass_config.py
- tilelang/utils/language.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1717
head_ref_oid: cbf1425ec2b97d9ffe0e44d40dd3371f04434ae8
---
# [Enhancement] Add explicit global memory load/store intrinsics (ldg/stg 32/64/128)

- Repository: `tile-ai/tilelang`
- PR: [1717](https://github.com/tile-ai/tilelang/pull/1717)
- Decision: `Treat compiler/runtime integration as part of the kernel: launch wrapper, cache/JIT behavior, generated code, and build flags need tests.`
- Reason: Treat compiler/runtime integration as part of the kernel: launch wrapper, cache/JIT behavior, generated code, and build flags need tests.

## Changed Paths

- `3rdparty/tvm`
- `src/op/builtin.cc`
- `src/op/builtin.h`
- `src/target/codegen_cuda.cc`
- `src/tl_templates/cuda/copy.h`
- `src/tl_templates/cuda/copy_sm100.h`
- `src/transform/lower_ldg_stg.cc`
- `testing/python/language/test_tilelang_language_ldg.py`
- `testing/python/language/test_tilelang_language_stg.py`
- `testing/python/transform/test_tilelang_transform_lower_ldgstg.py`
- `tilelang/engine/phase.py`
- `tilelang/jit/__init__.py`
- `tilelang/language/__init__.py`
- `tilelang/language/builtin.py`
- `tilelang/language/utils.py`
- `tilelang/transform/__init__.py`
- `tilelang/transform/pass_config.py`
- `tilelang/utils/language.py`

## Evidence Bundle

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