1 | //===- NVGPUDialect.h - MLIR Dialect for NVGPU ------------------*- 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 | // This file declares the Target dialect for NVGPU in MLIR. |
10 | // |
11 | //===----------------------------------------------------------------------===// |
12 | |
13 | #ifndef MLIR_DIALECT_NVGPU_NVGPUDIALECT_H_ |
14 | #define MLIR_DIALECT_NVGPU_NVGPUDIALECT_H_ |
15 | |
16 | #include "mlir/Bytecode/BytecodeOpInterface.h" |
17 | #include "mlir/IR/BuiltinTypes.h" |
18 | #include "mlir/IR/Dialect.h" |
19 | #include "mlir/IR/OpDefinition.h" |
20 | #include "mlir/Interfaces/SideEffectInterfaces.h" |
21 | |
22 | #include "mlir/Dialect/NVGPU/IR/NVGPUEnums.h.inc" |
23 | |
24 | constexpr int kWarpSize = 32; |
25 | |
26 | /// M size of wgmma.mma_async instruction |
27 | constexpr int kWgmmaSizeM = 64; |
28 | |
29 | /// Maximum TMA tile dimension (tensorRank) must be non-zero and less than or |
30 | /// equal to the maximum supported dimensionality of 5. |
31 | constexpr unsigned kMaxTMATensorDimension = 5; |
32 | /// Maximum TMA tile size (boxDim), which specifies number of elements |
33 | /// to be traversed along each of the kMaxTMATensorDimension (tensorRank) |
34 | /// dimensions, must be non-zero and less than or equal to 256. |
35 | constexpr unsigned kMaxTMADimension = 256; |
36 | /// Last dimension of 2D+ TMA must be 128 bytes |
37 | constexpr unsigned kMaxTMALastdimByte = 128; |
38 | |
39 | #define GET_ATTRDEF_CLASSES |
40 | #include "mlir/Dialect/NVGPU/IR/NVGPUAttrDefs.h.inc" |
41 | |
42 | #define GET_TYPEDEF_CLASSES |
43 | #include "mlir/Dialect/NVGPU/IR/NVGPUTypes.h.inc" |
44 | |
45 | #include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h.inc" |
46 | |
47 | #define GET_OP_CLASSES |
48 | #include "mlir/Dialect/NVGPU/IR/NVGPU.h.inc" |
49 | |
50 | #endif // MLIR_DIALECT_NVGPU_NVGPUDIALECT_H_ |
51 | |