github.com/johnnyeven/libtools@v0.0.0-20191126065708-61829c1adf46/third_party/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (about)

     1  //===- NVVMDialect.cpp - NVVM IR Ops and Dialect registration -------------===//
     2  //
     3  // Copyright 2019 The MLIR Authors.
     4  //
     5  // Licensed under the Apache License, Version 2.0 (the "License");
     6  // you may not use this file except in compliance with the License.
     7  // You may obtain a copy of the License at
     8  //
     9  //   http://www.apache.org/licenses/LICENSE-2.0
    10  //
    11  // Unless required by applicable law or agreed to in writing, software
    12  // distributed under the License is distributed on an "AS IS" BASIS,
    13  // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
    14  // See the License for the specific language governing permissions and
    15  // limitations under the License.
    16  // =============================================================================
    17  //
    18  // This file defines the types and operation details for the NVVM IR dialect in
    19  // MLIR, and the LLVM IR dialect.  It also registers the dialect.
    20  //
    21  // The NVVM dialect only contains GPU specific additions on top of the general
    22  // LLVM dialect.
    23  //
    24  //===----------------------------------------------------------------------===//
    25  
    26  #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
    27  
    28  #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
    29  #include "mlir/IR/Builders.h"
    30  #include "mlir/IR/MLIRContext.h"
    31  #include "mlir/IR/Operation.h"
    32  #include "mlir/IR/StandardTypes.h"
    33  #include "llvm/AsmParser/Parser.h"
    34  #include "llvm/IR/Attributes.h"
    35  #include "llvm/IR/Function.h"
    36  #include "llvm/IR/Type.h"
    37  #include "llvm/Support/SourceMgr.h"
    38  
    39  namespace mlir {
    40  namespace NVVM {
    41  
    42  //===----------------------------------------------------------------------===//
    43  // Printing/parsing for NVVM ops
    44  //===----------------------------------------------------------------------===//
    45  
    46  static void printNVVMIntrinsicOp(OpAsmPrinter *p, Operation *op) {
    47    *p << op->getName() << " ";
    48    p->printOperands(op->getOperands());
    49    if (op->getNumResults() > 0)
    50      interleaveComma(op->getResultTypes(), *p << " : ");
    51  }
    52  
    53  // <operation> ::= `llvm.nvvm.XYZ` : type
    54  static ParseResult parseNVVMSpecialRegisterOp(OpAsmParser *parser,
    55                                                OperationState *result) {
    56    Type type;
    57    if (parser->parseOptionalAttributeDict(result->attributes) ||
    58        parser->parseColonType(type))
    59      return failure();
    60  
    61    result->addTypes(type);
    62    return success();
    63  }
    64  
    65  static LLVM::LLVMDialect *getLlvmDialect(OpAsmParser *parser) {
    66    return parser->getBuilder()
    67        .getContext()
    68        ->getRegisteredDialect<LLVM::LLVMDialect>();
    69  }
    70  
    71  // <operation> ::=
    72  //     `llvm.nvvm.shfl.sync.bfly %dst, %val, %offset, %clamp_and_mask`
    73  //     : result_type
    74  static ParseResult parseNVVMShflSyncBflyOp(OpAsmParser *parser,
    75                                             OperationState *result) {
    76    auto llvmDialect = getLlvmDialect(parser);
    77    auto int32Ty = LLVM::LLVMType::getInt32Ty(llvmDialect);
    78  
    79    SmallVector<OpAsmParser::OperandType, 8> ops;
    80    Type type;
    81    return failure(parser->parseOperandList(ops) ||
    82                   parser->parseOptionalAttributeDict(result->attributes) ||
    83                   parser->parseColonType(type) ||
    84                   parser->addTypeToList(type, result->types) ||
    85                   parser->resolveOperands(ops, {int32Ty, type, int32Ty, int32Ty},
    86                                           parser->getNameLoc(),
    87                                           result->operands));
    88  }
    89  
    90  // <operation> ::= `llvm.nvvm.vote.ballot.sync %mask, %pred` : result_type
    91  static ParseResult parseNVVMVoteBallotOp(OpAsmParser *parser,
    92                                           OperationState *result) {
    93    auto llvmDialect = getLlvmDialect(parser);
    94    auto int32Ty = LLVM::LLVMType::getInt32Ty(llvmDialect);
    95    auto int1Ty = LLVM::LLVMType::getInt1Ty(llvmDialect);
    96  
    97    SmallVector<OpAsmParser::OperandType, 8> ops;
    98    Type type;
    99    return failure(parser->parseOperandList(ops) ||
   100                   parser->parseOptionalAttributeDict(result->attributes) ||
   101                   parser->parseColonType(type) ||
   102                   parser->addTypeToList(type, result->types) ||
   103                   parser->resolveOperands(ops, {int32Ty, int1Ty},
   104                                           parser->getNameLoc(),
   105                                           result->operands));
   106  }
   107  
   108  //===----------------------------------------------------------------------===//
   109  // NVVMDialect initialization, type parsing, and registration.
   110  //===----------------------------------------------------------------------===//
   111  
   112  // TODO(herhut): This should be the llvm.nvvm dialect once this is supported.
   113  NVVMDialect::NVVMDialect(MLIRContext *context) : Dialect("nvvm", context) {
   114    addOperations<
   115  #define GET_OP_LIST
   116  #include "mlir/Dialect/LLVMIR/NVVMOps.cpp.inc"
   117        >();
   118  
   119    // Support unknown operations because not all NVVM operations are registered.
   120    allowUnknownOperations();
   121  }
   122  
   123  #define GET_OP_CLASSES
   124  #include "mlir/Dialect/LLVMIR/NVVMOps.cpp.inc"
   125  
   126  static DialectRegistration<NVVMDialect> nvvmDialect;
   127  
   128  } // namespace NVVM
   129  } // namespace mlir