[NVIDIA] Lower warp-specialization barriers with NVVM

[ajit283]

Summary

Use the NVVM dialect for the remaining NVIDIA non-mbarrier barrier path instead of manually constructing an LLVM intrinsic call.

Warp-specialization switch loops currently emit llvm.nvvm.barrier.cta.sync.all by name. Emitting nvvm.barrier keeps the target operation represented in MLIR, uses the dialect verifier, and delegates LLVM intrinsic selection to NVVM translation.

CTA barriers, warp barriers, and cluster arrive/wait barriers elsewhere already use NVVM dialect operations. This change makes the warp-specialization all-participant named barriers consistent with them.

At Triton's pinned LLVM revision, nvvm.barrier translates to the same llvm.nvvm.barrier.cta.sync.all intrinsic and NVPTX backend instruction, so the emitted bar.sync semantics are unchanged.

Testing

• pre-commit run --from-ref origin/main --to-ref HEAD
• make
• lit -v test/Conversion (81/81 passed)
• pytest -s --tb=short python/test/unit/language/test_compile_only.py::test_compile_only_dot

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.