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