[CUDA][FIX] Make shfl[_sync] for unsigned long long non-recursive

[jdoerfert]

Summary

A copy-paste error caused UB in the definition of the unsigned long long
versions of the shfl intrinsics. Reported and diagnosed by @trws.

[tra]

Oops. Thank you for fixing this.

clang/test/CodeGenCUDA/shuffle_long_long.cu

53

This crashes LLVM when we taget sm_70 where these instructions no longer exist. We should probably disable those sync wrappers when we compile for GPUs where they are not available, so we'd get a proper compiler error instead of a crash.

Also, we should probably make non-sync instruction use conditional on SYNC. https://godbolt.org/z/7n4vsb41v

[jdoerfert]

Address comments

[jdoerfert]

@tra, unsure about the crash. For me this passes fine (no gpu), is anything missing?

[tra]

In D129536#3663957, @jdoerfert wrote:

@tra, unsure about the crash. For me this passes fine (no gpu), is anything missing?

The tests in the patch are running with -emit-llvm, so they are not actually lowering to NVPTX and that's where the failure happens. https://godbolt.org/z/cchaWxrhn

clang/lib/Headers/__clang_cuda_intrinsics.h

238–239

Nit: this change is irrelevant to the patch and can be removed.

clang/test/CodeGenCUDA/shuffle_long_long.cu

10

This macro should not be set. If you do need something to be compiled for sm_30, you should've specified via -target-cpu sm_30.

15

Nit: this should be <...> as we want the include to be found in compiler's include paths.

[jdoerfert]

In D129536#3666257, @tra wrote:

In D129536#3663957, @jdoerfert wrote:

@tra, unsure about the crash. For me this passes fine (no gpu), is anything missing?

The tests in the patch are running with -emit-llvm, so they are not actually lowering to NVPTX and that's where the failure happens. https://godbolt.org/z/cchaWxrhn

The assertion is arguably not great but doesn't really matter, does it? How would I detect if they are supported?

clang/lib/Headers/__clang_cuda_intrinsics.h

238–239

me running clang format on the file. I'll push it nfc before.

[tra]

In D129536#3666860, @jdoerfert wrote:

The assertion is arguably not great but doesn't really matter, does it? How would I detect if they are supported?

The latest revision of the patch is fine in this regard. My comment pointing to compiler crash reproducer was only intended to address the "For me this passes fine" part.

The only remaining thing is the manual __CUDA_ARCH__ redifinition which looks suspect to me. Is there any reason not to use -target-cpu sm_30 instead?

[jdoerfert]

Address comments

[jdoerfert]

In D129536#3666884, @tra wrote:

In D129536#3666860, @jdoerfert wrote:

The assertion is arguably not great but doesn't really matter, does it? How would I detect if they are supported?

The latest revision of the patch is fine in this regard. My comment pointing to compiler crash reproducer was only intended to address the "For me this passes fine" part.

The only remaining thing is the manual __CUDA_ARCH__ redifinition which looks suspect to me. Is there any reason not to use -target-cpu sm_30 instead?

No reason. Changed.

[jdoerfert]

Use <...>

clang/lib/Headers/__clang_cuda_intrinsics.h

#define __MAKE_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, __Mask,    \
   }                                                                            \
   inline __device__ unsigned long __FnName(                                    \
       unsigned long __val, __Type __offset, int __width = warpSize) {          \
     return static_cast<unsigned long>(                                         \
         ::__FnName(static_cast<long>(__val), __offset, __width));              \
   }                                                                            \
   inline __device__ unsigned long long __FnName(                               \
       unsigned long long __val, __Type __offset, int __width = warpSize) {     \
-    return static_cast<unsigned long long>(::__FnName(                         \
-        static_cast<unsigned long long>(__val), __offset, __width));           \
+    return static_cast<unsigned long long>(                                    \
+        ::__FnName(static_cast<long long>(__val), __offset, __width));         \
   }                                                                            \
   inline __device__ double __FnName(double __val, __Type __offset,             \
                                     int __width = warpSize) {                  \
     long long __tmp;                                                           \
     _Static_assert(sizeof(__tmp) == sizeof(__val));                            \
     memcpy(&__tmp, &__val, sizeof(__val));                                     \
     __tmp = ::__FnName(__tmp, __offset, __width);                              \
     double __ret;                                                              \
inline __device__ long long __FnName(unsigned int __mask, long long __val,   \
     __tmp.__b = ::__FnName(__mask, __tmp.__b, __offset, __width);              \
     long long __ret;                                                           \
     memcpy(&__ret, &__tmp, sizeof(__tmp));                                     \
     return __ret;                                                              \
   }                                                                            \
   inline __device__ unsigned long long __FnName(                               \
       unsigned int __mask, unsigned long long __val, __Type __offset,          \
       int __width = warpSize) {                                                \
-    return static_cast<unsigned long long>(::__FnName(                         \
-        __mask, static_cast<unsigned long long>(__val), __offset, __width));   \
+    return static_cast<unsigned long long>(                                    \
+        ::__FnName(__mask, static_cast<long long>(__val), __offset, __width)); \
   }                                                                            \
   inline __device__ long __FnName(unsigned int __mask, long __val,             \
                                   __Type __offset, int __width = warpSize) {   \
     _Static_assert(sizeof(long) == sizeof(long long) ||                        \
                    sizeof(long) == sizeof(int));                               \
     if (sizeof(long) == sizeof(long long)) {                                   \
       return static_cast<long>(::__FnName(                                     \
           __mask, static_cast<long long>(__val), __offset, __width));          \

 // Define __match* builtins CUDA-9 headers expect to see.
 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 700
 inline __device__ unsigned int __match32_any_sync(unsigned int mask,
                                                   unsigned int value) {
   return __nvvm_match_any_sync_i32(mask, value);
 }
 
 inline __device__ unsigned int
 __match64_any_sync(unsigned int mask, unsigned long long value) {
   return __nvvm_match_any_sync_i64(mask, value);
 }
 
 inline __device__ unsigned int
 __match32_all_sync(unsigned int mask, unsigned int value, int *pred) {
   return __nvvm_match_all_sync_i32p(mask, value, pred);
 }
 
 inline __device__ unsigned int

clang/test/CodeGenCUDA/shuffle_long_long.cu

+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -target-cpu sm_30 %s -o - | FileCheck %s --check-prefix=NO_SYNC
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -target-cpu sm_30 -target-feature +ptx70 -DSYNC -DCUDA_VERSION=9000 %s -o - | FileCheck %s --check-prefix=SYNC
+
+#include "Inputs/cuda.h"
+
+__device__ void *memcpy(void *dest, const void *src, size_t n);
+
+#define warpSize 32
+#include <__clang_cuda_intrinsics.h>
+
+__device__ void use(unsigned long long, long long);
+
+// Test function, 4 shfl calls.
+// NO_SYNC: define{{.*}} @_Z14test_long_longv
+// NO_SYNC:     call noundef i64 @_Z6__shflyii(
+// NO_SYNC:     call noundef i64 @_Z6__shflxii(
+
+// SYNC: define{{.*}} @_Z14test_long_longv
+// SYNC:        call noundef i64 @_Z11__shfl_syncjyii(
+// SYNC:        call noundef i64 @_Z11__shfl_syncjxii(
+
+// unsigned long long -> long long
+// NO_SYNC: define{{.*}} @_Z6__shflyii
+// NO_SYNC:     call noundef i64 @_Z6__shflxii(
+
+// long long -> int + int
+// NO_SYNC: define{{.*}} @_Z6__shflxii
+// NO_SYNC:     call noundef i32 @_Z6__shfliii(
+// NO_SYNC:     call noundef i32 @_Z6__shfliii(
+
+// NO_SYNC: define{{.*}} @_Z6__shfliii
+// NO_SYNC:   call i32 @llvm.nvvm.shfl.idx.i32
+
+// unsigned long long -> long long
+// SYNC: _Z11__shfl_syncjyii
+// SYNC:     call noundef i64 @_Z11__shfl_syncjxii(
+
+// long long -> int + int
+// SYNC: define{{.*}} @_Z11__shfl_syncjxii
+// SYNC:     call noundef i32 @_Z11__shfl_syncjiii(
+// SYNC:     call noundef i32 @_Z11__shfl_syncjiii(
+
+// SYNC: define{{.*}} @_Z11__shfl_syncjiii
+// SYNC:      call i32 @llvm.nvvm.shfl.sync.idx.i32
+
+__device__ void test_long_long() {
+  unsigned long long ull = 13;
+  long long ll = 17;
+#ifndef SYNC
+  ull = __shfl(ull, 7, 32);
+  ll = __shfl(ll, 7, 32);
+  use(ull, ll);
+#else
+  ull = __shfl_sync(0x11, ull, 7, 32);
+  ll = __shfl_sync(0x11, ll, 7, 32);
+  use(ull, ll);
+#endif
+}
+