Support for tritongpu.upcast_mxfp operation
#2700
Conversation
Signed-off-by: Tiotto, Ettore <[email protected]>
Signed-off-by: Tiotto, Ettore <[email protected]>
|
Note: Merging upstream to |
| // standalone values and returns them as a pair for (high 4 bits, low 4 bits). | ||
| std::pair<Value, Value> convertMxfp4x2ToBf16x2(RewriterBase &rewriter, | ||
| Location loc, Value v); | ||
| // Convert each value, which is an int8 containing 2 packed mxfp4 values, |
There was a problem hiding this comment.
Note: This is identical to upstream code as of commit 1cf7b1b
|
|
||
| return {v0, v1}; | ||
| } | ||
| SmallVector<Value> convertMxfp4x2ToBf16x2(RewriterBase &rewriter, Location loc, |
There was a problem hiding this comment.
Note: This is identical to upstream code as of commit 1cf7b1b
| @@ -19,17 +19,6 @@ using namespace mlir::triton::gpu; | |||
|
|
|||
| namespace { | |||
|
|
|||
| Value mxfpScaleBf16(RewriterBase &rewriter, Location loc, Value v, | |||
There was a problem hiding this comment.
Note: This is identical to upstream code as of commit 1cf7b1b
| @@ -30,47 +30,6 @@ class UpcastMXFPOpPattern : public ConvertOpToLLVMPattern<UpcastMXFPOp> { | |||
| : ConvertOpToLLVMPattern<UpcastMXFPOp>(typeConverter, benefit), | |||
| targetInfo(targetInfo) {} | |||
|
|
|||
| llvm::SmallVector<Value> unpackFP4Elements(Location loc, | |||
etiotto
requested review from
chengjunlu,
a team,
Dewei-Wang-sh,
alexbaden,
LiyangLingIntel and
quintinwang5
Merging in #2707. |
lib/Dialect/TritonGPU/IR/Ops.cpp
Outdated
| auto parentEncoding = oldEncoding.getParent(); | ||
|
|
||
| // Note: For Intel the dot operands layout's kWidth parameter must | ||
| // match the parent's dpas layout opsPerChannel. Given that the kWidth |
There was a problem hiding this comment.
opsPerChannel is defined by the HW DPAS instruction.
I think we should align the opsPerChannel to the result scalar type of the UpcastMXFPOp instead of double the size.
fp16/bf16 -> opsPerChannel=2
Otherwise there might be ambiguous in the lowering of UpcastMXFPOp.
| if (fpType == ScaleDotElemType::E2M1) | ||
| xVals = LLVM::convertMxfp4x2ToBf16x2(rewriter, loc, xVals); | ||
|
|
||
| // Each thread owns elements of 4 mxfp vectors so we need 4 scales |
There was a problem hiding this comment.
The better we need to make sure the layout are expected of the DotOp with the DPAS as parent and the layout conversion from the source operand to the dest operand are supported as well.
Add initial support for the new tritongpu
upcast_mxfpoperation.