Implement scaled_dot(mxfp8, fp8) via mma #4795
Conversation
Initial implementation using mma. Missing to test that it plays ball with the pipeliner.
|
|
||
| let results = (outs TT_FloatTensor:$d); | ||
|
|
||
| // Not sure why I need to fully specify the optional group, but otherwise it complains when loading the mlir file |
There was a problem hiding this comment.
The one referring to the scale of the rhs being an std::optional
There was a problem hiding this comment.
I guess I still don't understand the meaning of the comment :) What does "I need to fully specify the optional group" means then?
There was a problem hiding this comment.
I meant that I had to specify all this funny syntax: (`,`) : (`,` $rhs_scale^ `,`)?. In the other ops that we use optional args, we pretty much handle them in the frontend where we call different builders explicitly.
Here, I tried to handle the optional arg at the MLIR level, and it so happened that just writing (`,` $rhs_scale^)? wouldn't work. I had to write the whole if-else clause for the optional, adding another comma inside it.
|
|
||
| namespace mlir::triton::gpu { | ||
|
|
||
| LogicalResult UpcastMXFPOp::verify() { |
There was a problem hiding this comment.
nit: The verifiers and other methods for this dialect are currently in Dialect.cpp. I would keep it there for consistency. We can have a separate PR to move it for cleanup if you want.
No, it's just that I haven't tested if it works as expected just yet. |
|
BTW, it would be nice to have lit tests for accelerate_matmul as well lowering to of upscale to llvm |
|
@ThomasRaoux addressed the review (at long last) |
Initial implementation using mma. Missing to test that it plays ball with the pipeliner.
Initial implementation using mma. Missing to test that it plays ball with the pipeliner.
Initial implementation using mma.
Missing to test that it plays ball with the pipeliner.