[CUDA] Fix order of memcpy arguments in __shfl_*(<64-bit type>).

[tra]

Summary

Wrong argument order resulted in broken shfl ops for 64-bit types.

[timshen]

What's the test situation for these headers?

[tra]

Landed in https://github.com/llvm/llvm-project/commit/cc14de88da27a8178976972bdc8211c31f7ca9ae
@hans -- can we cherry-pick it into 10?

[tra]

In D73237#1836978, @timshen wrote:

What's the test situation for these headers?

They are not testable in clang/llvm tree as they need CUDA SDK headers. I should be able to add some tests in the test-suite.

[hans]

In D73237#1837077, @tra wrote:

Landed in https://github.com/llvm/llvm-project/commit/cc14de88da27a8178976972bdc8211c31f7ca9ae
@hans -- can we cherry-pick it into 10?

Yes, go ahead and "git cherry-pick -x" it and push to the branch.

[tra]

In D73237#1839216, @hans wrote:

In D73237#1837077, @tra wrote:

Landed in https://github.com/llvm/llvm-project/commit/cc14de88da27a8178976972bdc8211c31f7ca9ae
@hans -- can we cherry-pick it into 10?

Yes, go ahead and "git cherry-pick -x" it and push to the branch.

Done: https://github.com/llvm/llvm-project/commit/0df13627c6a4006de39e5f01d81a338793b0e82b

clang/lib/Headers/__clang_cuda_intrinsics.h

#define __MAKE_SHUFFLES(__FnName, __IntIntrinsic, __FloatIntrinsic, __Mask,    \
   inline __device__ long long __FnName(long long __val, __Type __offset,       \
                                        int __width = warpSize) {               \
     struct __Bits {                                                            \
       int __a, __b;                                                            \
     };                                                                         \
     _Static_assert(sizeof(__val) == sizeof(__Bits));                           \
     _Static_assert(sizeof(__Bits) == 2 * sizeof(int));                         \
     __Bits __tmp;                                                              \
-    memcpy(&__val, &__tmp, sizeof(__val));                                     \
+    memcpy(&__tmp, &__val, sizeof(__val));                                \
     __tmp.__a = ::__FnName(__tmp.__a, __offset, __width);                      \
     __tmp.__b = ::__FnName(__tmp.__b, __offset, __width);                      \
     long long __ret;                                                           \
     memcpy(&__ret, &__tmp, sizeof(__tmp));                                     \
     return __ret;                                                              \
   }                                                                            \
   inline __device__ long __FnName(long __val, __Type __offset,                 \
                                   int __width = warpSize) {                    \
inline __device__ long long __FnName(unsigned int __mask, long long __val,   \
                                        __Type __offset,                        \
                                        int __width = warpSize) {               \
     struct __Bits {                                                            \
       int __a, __b;                                                            \
     };                                                                         \
     _Static_assert(sizeof(__val) == sizeof(__Bits));                           \
     _Static_assert(sizeof(__Bits) == 2 * sizeof(int));                         \
     __Bits __tmp;                                                              \
-    memcpy(&__val, &__tmp, sizeof(__val));                                     \
+    memcpy(&__tmp, &__val, sizeof(__val));                                     \
     __tmp.__a = ::__FnName(__mask, __tmp.__a, __offset, __width);              \
     __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,          \