[mlir][nvvm] Fix mov.u32 to mov.pred #70027
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This PR fixes the incorrect PTX instruction. We actually move a predicate not u32, so the correct instruction should be `mov.pred`.
|
@llvm/pr-subscribers-mlir-llvm @llvm/pr-subscribers-mlir Author: Guray Ozen (grypp) ChangesThis PR fixes the incorrect Full diff: https://github.com/llvm/llvm-project/pull/70027.diff 2 Files Affected:
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 9cda7862ccb0fe3..6947cf10e3600d4 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -469,9 +469,9 @@ def NVVM_ElectSyncOp : NVVM_Op<"elect.sync",
"{ \n"
".reg .u32 rx; \n"
".reg .pred px; \n"
- " mov.u32 %0, 0; \n"
+ " mov.pred %0, 0; \n"
" elect.sync rx | px, 0xFFFFFFFF;\n"
- "@px mov.u32 %0, 1; \n"
+ "@px mov.pred %0, 1; \n"
"}\n"
);
}
diff --git a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
index 3bb0ab90775edf5..b907a86ebc48072 100644
--- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
+++ b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
@@ -497,9 +497,9 @@ func.func @elect_one_leader_sync() {
// CHECK: llvm.inline_asm has_side_effects asm_dialect = att "{
// CHECK-SAME: .reg .u32 rx;
// CHECK-SAME: .reg .pred px;
- // CHECK-SAME: mov.u32 $0, 0;
+ // CHECK-SAME: mov.pred $0, 0;
// CHECK-SAME: elect.sync rx | px, 0xFFFFFFFF;
- // CHECK-SAME: @px mov.u32 $0, 1;
+ // CHECK-SAME: @px mov.pred $0, 1;
// CHECK-SAME: "=b" : () -> i1
%cnd = nvvm.elect.sync -> i1
return
|
nicolasvasilache
approved these changes
Oct 24, 2023
qcolombet
approved these changes
Oct 24, 2023
zahiraam
pushed a commit
to zahiraam/llvm-project
that referenced
this pull request
Oct 26, 2023
This PR fixes the incorrect `mov` instruction in PTX. We actually move a predicate here, not u32, so the correct instruction should be `mov.pred`.
grypp
added a commit
that referenced
this pull request
Nov 10, 2023
…#70028) PR #69913 added a GEMM test (128x128x128 F32 += F16 * F16) with if-statement. This PR adds the same test using predicates in PTX. Predicate support is enabled using _BasicPtxBuilderInterface_ `(nvgpu.opcode ..., predicate = %pred)`. The predicate condition is computed in `Step 2. [GPU] Elect fastest thread in CTA` inspired by cutlass. It is as follows: ``` lane_predicate = nvvm.elect.sync warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0) warp_idx_in_warp_group = warp_idx % 4 predicate = (lane_predicate & warp_idx_in_warp_group) ``` Depends on #70027 #69934 #69935 #69584
zahiraam
pushed a commit
to zahiraam/llvm-project
that referenced
this pull request
Nov 20, 2023
…llvm#70028) PR llvm#69913 added a GEMM test (128x128x128 F32 += F16 * F16) with if-statement. This PR adds the same test using predicates in PTX. Predicate support is enabled using _BasicPtxBuilderInterface_ `(nvgpu.opcode ..., predicate = %pred)`. The predicate condition is computed in `Step 2. [GPU] Elect fastest thread in CTA` inspired by cutlass. It is as follows: ``` lane_predicate = nvvm.elect.sync warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0) warp_idx_in_warp_group = warp_idx % 4 predicate = (lane_predicate & warp_idx_in_warp_group) ``` Depends on llvm#70027 llvm#69934 llvm#69935 llvm#69584
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This PR fixes the incorrect
movinstruction in PTX. We actually move a predicate here, not u32, so the correct instruction should bemov.pred.