[all-commits] [llvm/llvm-project] dd080c: [mlir][nvvm] Add NVVMToLLVM Pass

[Guray Ozen via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: dd080c7579e0c0e1d41dafb5dbea01343d6e0dc1
      https://github.com/llvm/llvm-project/commit/dd080c7579e0c0e1d41dafb5dbea01343d6e0dc1
  Author: Guray Ozen <guray.ozen at gmail.com>
  Date:   2023-07-11 (Tue, 11 Jul 2023)

  Changed paths:
    A mlir/include/mlir/Conversion/NVVMToLLVM/NVVMToLLVM.h
    M mlir/include/mlir/Conversion/Passes.h
    M mlir/include/mlir/Conversion/Passes.td
    M mlir/include/mlir/Dialect/LLVMIR/CMakeLists.txt
    M mlir/include/mlir/Dialect/LLVMIR/NVVMDialect.h
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/lib/Conversion/CMakeLists.txt
    A mlir/lib/Conversion/NVVMToLLVM/CMakeLists.txt
    A mlir/lib/Conversion/NVVMToLLVM/NVVMToLLVM.cpp
    A mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir

  Log Message:
  -----------
  [mlir][nvvm] Add NVVMToLLVM Pass

It introduces an NVVMToLLVM Pass and a `BasicPtxBuilderOpInterface` interface. The Pass performs pattern matching on all the NVVM Ops that implement the BasicPtxBuilderOpInterface interface to generate LLVM Inline Assembly Ops.

The BasicPtxBuilderOpInterface interface is utilized in the convert-nvvm-to-llvm pass, which lowers Ops that support this interface to inline assembly Ops. The interface provides several methods that are used for this lowering.

The `getPtx` method returns PTX code. The `hasSideEffect` method is used to determine whether the op has any side effects on the memory. The `hasIntrinsic` method indicates whether the operation has intrinsic support in LLVM. This is particularly useful for Ops that don't have intrinsic support for each case. The `getAsmValues` method returns the arguments to be passed to the PTX code. The order of arguments starts with the results and they are used for write operations, followed by the operands and attributes.

Example:

If we have the following Op definition that returns PTX code through getPtx:
```tablegen
def NVVM_MBarrierArriveExpectTxOp : NVVM_Op<\"mbarrier.arrive.expect_tx\",
                    [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]>,
  Results<(outs LLVM_Type:$res)>, Arguments<(ins LLVM_i64ptr_any:$addr, I32:$txcount)> {
  ...
  let extraClassDefinition = [{
    const char* $cppClass::getPtx() { return \"mbarrier.arrive.expect_tx.b64 %0, [%1], %2;\"; }
  }\];
}
```

The NVVM Op will look like below:
```mlir
  %0 = nvvm.mbarrier.arrive.expect_tx %barrier, %txcount : !llvm.ptr, i32 -> i32
```

The `convert-nvvm-to-llvm` Pass generates the following PTX code, while keeping the order of arguments the same. The read/write modifiers are set based on the input and result types.
```mlir
  %0 = llvm.inline_asm has_side_effects asm_dialect = att "mbarrier.arrive.expect_tx.b64 %0, [%1], %2;", "=r,l,r" %arg0, %arg1 : (!llvm.ptr, i32) -> i32
```

Reviewed By: nicolasvasilache

Differential Revision: https://reviews.llvm.org/D154060