62 lines · plain
1// REQUIRES: nvptx-registered-target2 3// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -target-cpu sm_30 %s -o - | FileCheck %s --check-prefix=NO_SYNC4// 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=SYNC5 6#include "Inputs/cuda.h"7 8__device__ void *memcpy(void *dest, const void *src, size_t n);9 10#define warpSize 3211#include <__clang_cuda_intrinsics.h>12 13__device__ void use(unsigned long long, long long);14 15// Test function, 4 shfl calls.16// NO_SYNC: define{{.*}} @_Z14test_long_longv17// NO_SYNC: call noundef i64 @_Z6__shflyii(18// NO_SYNC: call noundef i64 @_Z6__shflxii(19 20// SYNC: define{{.*}} @_Z14test_long_longv21// SYNC: call noundef i64 @_Z11__shfl_syncjyii(22// SYNC: call noundef i64 @_Z11__shfl_syncjxii(23 24// unsigned long long -> long long25// NO_SYNC: define{{.*}} @_Z6__shflyii26// NO_SYNC: call noundef i64 @_Z6__shflxii(27 28// long long -> int + int29// NO_SYNC: define{{.*}} @_Z6__shflxii30// NO_SYNC: call noundef i32 @_Z6__shfliii(31// NO_SYNC: call noundef i32 @_Z6__shfliii(32 33// NO_SYNC: define{{.*}} @_Z6__shfliii34// NO_SYNC: call i32 @llvm.nvvm.shfl.idx.i3235 36// unsigned long long -> long long37// SYNC: _Z11__shfl_syncjyii38// SYNC: call noundef i64 @_Z11__shfl_syncjxii(39 40// long long -> int + int41// SYNC: define{{.*}} @_Z11__shfl_syncjxii42// SYNC: call noundef i32 @_Z11__shfl_syncjiii(43// SYNC: call noundef i32 @_Z11__shfl_syncjiii(44 45// SYNC: define{{.*}} @_Z11__shfl_syncjiii46// SYNC: call i32 @llvm.nvvm.shfl.sync.idx.i3247 48__device__ void test_long_long() {49 unsigned long long ull = 13;50 long long ll = 17;51#ifndef SYNC52 ull = __shfl(ull, 7, 32);53 ll = __shfl(ll, 7, 32);54 use(ull, ll);55#else56 ull = __shfl_sync(0x11, ull, 7, 32);57 ll = __shfl_sync(0x11, ll, 7, 32);58 use(ull, ll);59#endif60}61 62