MLIR  21.0.0git
NVVMDialect.h
Go to the documentation of this file.
1 //===- NVVMDialect.h - MLIR NVVM IR 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 // This file defines the NVVM IR dialect in MLIR, containing NVVM operations and
10 // NVVM specific extensions to the LLVM type system.
11 //
12 //===----------------------------------------------------------------------===//
13 
14 #ifndef MLIR_DIALECT_LLVMIR_NVVMDIALECT_H_
15 #define MLIR_DIALECT_LLVMIR_NVVMDIALECT_H_
16 
20 #include "mlir/IR/Dialect.h"
21 #include "mlir/IR/OpDefinition.h"
25 #include "llvm/IR/IntrinsicsNVPTX.h"
26 
27 #include "mlir/Dialect/LLVMIR/NVVMOpsEnums.h.inc"
28 
29 namespace mlir {
30 namespace NVVM {
31 
32 // Shared memory has 128-bit alignment
33 constexpr int kSharedMemoryAlignmentBit = 128;
34 
35 /// NVVM memory space identifiers.
37  /// Global memory space identifier.
39  /// Shared memory space identifier.
41  /// Constant memory space identifier.
43  /// Tensor memory space identifier.
44  /// Tensor memory is available only in arch-accelerated
45  /// variants from sm100 onwards.
47  /// Distributed shared memory space identifier.
48  /// Distributed shared memory is available only in sm90+.
50 };
51 
52 /// Return the element type and number of elements associated with a wmma matrix
53 /// of given chracteristics. This matches the logic in IntrinsicsNVVM.td
54 /// WMMA_REGS structure.
55 std::pair<mlir::Type, unsigned> inferMMAType(mlir::NVVM::MMATypes type,
56  mlir::NVVM::MMAFrag frag, int nRow,
57  int nCol,
58  mlir::MLIRContext *context);
59 } // namespace NVVM
60 } // namespace mlir
61 
62 ///// Ops /////
63 #define GET_ATTRDEF_CLASSES
64 #include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.h.inc"
65 
66 #define GET_OP_CLASSES
67 #include "mlir/Dialect/LLVMIR/NVVMOps.h.inc"
68 
69 #include "mlir/Dialect/LLVMIR/NVVMOpsDialect.h.inc"
70 
71 #endif /* MLIR_DIALECT_LLVMIR_NVVMDIALECT_H_ */
MLIRContext is the top-level object for a collection of MLIR operations.
Definition: MLIRContext.h:60
constexpr int kSharedMemoryAlignmentBit
Definition: NVVMDialect.h:33
NVVMMemorySpace
NVVM memory space identifiers.
Definition: NVVMDialect.h:36
@ kGlobalMemorySpace
Global memory space identifier.
Definition: NVVMDialect.h:38
@ kConstantMemorySpace
Constant memory space identifier.
Definition: NVVMDialect.h:42
@ kSharedClusterMemorySpace
Distributed shared memory space identifier.
Definition: NVVMDialect.h:49
@ kTensorMemorySpace
Tensor memory space identifier.
Definition: NVVMDialect.h:46
@ kSharedMemorySpace
Shared memory space identifier.
Definition: NVVMDialect.h:40
std::pair< mlir::Type, unsigned > inferMMAType(mlir::NVVM::MMATypes type, mlir::NVVM::MMAFrag frag, int nRow, int nCol, mlir::MLIRContext *context)
Return the element type and number of elements associated with a wmma matrix of given chracteristics.
Include the generated interface declarations.