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
24constexpr int kWarpSize = 32;
25
26/// M size of wgmma.mma_async instruction
27constexpr 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.
31constexpr 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.
35constexpr unsigned kMaxTMADimension = 256;
36/// Last dimension of 2D+ TMA must be 128 bytes
37constexpr 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

source code of mlir/include/mlir/Dialect/NVGPU/IR/NVGPUDialect.h