Skip to content

[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

Merged
merged 1 commit into from
May 23, 2024
Merged

[mlir][nvvm] Remove unused check-ptx #93147

merged 1 commit into from
May 23, 2024

Conversation

grypp
Copy link
Member

@grypp grypp commented May 23, 2024

The test used the check generated ptx with CHECK-PTX, but does not check that anymore. The PR removes these lines.

@llvmbot
Copy link
Member

llvmbot commented May 23, 2024

@llvm/pr-subscribers-mlir

Author: Guray Ozen (grypp)

Changes

Full diff: https://github.com/llvm/llvm-project/pull/93147.diff

1 Files Affected:

  • (modified) mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir (-9)
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 \

@llvmbot
Copy link
Member

llvmbot commented May 23, 2024

@llvm/pr-subscribers-mlir-gpu

Author: Guray Ozen (grypp)

Changes

Full diff: https://github.com/llvm/llvm-project/pull/93147.diff

1 Files Affected:

  • (modified) mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x8_8x128_noswizzle.mlir (-9)
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 \

@grypp grypp requested a review from matthias-springer May 23, 2024 10:01
@grypp grypp merged commit 7c137f7 into llvm:main May 23, 2024
6 checks passed
@grypp grypp deleted the remove-ptx branch May 23, 2024 10:09
@@ -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
Copy link
Contributor

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?

Copy link
Member Author

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.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants