Skip to content

Commit

Permalink
[Bugfix][Codegen][CUDA] Wrong casting in ASM (#14782)
Browse files Browse the repository at this point in the history
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.
  • Loading branch information
junrushao authored May 6, 2023
1 parent 298a0a4 commit c224a12
Showing 1 changed file with 2 additions and 2 deletions.
4 changes: 2 additions & 2 deletions src/target/source/codegen_cuda.cc
Original file line number Diff line number Diff line change
Expand Up @@ -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<IntImm>(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 (
Expand Down

0 comments on commit c224a12

Please sign in to comment.