[mlir][nvgpu] Remove strict verifiers on warpgroup.generate.descriptor
#69935
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 PR relaxes some rules in the verifier. I found this to be overly restrictive. It's certainly possible to work around these rules, for example one way is to generate additional subview and etc., but this just bloats the IR.
|
@llvm/pr-subscribers-mlir @llvm/pr-subscribers-mlir-gpu Author: Guray Ozen (grypp) ChangesThis PR relaxes some rules in the verifier. I found this to be overly restrictive. It's certainly possible to work around these rules, for example one way is to generate additional subview and etc., but this just bloats the IR. Full diff: https://github.com/llvm/llvm-project/pull/69935.diff 1 Files Affected:
diff --git a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
index f5b02fe1b515591..15eeba2839479d8 100644
--- a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
+++ b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
@@ -375,15 +375,9 @@ LogicalResult WarpgroupGenerateDescriptorOp::verify() {
MemRefType memrefType = getTensor().getType();
MemRefType tensorMapType = getTensorMap().getType().getTensor();
- if (memrefType != tensorMapType)
- return emitError() << "memref and tensor map type mismatch";
-
if (!memrefType.hasStaticShape() || !tensorMapType.hasStaticShape())
return emitError() << "supports only static shapes";
- if (memrefType.getRank() != 2)
- return emitError() << "supports only 2d memref is supported for now";
-
if (getTensorMap().getType().getSwizzle() !=
TensorMapSwizzleKind::SWIZZLE_128B) {
return emitError() << "supports only "
|
|
The test #69913 needs this PR. Can you fold this change into the PR that needs it ? |
qcolombet
reviewed
Oct 24, 2023
| @@ -375,15 +375,9 @@ LogicalResult WarpgroupGenerateDescriptorOp::verify() { | |||
| MemRefType memrefType = getTensor().getType(); | |||
| MemRefType tensorMapType = getTensorMap().getType().getTensor(); | |||
|
|
|||
| if (memrefType != tensorMapType) | |||
| return emitError() << "memref and tensor map type mismatch"; | |||
There was a problem hiding this comment.
What's the semantic of unmatched memref/tensor types?
Could you add a test case demonstrating this case?
There was a problem hiding this comment.
The issue is similar what we discuss here.
| if (!memrefType.hasStaticShape() || !tensorMapType.hasStaticShape()) | ||
| return emitError() << "supports only static shapes"; | ||
|
|
||
| if (memrefType.getRank() != 2) | ||
| return emitError() << "supports only 2d memref is supported for now"; |
There was a problem hiding this comment.
Since this is supposed to be used to feed into wgmma operations, why do we need to support more than 2ds?
(Sorry for the dumb questions x).)
|
I created #70923 that is a better cleanup and improvement. |
grypp
added a commit
that referenced
this pull request
Nov 10, 2023
…#70028) PR #69913 added a GEMM test (128x128x128 F32 += F16 * F16) with if-statement. This PR adds the same test using predicates in PTX. Predicate support is enabled using _BasicPtxBuilderInterface_ `(nvgpu.opcode ..., predicate = %pred)`. The predicate condition is computed in `Step 2. [GPU] Elect fastest thread in CTA` inspired by cutlass. It is as follows: ``` lane_predicate = nvvm.elect.sync warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0) warp_idx_in_warp_group = warp_idx % 4 predicate = (lane_predicate & warp_idx_in_warp_group) ``` Depends on #70027 #69934 #69935 #69584
zahiraam
pushed a commit
to zahiraam/llvm-project
that referenced
this pull request
Nov 20, 2023
…llvm#70028) PR llvm#69913 added a GEMM test (128x128x128 F32 += F16 * F16) with if-statement. This PR adds the same test using predicates in PTX. Predicate support is enabled using _BasicPtxBuilderInterface_ `(nvgpu.opcode ..., predicate = %pred)`. The predicate condition is computed in `Step 2. [GPU] Elect fastest thread in CTA` inspired by cutlass. It is as follows: ``` lane_predicate = nvvm.elect.sync warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0) warp_idx_in_warp_group = warp_idx % 4 predicate = (lane_predicate & warp_idx_in_warp_group) ``` Depends on llvm#70027 llvm#69934 llvm#69935 llvm#69584
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 PR relaxes some rules in the verifier. I found this to be overly restrictive. It's certainly possible to work around these rules, for example one way is to generate additional subview and etc., but this just bloats the IR.
The test #69913 needs this PR.