[all-commits] [llvm/llvm-project] c20428: [MLIR][NVVM] Fix the lowering of legacy mbar.arriv...

[Dhairyashil R G via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: c20428cbb5f551d54520d492b1d27011e4076f5b
      https://github.com/llvm/llvm-project/commit/c20428cbb5f551d54520d492b1d27011e4076f5b
  Author: Dhairyashil R G <dhairyashil25 at gmail.com>
  Date:   2026-09-16 (Wed, 16 Sep 2026)

  Changed paths:
    M mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
    M mlir/test/Target/LLVMIR/nvvm/mbar_arrive_drop.mlir

  Log Message:
  -----------
  [MLIR][NVVM] Fix the lowering of legacy mbar.arrive_drop (#222916)

`nvvm.mbarrier.arrive_drop` fails to compile for sm_80:

LLVM ERROR: Cannot select: intrinsic
%llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cta

The op always lowers to the `scope.*.space.*` intrinsics, which need
sm_90. The legacy `llvm.nvvm.mbarrier.arrive.drop{,.shared}` intrinsics
work from sm_80 but are never used.

This uses the legacy intrinsic for the basic case only, space=cta,
scope=cta, no relaxed, no explicit count, the same fix as #172476 for
`mbarrier.arrive`. The count forms need sm_90, so they are unchanged.

Assisted-by: Claude



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