[all-commits] [llvm/llvm-project] cc798f: [LLVM][NVPTX] Add Rubin extensions for G2S Tensor ...

[Rajat Bajpai via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: cc798fb63ca1a08f43cdcd314b66a96fa973e584
      https://github.com/llvm/llvm-project/commit/cc798fb63ca1a08f43cdcd314b66a96fa973e584
  Author: Rajat Bajpai <rbajpai at nvidia.com>
  Date:   2026-09-08 (Tue, 08 Sep 2026)

  Changed paths:
    M llvm/docs/NVPTXUsage.md
    M llvm/include/llvm/IR/IntrinsicsNVVM.td
    M llvm/include/llvm/IR/NVVMIntrinsicUtils.h
    M llvm/lib/IR/AutoUpgrade.cpp
    M llvm/lib/IR/NVVMIntrinsicUtils.cpp
    M llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
    M llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
    M llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
    M llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-1cta.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-2cta.ll
    A llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-cta-override.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-cta-sm100.ll
    A llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-cta-validate-data.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-gather4.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-im2colw.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-im2colw128.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-invalid.ll
    A llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-mc32.ll
    A llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-override.ll
    A llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s-validate-data.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-g2s.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g-im2colw.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g-scatter4.ll
    M llvm/test/CodeGen/NVPTX/cp-async-bulk-tensor-s2g.ll
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
    M mlir/test/Target/LLVMIR/nvvm/tma_load_cluster_im2col.mlir
    M mlir/test/Target/LLVMIR/nvvm/tma_load_cluster_tile.mlir
    M mlir/test/Target/LLVMIR/nvvm/tma_load_cta_im2col.mlir
    M mlir/test/Target/LLVMIR/nvvm/tma_load_cta_tile.mlir

  Log Message:
  -----------
  [LLVM][NVPTX] Add Rubin extensions for G2S Tensor intrinsics (#220029)

This change adds support for the Rubin-extended Global-to-Shared tensor
copy intrinsics for both the Cluster and CTA variants:

1. 32-bit multicast mask support: the G2S Cluster intrinsics are now
overloaded on the multicast mask type (i16/i32); an i32 mask lowers to
`.multicast::cluster::32b`.
2. Data-validity reporting through the mbarrier: all G2S Cluster and CTA
intrinsics, including the property-override forms, take a trailing `i32
%validate_pattern` immediate (range [0, 6)) that selects a
`.mbarrier::report::validity` qualifier. It defaults to 0 (disabled).
3. Override tensor map properties: the G2S Cluster and CTA intrinsics
now support `.override::global_address`, optionally followed by
`.override::global_dim` / `.override::global_dim_stride`, across all
load modes. The dim/stride overrides apply to tile mode only; im2col and
gather4 support the global-address override only.

The new `validate_pattern` argument changes the IR signatures of all G2S
intrinsics; legacy declarations are auto-upgraded with the missing
trailing arguments defaulted to 0, and the multicast mask type selects
the overload. The MLIR NVVM `CpAsyncBulkTensorGlobalToSharedClusterOp`
is updated to build the new signature.



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