This is an automated email from the ASF dual-hosted git repository. tqchen pushed a commit to branch tqchen/printer-refinement in repository https://gitbox.apache.org/repos/asf/tvm.git
commit 8d9ee4f90ddf26557487d4ebc40f163539a939af Author: Tianqi Chen <[email protected]> AuthorDate: Thu Oct 1 22:06:15 2026 +0000 [REFACTOR][Script] Preserve existing CUDA load emission --- src/backend/cuda/codegen/codegen_cuda.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/backend/cuda/codegen/codegen_cuda.cc b/src/backend/cuda/codegen/codegen_cuda.cc index 471cba6edc..8b365cdfa5 100644 --- a/src/backend/cuda/codegen/codegen_cuda.cc +++ b/src/backend/cuda/codegen/codegen_cuda.cc @@ -1295,14 +1295,14 @@ void CodeGenCUDA::Dispatch_(const CallNode* op, std::ostream& os) { std::string global_buffer = this->PrintExpr(addr_buffer->source.as_or_throw<tvm::tirx::BufferVar>().data()); std::string local_addr = this->PrintExpr(op->args[3]); - std::string destination = op->args[0].as<TensorLoadNode>() ? reg : reg + "[" + local_addr + "]"; this->stream << "asm volatile (\n"; this->stream << "\"{.reg .pred p;\\n\"\n"; this->stream << "\" setp.ne.b32 p, %2, 0;\\n\"\n"; this->stream << "\" @!p mov.b32 %0, 0;\\n\"\n"; this->stream << "\" @p ld.global.nc.f32 %0, [%1];}\\n\"\n"; // stream << "\" @p ld.global.nc.L2::128B.f32 %0, [%1];}\\n\"\n" ; - stream << ": \"=f\"(" << destination << ")\n"; + stream << ": \"=f\"(" << reg << "[" << local_addr << "]" + << ")\n"; stream << ": \"l\"((void*)(" << global_buffer << "+" << global_addr << ")), \"r\"((int)" << guard << ")\n"; stream << ");\n";
