This is an automated email from the ASF dual-hosted git repository. tqchen pushed a commit to branch tvm-further-cleanup-python-tests-followup in repository https://gitbox.apache.org/repos/asf/tvm.git
commit 79e2bd38665088df1654e28429372407a0166a4b Author: Tianqi Chen <[email protected]> AuthorDate: Sun Jul 5 18:30:01 2026 +0000 [FIX][TIR] Type raw string packed arguments --- src/tirx/op/op.cc | 6 ++++ src/tirx/transform/lower_tvm_builtin.cc | 1 + .../test_tir_transform_lower_tvm_builtin.py | 35 ++++++++++++++++++++++ 3 files changed, 42 insertions(+) diff --git a/src/tirx/op/op.cc b/src/tirx/op/op.cc index 363121933f..2344ba6137 100644 --- a/src/tirx/op/op.cc +++ b/src/tirx/op/op.cc @@ -562,6 +562,12 @@ PrimExpr reinterpret(PrimType t, PrimExpr value, Span span) { } Expr reinterpret(Type target_ty, Expr value, Span span) { + if (value.as<StringImmNode>()) { + TVM_FFI_CHECK(target_ty.as<PointerTypeNode>(), TypeError) + << "String reinterpret requires a pointer target, but got " << target_ty; + return Call(std::move(target_ty), tirx::builtin::reinterpret(), {std::move(value)}, {}, {}, + std::move(span)); + } if (auto target_dtype = target_ty.as<PrimType>()) { if (auto prim_value = value.as<PrimExpr>()) { return reinterpret(target_dtype.value(), prim_value.value(), std::move(span)); diff --git a/src/tirx/transform/lower_tvm_builtin.cc b/src/tirx/transform/lower_tvm_builtin.cc index 49ee7d306e..55790e51a4 100644 --- a/src/tirx/transform/lower_tvm_builtin.cc +++ b/src/tirx/transform/lower_tvm_builtin.cc @@ -553,6 +553,7 @@ class BuiltinLower : public StmtExprMutator { int arg_type_index; if (arg.as<StringImmNode>()) { arg_type_index = ffi::TypeIndex::kTVMFFIRawStr; + arg = reinterpret(PointerType::VoidPointerTy(), std::move(arg)); } else if (arg->ty.as<PointerTypeNode>()) { arg_type_index = IsArrayHandle(arg) ? ffi::TypeIndex::kTVMFFIDLTensorPtr : ffi::TypeIndex::kTVMFFIOpaquePtr; diff --git a/tests/python/tirx-transform/test_tir_transform_lower_tvm_builtin.py b/tests/python/tirx-transform/test_tir_transform_lower_tvm_builtin.py index df3e8d529c..aeab2eb9ed 100644 --- a/tests/python/tirx-transform/test_tir_transform_lower_tvm_builtin.py +++ b/tests/python/tirx-transform/test_tir_transform_lower_tvm_builtin.py @@ -113,6 +113,41 @@ def test_lower_call_packed(): tvm.ir.assert_structural_equal(After, Expected) [email protected](not env.has_llvm(), reason="need llvm") +def test_lower_call_packed_raw_string(): + @I.ir_module + class Before: + @T.prim_func(s_tir=True) + def main(): + T.func_attr({"target": tvm.target.Target("llvm")}) + T.call_packed("testing.echo", "payload") + + After = tvm.tirx.transform.LowerTVMBuiltin()(Before) + packed_values = [] + + def collect_packed_values(node): + if not isinstance(node, tvm.ir.Call): + return + if node.op != tvm.ir.Op.get("tirx.tvm_struct_set"): + return + if int(node.args[2]) == 15: # kTVMFFIAnyUnionValue + packed_values.append(node.args[3]) + + tvm.tirx.stmt_functor.post_order_visit(After["main"].body, collect_packed_values) + raw_string = next( + value + for value in packed_values + if isinstance(value, tvm.ir.Call) + and value.op == tvm.ir.Op.get("tirx.reinterpret") + and isinstance(value.args[0], tvm.tirx.StringImm) + ) + assert isinstance(raw_string.ty, tvm.ir.PointerType) + tvm.ir.assert_structural_equal(After, tvm.script.from_source(After.script())) + + # The typed pointer is required by the LLVM TVMFFIAny lowering. + tvm.compile(Before, target="llvm") + + @pytest.mark.skipif(not env.has_llvm(), reason="need llvm") def test_call_packed_return_non_i32(): # This call packed that return non i32 types
