[clang] [llvm] cuda clang: Fix argument order for __reduce_max_sync (PR #132881)

[Austin Schuh via llvm-commits]

https://github.com/AustinSchuh updated https://github.com/llvm/llvm-project/pull/132881

>From 381ba9f21b4c2a2c4028143b84e9f71c8a20692f Mon Sep 17 00:00:00 2001
From: Austin Schuh <austin.linux at gmail.com>
Date: Mon, 24 Mar 2025 21:42:45 -0700
Subject: [PATCH 1/2] cuda clang: Fix argument order for __reduce_max_sync

The following cuda kernel would crash with an "an illegal instruction
was encountered" message.

__global__ void testcode(const float* data, unsigned *max_value) {
    unsigned r = static_cast<unsigned>(data[threadIdx.x]);

    const unsigned mask = __ballot_sync(0xFFFFFFFF, true);

    unsigned mx = __reduce_max_sync(mask, r);
    atomicMax(max_value, mx);
}

Digging into the ptx from both nvcc and clang, I discovered that the
arguments for the mask and value were swapped.  This swaps them back.

Fixes: https://github.com/llvm/llvm-project/issues/131415
Signed-off-by: Austin Schuh <austin.linux at gmail.com>
---
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index b2e05a567b4fe..1943e94c3ee7a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -315,7 +315,7 @@ defm MATCH_ALLP_SYNC_64 : MATCH_ALLP_SYNC<Int64Regs, "b64", int_nvvm_match_all_s
 multiclass REDUX_SYNC<string BinOp, string PTXType, Intrinsic Intrin> {
   def : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$src, Int32Regs:$mask),
           "redux.sync." # BinOp # "." # PTXType # " $dst, $src, $mask;",
-          [(set i32:$dst, (Intrin i32:$src, Int32Regs:$mask))]>,
+          [(set i32:$dst, (Intrin i32:$mask, Int32Regs:$src))]>,
         Requires<[hasPTX<70>, hasSM<80>]>;
 }
 

>From aa9148ce6af19a799f37aea6d973087c8fb7c53c Mon Sep 17 00:00:00 2001
From: Austin Schuh <austin.linux at gmail.com>
Date: Tue, 25 Mar 2025 12:12:47 -0700
Subject: [PATCH 2/2] Swap in __clang_cuda_intrinsics.h instead

---
 clang/lib/Headers/__clang_cuda_intrinsics.h | 16 ++++++++--------
 llvm/lib/Target/NVPTX/NVPTXIntrinsics.td    |  2 +-
 2 files changed, 9 insertions(+), 9 deletions(-)

diff --git a/clang/lib/Headers/__clang_cuda_intrinsics.h b/clang/lib/Headers/__clang_cuda_intrinsics.h
index a04e8b6de44d0..8b230af6f6647 100644
--- a/clang/lib/Headers/__clang_cuda_intrinsics.h
+++ b/clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -515,32 +515,32 @@ __device__ inline cuuint32_t __nvvm_get_smem_pointer(void *__ptr) {
 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 800
 __device__ inline unsigned __reduce_add_sync(unsigned __mask,
                                              unsigned __value) {
-  return __nvvm_redux_sync_add(__mask, __value);
+  return __nvvm_redux_sync_add(__value, __mask);
 }
 __device__ inline unsigned __reduce_min_sync(unsigned __mask,
                                              unsigned __value) {
-  return __nvvm_redux_sync_umin(__mask, __value);
+  return __nvvm_redux_sync_umin(__value, __mask);
 }
 __device__ inline unsigned __reduce_max_sync(unsigned __mask,
                                              unsigned __value) {
-  return __nvvm_redux_sync_umax(__mask, __value);
+  return __nvvm_redux_sync_umax(__value, __mask);
 }
 __device__ inline int __reduce_min_sync(unsigned __mask, int __value) {
-  return __nvvm_redux_sync_min(__mask, __value);
+  return __nvvm_redux_sync_min(__value, __mask);
 }
 __device__ inline int __reduce_max_sync(unsigned __mask, int __value) {
-  return __nvvm_redux_sync_max(__mask, __value);
+  return __nvvm_redux_sync_max(__value, __mask);
 }
 __device__ inline unsigned __reduce_or_sync(unsigned __mask, unsigned __value) {
-  return __nvvm_redux_sync_or(__mask, __value);
+  return __nvvm_redux_sync_or(__value, __mask);
 }
 __device__ inline unsigned __reduce_and_sync(unsigned __mask,
                                              unsigned __value) {
-  return __nvvm_redux_sync_and(__mask, __value);
+  return __nvvm_redux_sync_and(__value, __mask);
 }
 __device__ inline unsigned __reduce_xor_sync(unsigned __mask,
                                              unsigned __value) {
-  return __nvvm_redux_sync_xor(__mask, __value);
+  return __nvvm_redux_sync_xor(__value, __mask);
 }
 
 __device__ inline void __nv_memcpy_async_shared_global_4(void *__dst,
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 1943e94c3ee7a..b2e05a567b4fe 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -315,7 +315,7 @@ defm MATCH_ALLP_SYNC_64 : MATCH_ALLP_SYNC<Int64Regs, "b64", int_nvvm_match_all_s
 multiclass REDUX_SYNC<string BinOp, string PTXType, Intrinsic Intrin> {
   def : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$src, Int32Regs:$mask),
           "redux.sync." # BinOp # "." # PTXType # " $dst, $src, $mask;",
-          [(set i32:$dst, (Intrin i32:$mask, Int32Regs:$src))]>,
+          [(set i32:$dst, (Intrin i32:$src, Int32Regs:$mask))]>,
         Requires<[hasPTX<70>, hasSM<80>]>;
 }