[MLIR][NVVM] [NFC] Update test cmd-lines and doc links #128207
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
For the NVVM Dialect tests under Target/LLVMIR/nvvm/ dir, we verify the lowering to the intrinsics using mlir-translate. For these tests, remove the -verify-diagnostics option from the cmd-line since all the verifier checks are tested through the nvvmir-invalid.mlir file. Similarly, remove the split-input-file option which is not relevant here. Fix a few remaining links in the NVVMOps.td file. All the reference links follow the same style now. Rename the tcgen05-barriers.mlir file to tcgen05-commit.mlir and move the wait/fence tests to a separate file. Signed-off-by: Durgadoss R <[email protected]>
|
@grypp , Please help with review |
|
@llvm/pr-subscribers-mlir-llvm @llvm/pr-subscribers-mlir Author: Durgadoss R (durga4github) ChangesFor the NVVM Dialect tests under Target/LLVMIR/nvvm/ dir, Update a few remaining links in the NVVMOps.td file. Rename the tcgen05-barriers.mlir file to tcgen05-commit.mlir Full diff: https://github.com/llvm/llvm-project/pull/128207.diff 10 Files Affected:
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 0692e8e32dbf8..633e4aaba5462 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -2648,8 +2648,7 @@ def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc"> {
the amount specified by `nCols` and writes the destination
address to the `addr` argument. The `nCols` operand specifies the
number of columns to be allocated and it must be a power-of-two.
- [For more information, refer to the PTX ISA]
- (https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-memory-alloc-manage-instructions)
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-memory-alloc-manage-instructions)
}];
let arguments = (ins
@@ -2679,8 +2678,7 @@ def NVVM_Tcgen05DeallocOp : NVVM_Op<"tcgen05.dealloc"> {
specified by `tmemAddr`, which must be from a previous tensor
memory allocation. The `nCols` operand specifies the number
of columns to be de-allocated, and it must be a power-of-two.
- [For more information, refer to the PTX ISA]
- (https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-memory-alloc-manage-instructions)
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-memory-alloc-manage-instructions)
}];
let arguments = (ins LLVM_PointerTensor:$taddr, I32:$nCols,
@@ -2708,8 +2706,7 @@ def NVVM_Tcgen05RelinquishAllocPermitOp : NVVM_Op<"tcgen05.relinquish_alloc_perm
of the executing thread is relinquishing the right to allocate
Tensor Memory. So, it is illegal for a CTA to perform `tcgen05.alloc`
after any of its constituent threads execute `tcgen05.relinquish_alloc_permit`.
- [For more information, refer to the PTX ISA]
- (https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-memory-alloc-manage-instructions)
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-memory-alloc-manage-instructions)
}];
let arguments = (ins
@@ -2733,8 +2730,7 @@ def NVVM_Tcgen05FenceOp : NVVM_Op<"tcgen05.fence"> {
The `tcgen05.fence<after>` orders all subsequent async tcgen05 operations
with respect to the prior tcgen05 and execution ordering operations.
- [For more information refer to the PTX ISA]
- (https://docs.nvidia.com/cuda/parallel-thread-execution/#tensorcore-5th-generation-instructions-tcgen05-fence)
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tensorcore-5th-generation-instructions-tcgen05-fence)
}];
let arguments = (ins Tcgen05FenceKindAttr:$kind);
@@ -2756,8 +2752,7 @@ def NVVM_Tcgen05WaitOp : NVVM_Op<"tcgen05.wait"> {
have completed. Similarly, the `tcgen05.wait<store>` causes the executing
thread to block until all prior `tcgen05.st` operations issued by the
executing thread have completed.
- [For more information refer PTX ISA]
- (https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-wait)
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-instructions-tcgen05-wait)
}];
let arguments = (ins Tcgen05WaitKindAttr:$kind);
@@ -2782,8 +2777,7 @@ def NVVM_Tcgen05CommitOp : NVVM_Op<"tcgen05.commit"> {
when present, specifies the destination CTAs in the cluster such
that each bit position in the 16-bit `multicastMask` operand
corresponds to the `nvvm.read.ptx.sreg.ctaid` of the destination CTA.
- [For more information refer PTX ISA]
- (https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen-async-sync-operations-commit)
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen-async-sync-operations-commit)
}];
let arguments = (ins
diff --git a/mlir/test/Target/LLVMIR/nvvm/cvt_tf32.mlir b/mlir/test/Target/LLVMIR/nvvm/cvt_tf32.mlir
index ff7bad0149d4c..2bce9e1a5d3e4 100644
--- a/mlir/test/Target/LLVMIR/nvvm/cvt_tf32.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/cvt_tf32.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-translate -mlir-to-llvmir %s -split-input-file --verify-diagnostics | FileCheck %s
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
// CHECK-LABEL: @convert_float_to_tf32_rna
llvm.func @convert_float_to_tf32_rna(%src : f32) -> i32 {
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir
index 781efa2567111..6a7e4ac515b81 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-alloc.mlir
@@ -1,5 +1,4 @@
-// RUN: mlir-opt -split-input-file -verify-diagnostics %s
-// RUN: mlir-translate -mlir-to-llvmir -split-input-file -verify-diagnostics %s | FileCheck %s --check-prefix=CHECK-LLVM
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s --check-prefix=CHECK-LLVM
// CHECK-LABEL: @llvm_nvvm_tcgen05_alloc
llvm.func @llvm_nvvm_tcgen05_alloc(%addr : !llvm.ptr, %ncols : i32) {
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-barriers.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir
similarity index 66%
rename from mlir/test/Target/LLVMIR/nvvm/tcgen05-barriers.mlir
rename to mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir
index 7536a4567e34e..80cf29f3704c2 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-barriers.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-commit.mlir
@@ -1,27 +1,4 @@
-// RUN: mlir-opt -split-input-file -verify-diagnostics %s
-// RUN: mlir-translate -mlir-to-llvmir -split-input-file -verify-diagnostics %s | FileCheck %s --check-prefix=CHECK-LLVM
-
-// CHECK-LABEL: @llvm_nvvm_tcgen05_fence
-llvm.func @llvm_nvvm_tcgen05_fence() {
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.fence.before.thread.sync()
- nvvm.tcgen05.fence #nvvm.tcgen05_fence<before>
-
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.fence.after.thread.sync()
- nvvm.tcgen05.fence #nvvm.tcgen05_fence<after>
-
- llvm.return
-}
-
-// CHECK-LABEL: @llvm_nvvm_tcgen05_wait
-llvm.func @llvm_nvvm_tcgen05_wait() {
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.wait.ld()
- nvvm.tcgen05.wait #nvvm.tcgen05_wait<load>
-
- // CHECK-LLVM: call void @llvm.nvvm.tcgen05.wait.st()
- nvvm.tcgen05.wait #nvvm.tcgen05_wait<store>
-
- llvm.return
-}
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s --check-prefix=CHECK-LLVM
// CHECK-LABEL: @llvm_nvvm_tcgen05_commit_generic
llvm.func @llvm_nvvm_tcgen05_commit_generic(%barrier : !llvm.ptr, %cta_mask : i16) {
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-cp.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-cp.mlir
index 91128cd00c873..bf72714d16de7 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-cp.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-cp.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-translate -mlir-to-llvmir -split-input-file %s | FileCheck %s
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
// CHECK-LABEL: @nvvm_tcgen05_cp_128x256b
llvm.func @nvvm_tcgen05_cp_128x256b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-fence-wait.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-fence-wait.mlir
new file mode 100644
index 0000000000000..ee4a517a4bffa
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-fence-wait.mlir
@@ -0,0 +1,23 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s --check-prefix=CHECK-LLVM
+
+// CHECK-LABEL: @llvm_nvvm_tcgen05_fence
+llvm.func @llvm_nvvm_tcgen05_fence() {
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.fence.before.thread.sync()
+ nvvm.tcgen05.fence #nvvm.tcgen05_fence<before>
+
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.fence.after.thread.sync()
+ nvvm.tcgen05.fence #nvvm.tcgen05_fence<after>
+
+ llvm.return
+}
+
+// CHECK-LABEL: @llvm_nvvm_tcgen05_wait
+llvm.func @llvm_nvvm_tcgen05_wait() {
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.wait.ld()
+ nvvm.tcgen05.wait #nvvm.tcgen05_wait<load>
+
+ // CHECK-LLVM: call void @llvm.nvvm.tcgen05.wait.st()
+ nvvm.tcgen05.wait #nvvm.tcgen05_wait<store>
+
+ llvm.return
+}
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-shift.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-shift.mlir
index 48753a3fdb21b..78c50cf96cf90 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-shift.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-shift.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-translate -mlir-to-llvmir -split-input-file %s | FileCheck %s
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
// CHECK-LABEL: @llvm_nvvm_tcgen05_shift
llvm.func @llvm_nvvm_tcgen05_shift(%taddr : !llvm.ptr<6>) {
diff --git a/mlir/test/Target/LLVMIR/nvvm/tma_bulk_copy.mlir b/mlir/test/Target/LLVMIR/nvvm/tma_bulk_copy.mlir
index aa2d680f5117e..0e3f98a134491 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tma_bulk_copy.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tma_bulk_copy.mlir
@@ -1,5 +1,4 @@
-// RUN: mlir-opt -split-input-file -verify-diagnostics %s
-// RUN: mlir-translate -mlir-to-llvmir -split-input-file -verify-diagnostics %s | FileCheck %s
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
// CHECK-LABEL: @llvm_nvvm_cp_async_bulk_global_to_shared_cluster
llvm.func @llvm_nvvm_cp_async_bulk_global_to_shared_cluster(%dst : !llvm.ptr<3>, %src : !llvm.ptr<1>, %mbar : !llvm.ptr<3>, %size : i32, %mc : i16, %ch : i64) {
diff --git a/mlir/test/Target/LLVMIR/nvvm/tma_prefetch.mlir b/mlir/test/Target/LLVMIR/nvvm/tma_prefetch.mlir
index 7be29fd616a6f..f1fa3b61f2dd9 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tma_prefetch.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tma_prefetch.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-translate -mlir-to-llvmir %s -split-input-file --verify-diagnostics | FileCheck %s
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
// CHECK-LABEL: @tma_prefetch_1d
llvm.func @tma_prefetch_1d(%tma_desc : !llvm.ptr, %d0 : i32, %ch : i64) {
diff --git a/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce.mlir b/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce.mlir
index 3809bc0bce897..6e0b48489e8b0 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tma_store_reduce.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-translate -mlir-to-llvmir -split-input-file --verify-diagnostics %s | FileCheck %s
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
// CHECK-LABEL: define void @tma_store_reduce_1d(
llvm.func @tma_store_reduce_1d(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %d0 : i32, %ch : i64) {
|
grypp
approved these changes
Feb 21, 2025
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.
For the NVVM Dialect tests under Target/LLVMIR/nvvm/ dir,
we verify the lowering to the intrinsics using mlir-translate.
Remove the -verify-diagnostics option from the cmd-line
for these tests since all the verifier checks are tested through
the nvvmir-invalid.mlir file. Similarly, remove the split-input-file
option which is not relevant here.
Update a few remaining links in the NVVMOps.td file.
All the reference links follow the same style now.
Rename the tcgen05-barriers.mlir file to tcgen05-commit.mlir
and move the wait/fence tests to a separate file.