Use default FMA implementation of Triton upstream (#4680)

Signed-off-by: Anatoly Myachev <[email protected]>

Author: anmyachev

third_party/intel/lib/TritonIntelGPUToLLVM/CMakeLists.txt

@@ -5,7 +5,6 @@ add_triton_library(TritonIntelGPUToLLVM
     ControlFlowOpToLLVM.cpp
     ConvertLayoutOpToLLVM.cpp
     DotOpToLLVM/DPAS.cpp
-    DotOpToLLVM/FMA.cpp
     DotOpToLLVM.cpp
     ElementwiseOpToLLVM.cpp
     Fp4ToFpOpToLLVM.cpp

third_party/intel/lib/TritonIntelGPUToLLVM/DotOpToLLVM.cpp

@@ -7,11 +7,6 @@ using ::mlir::triton::gpu::getShapePerCTA;
 using ::mlir::triton::gpu::intel::DpasEncodingAttr;
 
 namespace fma_details {
-LogicalResult
-convertFMADot(triton::DotOp op, triton::DotOp::Adaptor adaptor,
-              const TritonIntelGPUToLLVMTypeConverter *typeConverter,
-              ConversionPatternRewriter &rewriter);
-
 LogicalResult convertDPAS(triton::DotOp op, triton::DotOp::Adaptor adaptor,
                           TritonIntelGPUToLLVMTypeConverter *typeConverter,
                           ConversionPatternRewriter &rewriter);
@@ -44,8 +39,7 @@ struct DotOpConversion : public ConvertTritonGPUOpToLLVMPattern<triton::DotOp> {
 
     if (isa<BlockedEncodingAttr>(
             cast<RankedTensorType>(D.getType()).getEncoding()))
-      return fma_details::convertFMADot(op, adaptor, getTypeConverter(),
-                                        rewriter);
+      return convertFMADot(op, adaptor, getTypeConverter(), rewriter);
 
     llvm::report_fatal_error(
         "Unsupported DotOp found when converting TritonGPU to LLVM.");