[all-commits] [llvm/llvm-project] 2c5739: [mlir][nvgpu] Implement `nvgpu.device_async_copy` ...

[Guray Ozen via All-commits]

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

  Changed paths:
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
    M mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
    M mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
    M mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir
    M mlir/test/Conversion/NVGPUToNVVM/typed-pointers.mlir
    M mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
    M mlir/test/Dialect/LLVMIR/invalid-typed-pointers.mlir
    M mlir/test/Dialect/LLVMIR/invalid.mlir
    M mlir/test/Dialect/LLVMIR/nvvm-typed-pointers.mlir
    M mlir/test/Dialect/LLVMIR/nvvm.mlir
    M mlir/test/Dialect/NVGPU/invalid.mlir
    M mlir/test/Target/LLVMIR/nvvmir.mlir

  Log Message:
  -----------
  [mlir][nvgpu] Implement `nvgpu.device_async_copy` by NVVMToLLVM Pass

`nvgpu.device_async_copy` is lowered into `cp.async` PTX instruction. However, NVPTX backend does not support its all mode especially when zero padding is needed. Therefore, current MLIR implementation genereates inline assembly for that.

This work simplifies PTX generation for `nvgpu.device_async_copy`, and implements it by `NVVMToLLVM` Pass.

Depends on D154060

Reviewed By: nicolasvasilache, manishucsd

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