[NVIDIA] Lower tensor-map proxy fences with NVVM

[ajit283]

Summary

Replace the remaining supported compiler-generated inline NVIDIA fence with NVVM dialect operations.

The tensor-map acquire fence currently emits fence.proxy.tensormap::generic.acquire.gpu through inline PTX. Keeping this operation in NVVM exposes its synchronization semantics to MLIR and delegates PTX spelling and target validation to LLVM's NVPTX backend.

Details

• Replace inline fence.proxy.tensormap::generic.acquire.gpu with nvvm.fence.proxy.acquire.
• Preserve first-warp predication with explicit LLVM control flow because the NVVM fence intrinsic is not predicated.
• Represent the fence's generic address explicitly with an address-space cast from the global tensor-map pointer.
• Preserve the existing inline cp.async.bulk PTXAS workaround unchanged.
• Preserve the final CTA synchronization barrier.

Other compiler-generated fence paths already use structured lowering:

• async/shared proxy fences use nvvm.fence.proxy;
• the mbarrier initialization fence is handled by [NVIDIA] Lower SM90+ mbarriers with NVVM #10937;
• atomic thread fences are represented by LLVM atomic ordering;
• Triton does not currently emit an inline fence.sc.cluster path.

tensormap.cp_fenceproxy remains inline because the pinned LLVM revision has no corresponding NVVM operation or LLVM intrinsic. It is a tensor-map copy instruction and requires LLVM support before it can be migrated.

At Triton's pinned LLVM revision, the NVPTX backend maps the new operation to the same fence.proxy.tensormap::generic.acquire.gpu PTX instruction.

Testing

• pre-commit run --from-ref origin/main --to-ref HEAD
• make
• lit -v test/Conversion (81/81 passed)
• Lowered the tensor-map fence through llc and assembled it with CUDA 13.3 ptxas for SM90 and SM100a.

New contributor declaration

• I am not making a trivial change, such as fixing a typo in a comment.

• I have written a PR description following these
  rules.

• I have run pre-commit run --from-ref origin/main --to-ref HEAD.

• I have added tests.

  • /test for lit tests
  • /unittest for C++ tests
  • /python/test for end-to-end tests

• I have not added any lit tests.