[all-commits] [llvm/llvm-project] e56d67: [mlir][nvgpu] Add `tma.create.descriptor` to creat...

[Guray Ozen via All-commits]

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

  Changed paths:
    M mlir/include/mlir/Conversion/GPUCommon/GPUCommonPass.h
    M mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
    M mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp
    M mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
    M mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
    M mlir/lib/ExecutionEngine/CudaRuntimeWrappers.cpp
    M mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir

  Log Message:
  -----------
  [mlir][nvgpu] Add `tma.create.descriptor` to create tensor map descriptor

The Op creates a tensor map descriptor object representing tiled memory region. The descriptor is used by Tensor Memory Access (TMA). The `tensor` is the source tensor to be tiled. The `boxDimensions` is the size of the tiled memory region in each dimension.

The pattern here lowers `tma.create.descriptor` to a runtime function call that eventually calls calls CUDA Driver's `cuTensorMapEncodeTiled`. For more information see below:
https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TENSOR__MEMORY.html

Depends on D155453

Reviewed By: nicolasvasilache

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