1 //===- VectorToGPU.h - Convert vector to GPU dialect ------------*- C++ -*-===// 2 // 3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4 // See https://llvm.org/LICENSE.txt for license information. 5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6 // 7 //===----------------------------------------------------------------------===// 8 9 #ifndef MLIR_CONVERSION_VECTORTOGPU_VECTORTOGPU_H 10 #define MLIR_CONVERSION_VECTORTOGPU_VECTORTOGPU_H 11 12 #include "mlir/IR/PatternMatch.h" 13 14 namespace mlir { 15 class MLIRContext; 16 class Pass; 17 class RewritePatternSet; 18 19 #define GEN_PASS_DECL_CONVERTVECTORTOGPU 20 #include "mlir/Conversion/Passes.h.inc" 21 22 /// Patterns to transform vector ops into a canonical form to convert to MMA 23 /// matrix operations. If `useNvGpu` is true, then the patterns will populated 24 /// will prepare for conversion to `nvgpu` mma operations rather than the `gpu` 25 /// dialect WMMA operations. 26 void populatePrepareVectorToMMAPatterns(RewritePatternSet &patterns, 27 bool useNvGpu = false); 28 29 /// Convert vector ops to MMA matrix operations nested under `rootOp`. This will 30 /// convert slice of operations that can be legally converted to MMA operations. 31 /// The rest of the vector operations are left untouched. 32 LogicalResult convertVectorToMMAOps(RewriterBase &rewriter, Operation *rootOp); 33 34 /// Convert vector ops ops nested under `rootOp` to vector and GPU operaitons 35 /// compatible with the `nvvm.mma.sync` lowering path. This will convert a slice 36 /// of operations that can be legally lowered on this path while the rest of 37 /// the vector operations are left untouched. 38 LogicalResult convertVectorToNVVMCompatibleMMASync(RewriterBase &rewriter, 39 Operation *rootOp); 40 41 /// Convert from vector to GPU ops. 42 std::unique_ptr<Pass> createConvertVectorToGPUPass(bool useNvGpu = false); 43 44 } // namespace mlir 45 46 #endif // MLIR_CONVERSION_VECTORTOGPU_VECTORTOGPU_H 47