| 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 | |
|---|
| 17 | #include "mlir/Bytecode/BytecodeOpInterface.h" |
|---|
| 18 | #include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.h" |
|---|
| 19 | #include "mlir/Dialect/LLVMIR/LLVMDialect.h" |
|---|
| 20 | #include "mlir/IR/Dialect.h" |
|---|
| 21 | #include "mlir/IR/OpDefinition.h" |
|---|
| 22 | #include "mlir/Interfaces/SideEffectInterfaces.h" |
|---|
| 23 | #include "llvm/IR/IntrinsicsNVPTX.h" |
|---|
| 24 | |
|---|
| 25 | #include "mlir/Dialect/LLVMIR/NVVMOpsEnums.h.inc" |
|---|
| 26 | |
|---|
| 27 | namespace mlir { |
|---|
| 28 | namespace NVVM { |
|---|
| 29 | |
|---|
| 30 | // Shared memory has 128-bit alignment |
|---|
| 31 | constexpr int kSharedMemoryAlignmentBit = 128; |
|---|
| 32 | |
|---|
| 33 | /// NVVM memory space identifiers. |
|---|
| 34 | enum NVVMMemorySpace { |
|---|
| 35 | /// Global memory space identifier. |
|---|
| 36 | kGlobalMemorySpace = 1, |
|---|
| 37 | /// Shared memory space identifier. |
|---|
| 38 | kSharedMemorySpace = 3 |
|---|
| 39 | }; |
|---|
| 40 | |
|---|
| 41 | /// Return the element type and number of elements associated with a wmma matrix |
|---|
| 42 | /// of given chracteristics. This matches the logic in IntrinsicsNVVM.td |
|---|
| 43 | /// WMMA_REGS structure. |
|---|
| 44 | std::pair<mlir::Type, unsigned> inferMMAType(mlir::NVVM::MMATypes type, |
|---|
| 45 | mlir::NVVM::MMAFrag frag, int nRow, |
|---|
| 46 | int nCol, |
|---|
| 47 | mlir::MLIRContext *context); |
|---|
| 48 | } // namespace NVVM |
|---|
| 49 | } // namespace mlir |
|---|
| 50 | |
|---|
| 51 | ///// Ops ///// |
|---|
| 52 | #define GET_ATTRDEF_CLASSES |
|---|
| 53 | #include "mlir/Dialect/LLVMIR/NVVMOpsAttributes.h.inc" |
|---|
| 54 | |
|---|
| 55 | #define GET_OP_CLASSES |
|---|
| 56 | #include "mlir/Dialect/LLVMIR/NVVMOps.h.inc" |
|---|
| 57 | |
|---|
| 58 | #include "mlir/Dialect/LLVMIR/NVVMOpsDialect.h.inc" |
|---|
| 59 | |
|---|
| 60 | #endif /* MLIR_DIALECT_LLVMIR_NVVMDIALECT_H_ */ |
|---|
| 61 | |
|---|
