From c224a12e475695f5914699347401968f5c601b66 Mon Sep 17 00:00:00 2001 From: Junru Shao Date: Sat, 6 May 2023 07:21:43 -0700 Subject: [PATCH] [Bugfix][Codegen][CUDA] Wrong casting in ASM (#14782) The current HEAD will emit the inlined assembly that mistakenly include type casting, which is not allowed syntax, for example: ```c++ __asm__ __volatile__("cp.async.wait_group (int64_t)2;"); ```` This commit fixes this bug. --- src/target/source/codegen_cuda.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/target/source/codegen_cuda.cc b/src/target/source/codegen_cuda.cc index 0a84ed658034..d2131c522e38 100644 --- a/src/target/source/codegen_cuda.cc +++ b/src/target/source/codegen_cuda.cc @@ -931,8 +931,8 @@ void CodeGenCUDA::VisitExpr_(const CallNode* op, std::ostream& os) { } else if (op->op.same_as(builtin::ptx_commit_group())) { this->stream << "__asm__ __volatile__(\"cp.async.commit_group;\");\n\n"; } else if (op->op.same_as(builtin::ptx_wait_group())) { - std::string N = this->PrintExpr(op->args[0]); - this->stream << "__asm__ __volatile__(\"cp.async.wait_group " + N + ";\");\n\n"; + int n = Downcast(op->args[0])->value; + this->stream << "__asm__ __volatile__(\"cp.async.wait_group " << n << ";\");\n\n"; } else if (op->op.same_as(builtin::ptx_ldg32())) { /* asm volatile (