From 2d8d3676eda18bd3d8e6fa783399ff96d3cd4ded Mon Sep 17 00:00:00 2001 From: Lei Wang <34334180+LeiWang1999@users.noreply.github.com> Date: Wed, 14 Jan 2026 16:02:31 +0800 Subject: [PATCH] [CUDA] Enhance Broadcast Codegen for Symbolic Value (#1669) * Enhance CUDA code generation for BroadcastNode by implementing compile-time constant folding and runtime broadcasting for various lane configurations. Improved handling for 4-bit and 8-bit integer types, ensuring correct replication and type casting in output expressions. This update increases performance and correctness in CUDA kernel generation. * add test * lint fix * fix --- src/target/codegen_cuda.cc | 55 ++++++++++--------- .../test_tilelang_language_vectorize.py | 17 ++++++ 2 files changed, 46 insertions(+), 26 deletions(-) diff --git a/src/target/codegen_cuda.cc b/src/target/codegen_cuda.cc index 0dfb85341..1f7f202d0 100644 --- a/src/target/codegen_cuda.cc +++ b/src/target/codegen_cuda.cc @@ -3193,32 +3193,34 @@ void CodeGenTileLangCUDA::VisitExpr_(const BroadcastNode *op, std::ostream &os) { // NOLINT(*) int lanes = static_cast(Downcast(op->lanes)->value); if ((op->dtype.is_int() || op->dtype.is_uint()) && op->dtype.bits() == 8) { - if (lanes == 4) { - // make_int8x4 - const int64_t *p = as_const_int(op->value); - ICHECK(p); - int64_t v = *p & 0xFF; - v = (v << 24) | (v << 16) | (v << 8) | v; - if (op->dtype.is_uint()) { - os << "(uint)" << v; - } else { - os << "(int)" << v; + const int64_t *p = as_const_int(op->value); + if (p) { + if (lanes == 4) { + // make_int8x4 + ICHECK(p); + int64_t v = *p & 0xFF; + v = (v << 24) | (v << 16) | (v << 8) | v; + if (op->dtype.is_uint()) { + os << "(uint)" << v; + } else { + os << "(int)" << v; + } + return; + } else if (lanes == 32) { + // make_int8x32 + const int64_t *p = as_const_int(op->value); + ICHECK(p); + int64_t v = *p & 0xFF; + v = (v << 24) | (v << 16) | (v << 8) | v; + if (op->dtype.is_uint()) { + os << "make_ulonglong4(" << v << ", " << v << ", " << v << ", " << v + << ")"; + } else { + os << "make_longlong4(" << v << ", " << v << ", " << v << ", " << v + << ")"; + } + return; } - return; - } else if (lanes == 32) { - // make_int8x32 - const int64_t *p = as_const_int(op->value); - ICHECK(p); - int64_t v = *p & 0xFF; - v = (v << 24) | (v << 16) | (v << 8) | v; - if (op->dtype.is_uint()) { - os << "make_ulonglong4(" << v << ", " << v << ", " << v << ", " << v - << ")"; - } else { - os << "make_longlong4(" << v << ", " << v << ", " << v << ", " << v - << ")"; - } - return; } } @@ -3284,7 +3286,8 @@ void CodeGenTileLangCUDA::VisitExpr_(const BroadcastNode *op, if ((op->dtype.is_int() || op->dtype.is_uint()) && op->dtype.bits() == 4) { bool fail = false; const int64_t *p = as_const_int(op->value); - ICHECK(p); + ICHECK(p) << "BroadcastNode " << op << " value: " << op->value + << " is not a constant"; int64_t v = *p & 0xF; if (lanes == 4) { diff --git a/testing/python/language/test_tilelang_language_vectorize.py b/testing/python/language/test_tilelang_language_vectorize.py index 577527ae6..3ef107002 100644 --- a/testing/python/language/test_tilelang_language_vectorize.py +++ b/testing/python/language/test_tilelang_language_vectorize.py @@ -148,5 +148,22 @@ def test_vectorize_all_dtypes(dtype, vec_num): kernel(x) +@tilelang.jit +def vectorize_broadcast_int8(vec_num): + with T.Kernel(1, threads=128): + a = T.alloc_local((64,), "int8") + b = T.alloc_var("int8") + + for i in T.vectorized(vec_num): + a[i] = b + + +@tilelang.testing.requires_cuda +@pytest.mark.parametrize("vec_num", [4, 32]) +def test_vectorize_broadcast_int8(vec_num): + """Test broadcasting a non-constant int8 value to a vectorized store.""" + vectorize_broadcast_int8.compile(vec_num=vec_num) + + if __name__ == "__main__": tilelang.testing.main()