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 /// Patterns to transform vector ops into a canonical form to convert to MMA
20 /// matrix operations. If `useNvGpu` is true, then the patterns will populated
21 /// will prepare for conversion to `nvgpu` mma operations rather than the `gpu`
22 /// dialect WMMA operations.
23 void populatePrepareVectorToMMAPatterns(RewritePatternSet &patterns,
24                                         bool useNvGpu = false);
25 
26 /// Convert vector ops to MMA matrix operations nested under `rootOp`. This will
27 /// convert slice of operations that can be legally converted to MMA operations.
28 /// The rest of the vector operations are left untouched.
29 void convertVectorToMMAOps(Operation *rootOp);
30 
31 /// Convert vector ops ops nested under `rootOp` to vector and GPU operaitons
32 /// compatible with the `nvvm.mma.sync` lowering path. This will convert a slice
33 /// of operations that can be legally lowered on this path while the rest of
34 /// the vector operations are left untouched.
35 LogicalResult convertVectorToNVVMCompatibleMMASync(Operation *rootOp);
36 
37 /// Convert from vector to GPU ops.
38 std::unique_ptr<Pass> createConvertVectorToGPUPass(bool useNvGpu = false);
39 
40 } // namespace mlir
41 
42 #endif // MLIR_CONVERSION_VECTORTOGPU_VECTORTOGPU_H
43