[all-commits] [llvm/llvm-project] 40d005: [NVPTX] Add TMA bulk tensor reduction intrinsics (...

[Durgadoss R via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 40d0058e6ad16d4716192a0e331ac58a6ed9def3
      https://github.com/llvm/llvm-project/commit/40d0058e6ad16d4716192a0e331ac58a6ed9def3
  Author: Durgadoss R <durgadossr at nvidia.com>
  Date:   2024-11-27 (Wed, 27 Nov 2024)

  Changed paths:
    M llvm/docs/NVPTXUsage.rst
    M llvm/include/llvm/IR/IntrinsicsNVVM.td
    A llvm/include/llvm/IR/NVVMIntrinsicFlags.h
    M llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
    M llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
    M llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
    M llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
    M llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
    A llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-reduce.ll

  Log Message:
  -----------
  [NVPTX] Add TMA bulk tensor reduction intrinsics (#116854)

This patch adds NVVM intrinsics and NVPTX codegen for:
* cp.async.bulk.tensor.reduce.1D -> 5D variants, supporting both Tile
   and Im2Col modes.
* These intrinsics optionally support cache_hints as indicated by the
   boolean flag argument.
* Lit tests are added for all combinations of these intrinsics in
   cp-async-bulk-tensor-reduce.ll.
* The generated PTX is verified with a 12.3 ptxas executable.
* Added docs for these intrinsics in NVPTXUsage.rst file.

PTX Spec reference:
https://docs.nvidia.com/cuda/parallel-thread-execution/#data-movement-and-conversion-instructions-cp-reduce-async-bulk-tensor

Signed-off-by: Durgadoss R <durgadossr at nvidia.com>



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications