diff options
| author | Pradeep Kumar <pradeepisro49@gmail.com> | 2024-01-09 21:01:51 +0530 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2024-01-09 16:31:51 +0100 |
| commit | 0242d27dc89ff19e331ae4945933cdb360c7d4cf (patch) | |
| tree | 41dd46f780d116beb0e9f58710c815c35688b677 | |
| parent | db1d9ad109d8e0f17acb2de60c8b57085fe2de77 (diff) | |
[MLIR][NVVM] Add missing `;` when lowering stmatrix Op (#77471)
| -rw-r--r-- | mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 4 | ||||
| -rw-r--r-- | mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir | 8 |
2 files changed, 6 insertions, 6 deletions
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td index 52857164ffaf..3a6c6e5438c6 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td @@ -1345,8 +1345,8 @@ def NVVM_StMatrixOp: NVVM_PTXBuilder_Op<"stmatrix">, ptx += ".x" + std::to_string(d); if (getLayout() == NVVM::MMALayout::col) ptx += ".trans"; - if(d == 1) ptx += ".m8n8.shared.b16 [%0], {%1}"; - if(d == 2) ptx += ".m8n8.shared.b16 [%0], {%1, %2}"; + if(d == 1) ptx += ".m8n8.shared.b16 [%0], {%1};"; + if(d == 2) ptx += ".m8n8.shared.b16 [%0], {%1, %2};"; if(d == 4) ptx += ".m8n8.shared.b16 [%0], {%1, %2, %3, %4};"; return ptx; } diff --git a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir index 74186138c3a9..7e08ec6ffcbd 100644 --- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir +++ b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir @@ -599,11 +599,11 @@ func.func @elect_one_leader_sync() { // CHECK-SAME: %[[arg3:[a-zA-Z0-9_]+]]: i32, // CHECK-SAME: %[[arg4:[a-zA-Z0-9_]+]]: i32) llvm.func @stmatrix(%arg0 : !llvm.ptr<3>, %m1 : i32, %m2 : i32, %m3 : i32, %m4 : i32) { -// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "stmatrix.sync.aligned.x1.m8n8.shared.b16 [$0], {$1}", "r,r" %[[arg0]], %[[arg1]] : (!llvm.ptr<3>, i32) -> () -// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "stmatrix.sync.aligned.x2.m8n8.shared.b16 [$0], {$1, $2}", "r,r,r" %[[arg0]], %[[arg1]], %[[arg2]] : (!llvm.ptr<3>, i32, i32) -> () +// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "stmatrix.sync.aligned.x1.m8n8.shared.b16 [$0], {$1};", "r,r" %[[arg0]], %[[arg1]] : (!llvm.ptr<3>, i32) -> () +// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "stmatrix.sync.aligned.x2.m8n8.shared.b16 [$0], {$1, $2};", "r,r,r" %[[arg0]], %[[arg1]], %[[arg2]] : (!llvm.ptr<3>, i32, i32) -> () // CHECK: llvm.inline_asm has_side_effects asm_dialect = att "stmatrix.sync.aligned.x4.m8n8.shared.b16 [$0], {$1, $2, $3, $4};", "r,r,r,r,r" %[[arg0]], %[[arg1]], %[[arg2]], %[[arg3]], %[[arg4]] : (!llvm.ptr<3>, i32, i32, i32, i32) -> () -// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "stmatrix.sync.aligned.x1.trans.m8n8.shared.b16 [$0], {$1}", "r,r" %[[arg0]], %[[arg1]] : (!llvm.ptr<3>, i32) -> () -// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "stmatrix.sync.aligned.x2.trans.m8n8.shared.b16 [$0], {$1, $2}", "r,r,r" %[[arg0]], %[[arg1]], %[[arg2]] : (!llvm.ptr<3>, i32, i32) -> () +// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "stmatrix.sync.aligned.x1.trans.m8n8.shared.b16 [$0], {$1};", "r,r" %[[arg0]], %[[arg1]] : (!llvm.ptr<3>, i32) -> () +// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "stmatrix.sync.aligned.x2.trans.m8n8.shared.b16 [$0], {$1, $2};", "r,r,r" %[[arg0]], %[[arg1]], %[[arg2]] : (!llvm.ptr<3>, i32, i32) -> () // CHECK: llvm.inline_asm has_side_effects asm_dialect = att "stmatrix.sync.aligned.x4.trans.m8n8.shared.b16 [$0], {$1, $2, $3, $4};", "r,r,r,r,r" %[[arg0]], %[[arg1]], %[[arg2]], %[[arg3]], %[[arg4]] : (!llvm.ptr<3>, i32, i32, i32, i32) -> () nvvm.stmatrix %arg0, %m1 {layout = #nvvm.mma_layout<row>} : !llvm.ptr<3>, i32 nvvm.stmatrix %arg0, %m1, %m2 {layout = #nvvm.mma_layout<row>} : !llvm.ptr<3>, i32, i32 |