jdoerfert created this revision.
jdoerfert added a reviewer: tra.
Herald added subscribers: mattd, bollu, yaxunl.
Herald added a project: All.
jdoerfert requested review of this revision.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.
A copy-paste error caused UB in the definition of the unsigned long long
versions of the shfl intrinsics. Reported and diagnosed by @trws.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D129536
Files:
clang/lib/Headers/__clang_cuda_intrinsics.h
clang/test/CodeGenCUDA/shuffle_long_long.cu
Index: clang/test/CodeGenCUDA/shuffle_long_long.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/shuffle_long_long.cu
@@ -0,0 +1,61 @@
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm %s -o - | FileCheck %s --check-prefix=NO_SYNC
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -target-feature +ptx70 -DSYNC -DCUDA_VERSION=9000 %s -o - | FileCheck %s --check-prefixes=NO_SYNC,SYNC
+
+#include "Inputs/cuda.h"
+
+#undef __CUDA_ARCH__
+#define __CUDA_ARCH__ 300
+
+__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: 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(
+
+// 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(
+
+// NO_SYNC: define{{.*}} @_Z6__shfliii
+// NO_SYNC: call i32 @llvm.nvvm.shfl.idx.i32
+
+// 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;
+ ull = __shfl(ull, 7, 32);
+ ll = __shfl(ll, 7, 32);
+ use(ull, ll);
+#ifdef SYNC
+ ull = __shfl_sync(0x11, ull, 7, 32);
+ ll = __shfl_sync(0x11, ll, 7, 32);
+ use(ull, ll);
+#endif
+}
+
Index: clang/lib/Headers/__clang_cuda_intrinsics.h
===================================================================
--- clang/lib/Headers/__clang_cuda_intrinsics.h
+++ clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -45,7 +45,7 @@
_Static_assert(sizeof(__val) == sizeof(__Bits)); \
_Static_assert(sizeof(__Bits) == 2 * sizeof(int)); \
__Bits __tmp; \
- memcpy(&__tmp, &__val, sizeof(__val)); \
+ memcpy(&__tmp, &__val, sizeof(__val)); \
__tmp.__a = ::__FnName(__tmp.__a, __offset, __width); \
__tmp.__b = ::__FnName(__tmp.__b, __offset, __width); \
long long __ret; \
@@ -71,8 +71,8 @@
} \
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) { \
@@ -139,8 +139,8 @@
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) { \
@@ -234,8 +234,8 @@
return __nvvm_match_any_sync_i32(mask, value);
}
-inline __device__ unsigned int
-__match64_any_sync(unsigned int mask, unsigned long long value) {
+inline __device__ unsigned int __match64_any_sync(unsigned int mask,
+ unsigned long long value) {
return __nvvm_match_any_sync_i64(mask, value);
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
