[MLIR][NVGPU] Fix nvgpu_arrive syntax in matmulBuilder.py #113713
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 patch updates the syntax for nvgpu_arrive Op in matmulBuilder.py. This fixes the compilation error for this test. For the warp-specialized matmul_kernel implementation, removing the WaitGroupSyncOp (after the mma-main-loop) fixes the hang observed. With these two fixes, the test compiles and executes successfully on an sm90a machine. Signed-off-by: Durgadoss R <[email protected]>
|
@llvm/pr-subscribers-mlir Author: Durgadoss R (durga4github) ChangesThis patch updates the syntax for nvgpu_arrive Op For the warp-specialized matmul_kernel implementation, With these two fixes, the test compiles and Full diff: https://github.com/llvm/llvm-project/pull/113713.diff 1 Files Affected:
diff --git a/mlir/test/Integration/GPU/CUDA/sm90/python/tools/matmulBuilder.py b/mlir/test/Integration/GPU/CUDA/sm90/python/tools/matmulBuilder.py
index 75f0dc947e0681..5394d4a3272555 100644
--- a/mlir/test/Integration/GPU/CUDA/sm90/python/tools/matmulBuilder.py
+++ b/mlir/test/Integration/GPU/CUDA/sm90/python/tools/matmulBuilder.py
@@ -568,9 +568,7 @@ def generate_matmul_ws(
barId,
predicate=consumerPrimaryThread,
)
- nvgpu.mbarrier_arrive(
- ir.Type.parse("!nvgpu.mbarrier.token"), mbarDONE, barId
- )
+ nvgpu.mbarrier_arrive(mbarDONE, barId)
debug_print(
"[cons] iv={} | mbarDONE[{}] arrive [done]",
iv,
@@ -589,14 +587,9 @@ def generate_matmul_ws(
# Step 6.3.5. Yield
scf.yield_([new_acc, phaseParity])
- # Step 6.3. Wait All WGMMA
- nvvm.WgmmaWaitGroupSyncOp(0)
-
with ir.InsertionPoint(scf.IfOp(consumerPrimaryThread).then_block):
barId = c((K // BLOCK_K) % num_stages)
- nvgpu.mbarrier_arrive(
- ir.Type.parse("!nvgpu.mbarrier.token"), mbarDONE, barId
- )
+ nvgpu.mbarrier_arrive(mbarDONE, barId)
scf.yield_([])
# Step 6.4. Epilogue (registers --> shared memory)
|
|
@llvm/pr-subscribers-mlir-gpu Author: Durgadoss R (durga4github) ChangesThis patch updates the syntax for nvgpu_arrive Op For the warp-specialized matmul_kernel implementation, With these two fixes, the test compiles and Full diff: https://github.com/llvm/llvm-project/pull/113713.diff 1 Files Affected:
diff --git a/mlir/test/Integration/GPU/CUDA/sm90/python/tools/matmulBuilder.py b/mlir/test/Integration/GPU/CUDA/sm90/python/tools/matmulBuilder.py
index 75f0dc947e0681..5394d4a3272555 100644
--- a/mlir/test/Integration/GPU/CUDA/sm90/python/tools/matmulBuilder.py
+++ b/mlir/test/Integration/GPU/CUDA/sm90/python/tools/matmulBuilder.py
@@ -568,9 +568,7 @@ def generate_matmul_ws(
barId,
predicate=consumerPrimaryThread,
)
- nvgpu.mbarrier_arrive(
- ir.Type.parse("!nvgpu.mbarrier.token"), mbarDONE, barId
- )
+ nvgpu.mbarrier_arrive(mbarDONE, barId)
debug_print(
"[cons] iv={} | mbarDONE[{}] arrive [done]",
iv,
@@ -589,14 +587,9 @@ def generate_matmul_ws(
# Step 6.3.5. Yield
scf.yield_([new_acc, phaseParity])
- # Step 6.3. Wait All WGMMA
- nvvm.WgmmaWaitGroupSyncOp(0)
-
with ir.InsertionPoint(scf.IfOp(consumerPrimaryThread).then_block):
barId = c((K // BLOCK_K) % num_stages)
- nvgpu.mbarrier_arrive(
- ir.Type.parse("!nvgpu.mbarrier.token"), mbarDONE, barId
- )
+ nvgpu.mbarrier_arrive(mbarDONE, barId)
scf.yield_([])
# Step 6.4. Epilogue (registers --> shared memory)
|
|
@grypp, Please help with a review |
grypp
approved these changes
Oct 25, 2024
|
Can change the prefix of the title to [MLIR] |
durga4github
changed the title
Sure, updated. Builds are clean, submitting this. |
winner245
pushed a commit
to winner245/llvm-project
that referenced
this pull request
Oct 26, 2024
This patch updates the syntax for nvgpu_arrive Op in matmulBuilder.py. This fixes the compilation error for this test. For the warp-specialized matmul_kernel implementation, removing the WaitGroupSyncOp (after the mma-main-loop) fixes the hang observed. With these two fixes, the test compiles and executes successfully on an sm90a machine. Signed-off-by: Durgadoss R <[email protected]>
NoumanAmir657
pushed a commit
to NoumanAmir657/llvm-project
that referenced
this pull request
Nov 4, 2024
This patch updates the syntax for nvgpu_arrive Op in matmulBuilder.py. This fixes the compilation error for this test. For the warp-specialized matmul_kernel implementation, removing the WaitGroupSyncOp (after the mma-main-loop) fixes the hang observed. With these two fixes, the test compiles and executes successfully on an sm90a machine. Signed-off-by: Durgadoss R <[email protected]>
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 patch updates the syntax for nvgpu_arrive Op
in matmulBuilder.py. This fixes the compilation
error for this test.
For the warp-specialized matmul_kernel implementation,
removing the WaitGroupSyncOp (after the mma-main-loop)
fixes the hang observed.
With these two fixes, the test compiles and
executes successfully on an sm90a machine.