Conversation
|
@llvm/pr-subscribers-mlir Author: Guray Ozen (grypp) ChangesFull diff: https://github.com/llvm/llvm-project/pull/93147.diff 1 Files Affected:
diff --git a/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir b/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir
index 2e59b7234e53d..391fda82e1e19 100644
--- a/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir
+++ b/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir
@@ -6,15 +6,6 @@
// RUN: --entry-point-result=void \
// RUN: | FileCheck %s
-// Basic PTX check to make sure we are generating the right instructions.
-
-// CHECK-PTX: mbarrier.init.shared.b64
-// CHECK-PTX: mbarrier.arrive.expect_tx.shared.b64
-// CHECK-PTX: cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes
-// CHECK-PTX: cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes
-// CHECK-PTX: mbarrier.arrive.expect_tx.shared.b64
-// CHECK-PTX: mbarrier.try_wait.parity.shared.b64
-
// RUN: mlir-opt %s --convert-nvgpu-to-nvvm \
// RUN: -gpu-kernel-outlining \
// RUN: -convert-nvvm-to-llvm \
|
|
@llvm/pr-subscribers-mlir-gpu Author: Guray Ozen (grypp) ChangesFull diff: https://github.com/llvm/llvm-project/pull/93147.diff 1 Files Affected:
diff --git a/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir b/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir
index 2e59b7234e53d..391fda82e1e19 100644
--- a/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir
+++ b/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir
@@ -6,15 +6,6 @@
// RUN: --entry-point-result=void \
// RUN: | FileCheck %s
-// Basic PTX check to make sure we are generating the right instructions.
-
-// CHECK-PTX: mbarrier.init.shared.b64
-// CHECK-PTX: mbarrier.arrive.expect_tx.shared.b64
-// CHECK-PTX: cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes
-// CHECK-PTX: cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes
-// CHECK-PTX: mbarrier.arrive.expect_tx.shared.b64
-// CHECK-PTX: mbarrier.try_wait.parity.shared.b64
-
// RUN: mlir-opt %s --convert-nvgpu-to-nvvm \
// RUN: -gpu-kernel-outlining \
// RUN: -convert-nvvm-to-llvm \
|
|
|
||
| // Basic PTX check to make sure we are generating the right instructions. | ||
|
|
||
| // CHECK-PTX: mbarrier.init.shared.b64 |
There was a problem hiding this comment.
This was changed here --check-prefixes=CHECK-PTX https://github.com/llvm/llvm-project/pull/68184/files#diff-468d25896ab3abaf617d05a64f7d107372665c6a6a9af470207dddd0f1f66a6d; should be CHECK-PTX prefix added back or kill all lines 1-8?
There was a problem hiding this comment.
There are merits of checking the PTX assembly. But this is an integration test. It's not the best place the test PTX. I've added that check earlier, but we now test the generated PTX in the conversion tests.
The test used the check generated ptx with
CHECK-PTX, but does not check that anymore. The PR removes these lines.