[llvm] [NVPTX] Add elect.sync Intrinsic (PR #104780)

[via llvm-commits]

================
@@ -0,0 +1,42 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx80 | FileCheck %s
+; RUN: %if ptxas-12.0 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx80 | %ptxas-verify -arch=sm_90 %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+declare {i32, i1} @llvm.nvvm.elect.sync(i32)
+
+define {i32, i1} @elect_sync(i32 %mask) {
+; CHECK-LABEL: elect_sync(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b16 %rs<2>;
+; CHECK-NEXT:    .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.u32 %r1, [elect_sync_param_0];
+; CHECK-NEXT:    elect.sync %r2|%p1, %r1;
+; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r2;
+; CHECK-NEXT:    selp.u16 %rs1, -1, 0, %p1;
+; CHECK-NEXT:    st.param.b8 [func_retval0+4], %rs1;
+; CHECK-NEXT:    ret;
+  %val = call {i32, i1} @llvm.nvvm.elect.sync(i32 %mask)
+  ret {i32, i1} %val
+}
+
+define {i32, i1} @elect_sync_imm() {
+; CHECK-LABEL: elect_sync_imm(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b16 %rs<2>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    elect.sync %r1|%p1, -1;
+; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r1;
+; CHECK-NEXT:    selp.u16 %rs1, -1, 0, %p1;
+; CHECK-NEXT:    st.param.b8 [func_retval0+4], %rs1;
+; CHECK-NEXT:    ret;
+  %val = call {i32, i1} @llvm.nvvm.elect.sync(i32 u0xffffffff)
+  ret {i32, i1} %val
+}
----------------
gonzalobg wrote:

It would be nice to have a test with two `elect.sync` back to back on the same mask, that CHECKs that the second one is present. If we ever decide to optimize this further, that test may need updating.

https://github.com/llvm/llvm-project/pull/104780