MLIR  22.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 
23 #include "mlir/IR/Dialect.h"
24 #include "mlir/IR/OpDefinition.h"
28 #include "llvm/IR/IntrinsicsNVPTX.h"
29 
30 #include "mlir/Dialect/LLVMIR/NVVMOpsEnums.h.inc"
31 
32 namespace mlir {
33 namespace NVVM {
34 /// Utility functions to compare NVVMMemorySpace with unsigned values.
35 inline bool operator==(unsigned as, NVVMMemorySpace memSpace) {
36  return as == static_cast<unsigned>(memSpace);
37 }
38 inline bool operator==(NVVMMemorySpace memSpace, unsigned as) {
39  return static_cast<unsigned>(memSpace) == as;
40 }
41 inline bool operator!=(unsigned as, NVVMMemorySpace memSpace) {
42  return as != static_cast<unsigned>(memSpace);
43 }
44 inline bool operator!=(NVVMMemorySpace memSpace, unsigned as) {
45  return static_cast<unsigned>(memSpace) != as;
46 }
47 
48 // Shared memory has 128-bit alignment
49 constexpr int kSharedMemoryAlignmentBit = 128;
50 
51 /// A pair type of LLVM's Intrinsic ID and args (which are llvm values).
52 /// This type is returned by the getIntrinsicIDAndArgs() methods.
53 using IDArgPair =
54  std::pair<llvm::Intrinsic::ID, llvm::SmallVector<llvm::Value *>>;
55 
56 /// Return the element type and number of elements associated with a wmma matrix
57 /// of given chracteristics. This matches the logic in IntrinsicsNVVM.td
58 /// WMMA_REGS structure.
59 std::pair<mlir::Type, unsigned> inferMMAType(mlir::NVVM::MMATypes type,
60  mlir::NVVM::MMAFrag frag, int nRow,
61  int nCol,
62  mlir::MLIRContext *context);
63 } // namespace NVVM
64 } // namespace mlir
65 
66 ///// Ops /////
67 #define GET_ATTRDEF_CLASSES
68 #include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.h.inc"
69 
70 #define GET_OP_CLASSES
71 #include "mlir/Dialect/LLVMIR/NVVMOps.h.inc"
72 
73 #include "mlir/Dialect/LLVMIR/NVVMOpsDialect.h.inc"
74 
75 #endif /* MLIR_DIALECT_LLVMIR_NVVMDIALECT_H_ */
MLIRContext is the top-level object for a collection of MLIR operations.
Definition: MLIRContext.h:63
std::pair< llvm::Intrinsic::ID, llvm::SmallVector< llvm::Value * > > IDArgPair
A pair type of LLVM's Intrinsic ID and args (which are llvm values).
Definition: NVVMDialect.h:54
constexpr int kSharedMemoryAlignmentBit
Definition: NVVMDialect.h:49
bool operator==(unsigned as, NVVMMemorySpace memSpace)
Utility functions to compare NVVMMemorySpace with unsigned values.
Definition: NVVMDialect.h:35
bool operator!=(unsigned as, NVVMMemorySpace memSpace)
Definition: NVVMDialect.h:41
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.