Support for tritongpu.upcast_mxfp operation
#2700
Conversation
Signed-off-by: Tiotto, Ettore <ettore.tiotto@intel.com>
Signed-off-by: Tiotto, Ettore <ettore.tiotto@intel.com>
|
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 1cf7b1b31cde8c62611e421becd4648c7284d76c
|
|
||
| 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 1cf7b1b31cde8c62611e421becd4648c7284d76c
| @@ -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 1cf7b1b31cde8c62611e421becd4648c7284d76c
| @@ -30,47 +30,6 @@ class UpcastMXFPOpPattern : public ConvertOpToLLVMPattern<UpcastMXFPOp> { | |||
| : ConvertOpToLLVMPattern<UpcastMXFPOp>(typeConverter, benefit), | |||
| targetInfo(targetInfo) {} | |||
|
|
|||
| llvm::SmallVector<Value> unpackFP4Elements(Location loc, | |||
There was a problem hiding this comment.
Removed in 1cf7b1b31cde8c62611e421becd4648c7284d76c
etiotto
requested review from
chengjunlu,
a team,
Dewei-Wang-sh,
alexbaden,
LiyangLingIntel and
quintinwang5
Merging in #2707. |
Add initial support for the new tritongpu
upcast_mxfpoperation.