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 };
48 
49 /// Return the element type and number of elements associated with a wmma matrix
50 /// of given chracteristics. This matches the logic in IntrinsicsNVVM.td
51 /// WMMA_REGS structure.
52 std::pair<mlir::Type, unsigned> inferMMAType(mlir::NVVM::MMATypes type,
53  mlir::NVVM::MMAFrag frag, int nRow,
54  int nCol,
55  mlir::MLIRContext *context);
56 } // namespace NVVM
57 } // namespace mlir
58 
59 ///// Ops /////
60 #define GET_ATTRDEF_CLASSES
61 #include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.h.inc"
62 
63 #define GET_OP_CLASSES
64 #include "mlir/Dialect/LLVMIR/NVVMOps.h.inc"
65 
66 #include "mlir/Dialect/LLVMIR/NVVMOpsDialect.h.inc"
67 
68 #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
@ 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.