-
Notifications
You must be signed in to change notification settings - Fork 13.5k
[mlir][nvvm] Remove unused check-ptx #93147
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
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 \
|
@@ -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 |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.