diff options
Diffstat (limited to 'clang/test/CodeGen/builtins-nvptx.c')
| -rw-r--r-- | clang/test/CodeGen/builtins-nvptx.c | 266 |
1 files changed, 260 insertions, 6 deletions
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c index cd21361140b..b0d646a51fe 100644 --- a/clang/test/CodeGen/builtins-nvptx.c +++ b/clang/test/CodeGen/builtins-nvptx.c @@ -1,8 +1,12 @@ // REQUIRES: nvptx-registered-target -// RUN: %clang_cc1 -triple nvptx-unknown-unknown -fcuda-is-device -S -emit-llvm -o - -x cuda %s | \ -// RUN: FileCheck -check-prefix=CHECK -check-prefix=LP32 %s -// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device -S -emit-llvm -o - -x cuda %s | \ -// RUN: FileCheck -check-prefix=CHECK -check-prefix=LP64 %s +// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_60 \ +// RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \ +// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP32 %s +// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_60 \ +// RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \ +// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s +// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_53 \ +// RUN: -DERROR_CHECK -fcuda-is-device -S -o /dev/null -x cuda -verify %s #define __device__ __attribute__((device)) #define __global__ __attribute__((global)) @@ -191,8 +195,9 @@ __shared__ long long sll; // Check for atomic intrinsics // CHECK-LABEL: nvvm_atom -__device__ void nvvm_atom(float *fp, float f, int *ip, int i, unsigned int *uip, unsigned ui, long *lp, long l, - long long *llp, long long ll) { +__device__ void nvvm_atom(float *fp, float f, double *dfp, double df, int *ip, + int i, unsigned int *uip, unsigned ui, long *lp, + long l, long long *llp, long long ll) { // CHECK: atomicrmw add __nvvm_atom_add_gen_i(ip, i); // CHECK: atomicrmw add @@ -280,6 +285,255 @@ __device__ void nvvm_atom(float *fp, float f, int *ip, int i, unsigned int *uip, // CHECK: call i32 @llvm.nvvm.atomic.load.dec.32.p0i32 __nvvm_atom_dec_gen_ui(uip, ui); + + ////////////////////////////////////////////////////////////////// + // Atomics with scope (only supported on sm_60+). + +#if ERROR_CHECK || __CUDA_ARCH__ >= 600 + + // CHECK: call i32 @llvm.nvvm.atomic.add.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_add_gen_i' needs target feature satom}} + __nvvm_atom_cta_add_gen_i(ip, i); + // LP32: call i32 @llvm.nvvm.atomic.add.gen.i.cta.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.add.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_add_gen_l' needs target feature satom}} + __nvvm_atom_cta_add_gen_l(&dl, l); + // CHECK: call i64 @llvm.nvvm.atomic.add.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_add_gen_ll' needs target feature satom}} + __nvvm_atom_cta_add_gen_ll(&sll, ll); + // CHECK: call i32 @llvm.nvvm.atomic.add.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_add_gen_i' needs target feature satom}} + __nvvm_atom_sys_add_gen_i(ip, i); + // LP32: call i32 @llvm.nvvm.atomic.add.gen.i.sys.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.add.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_add_gen_l' needs target feature satom}} + __nvvm_atom_sys_add_gen_l(&dl, l); + // CHECK: call i64 @llvm.nvvm.atomic.add.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_add_gen_ll' needs target feature satom}} + __nvvm_atom_sys_add_gen_ll(&sll, ll); + + // CHECK: call float @llvm.nvvm.atomic.add.gen.f.cta.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_cta_add_gen_f' needs target feature satom}} + __nvvm_atom_cta_add_gen_f(fp, f); + // CHECK: call double @llvm.nvvm.atomic.add.gen.f.cta.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_cta_add_gen_d' needs target feature satom}} + __nvvm_atom_cta_add_gen_d(dfp, df); + // CHECK: call float @llvm.nvvm.atomic.add.gen.f.sys.f32.p0f32 + // expected-error@+1 {{'__nvvm_atom_sys_add_gen_f' needs target feature satom}} + __nvvm_atom_sys_add_gen_f(fp, f); + // CHECK: call double @llvm.nvvm.atomic.add.gen.f.sys.f64.p0f64 + // expected-error@+1 {{'__nvvm_atom_sys_add_gen_d' needs target feature satom}} + __nvvm_atom_sys_add_gen_d(dfp, df); + + // CHECK: call i32 @llvm.nvvm.atomic.exch.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_i' needs target feature satom}} + __nvvm_atom_cta_xchg_gen_i(ip, i); + // LP32: call i32 @llvm.nvvm.atomic.exch.gen.i.cta.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.exch.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_l' needs target feature satom}} + __nvvm_atom_cta_xchg_gen_l(&dl, l); + // CHECK: call i64 @llvm.nvvm.atomic.exch.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_xchg_gen_ll' needs target feature satom}} + __nvvm_atom_cta_xchg_gen_ll(&sll, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.exch.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_i' needs target feature satom}} + __nvvm_atom_sys_xchg_gen_i(ip, i); + // LP32: call i32 @llvm.nvvm.atomic.exch.gen.i.sys.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.exch.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_l' needs target feature satom}} + __nvvm_atom_sys_xchg_gen_l(&dl, l); + // CHECK: call i64 @llvm.nvvm.atomic.exch.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_xchg_gen_ll' needs target feature satom}} + __nvvm_atom_sys_xchg_gen_ll(&sll, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_i' needs target feature satom}} + __nvvm_atom_cta_max_gen_i(ip, i); + // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ui' needs target feature satom}} + __nvvm_atom_cta_max_gen_ui((unsigned int *)ip, i); + // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_l' needs target feature satom}} + __nvvm_atom_cta_max_gen_l(&dl, l); + // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.cta.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ul' needs target feature satom}} + __nvvm_atom_cta_max_gen_ul((unsigned long *)lp, l); + // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ll' needs target feature satom}} + __nvvm_atom_cta_max_gen_ll(&sll, ll); + // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_max_gen_ull' needs target feature satom}} + __nvvm_atom_cta_max_gen_ull((unsigned long long *)llp, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_i' needs target feature satom}} + __nvvm_atom_sys_max_gen_i(ip, i); + // CHECK: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ui' needs target feature satom}} + __nvvm_atom_sys_max_gen_ui((unsigned int *)ip, i); + // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_l' needs target feature satom}} + __nvvm_atom_sys_max_gen_l(&dl, l); + // LP32: call i32 @llvm.nvvm.atomic.max.gen.i.sys.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ul' needs target feature satom}} + __nvvm_atom_sys_max_gen_ul((unsigned long *)lp, l); + // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ll' needs target feature satom}} + __nvvm_atom_sys_max_gen_ll(&sll, ll); + // CHECK: call i64 @llvm.nvvm.atomic.max.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_max_gen_ull' needs target feature satom}} + __nvvm_atom_sys_max_gen_ull((unsigned long long *)llp, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_i' needs target feature satom}} + __nvvm_atom_cta_min_gen_i(ip, i); + // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ui' needs target feature satom}} + __nvvm_atom_cta_min_gen_ui((unsigned int *)ip, i); + // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_l' needs target feature satom}} + __nvvm_atom_cta_min_gen_l(&dl, l); + // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.cta.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ul' needs target feature satom}} + __nvvm_atom_cta_min_gen_ul((unsigned long *)lp, l); + // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ll' needs target feature satom}} + __nvvm_atom_cta_min_gen_ll(&sll, ll); + // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_min_gen_ull' needs target feature satom}} + __nvvm_atom_cta_min_gen_ull((unsigned long long *)llp, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_i' needs target feature satom}} + __nvvm_atom_sys_min_gen_i(ip, i); + // CHECK: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ui' needs target feature satom}} + __nvvm_atom_sys_min_gen_ui((unsigned int *)ip, i); + // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_l' needs target feature satom}} + __nvvm_atom_sys_min_gen_l(&dl, l); + // LP32: call i32 @llvm.nvvm.atomic.min.gen.i.sys.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ul' needs target feature satom}} + __nvvm_atom_sys_min_gen_ul((unsigned long *)lp, l); + // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ll' needs target feature satom}} + __nvvm_atom_sys_min_gen_ll(&sll, ll); + // CHECK: call i64 @llvm.nvvm.atomic.min.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_min_gen_ull' needs target feature satom}} + __nvvm_atom_sys_min_gen_ull((unsigned long long *)llp, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.inc.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_inc_gen_ui' needs target feature satom}} + __nvvm_atom_cta_inc_gen_ui((unsigned int *)ip, i); + // CHECK: call i32 @llvm.nvvm.atomic.inc.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_inc_gen_ui' needs target feature satom}} + __nvvm_atom_sys_inc_gen_ui((unsigned int *)ip, i); + + // CHECK: call i32 @llvm.nvvm.atomic.dec.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_dec_gen_ui' needs target feature satom}} + __nvvm_atom_cta_dec_gen_ui((unsigned int *)ip, i); + // CHECK: call i32 @llvm.nvvm.atomic.dec.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_dec_gen_ui' needs target feature satom}} + __nvvm_atom_sys_dec_gen_ui((unsigned int *)ip, i); + + // CHECK: call i32 @llvm.nvvm.atomic.and.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_and_gen_i' needs target feature satom}} + __nvvm_atom_cta_and_gen_i(ip, i); + // LP32: call i32 @llvm.nvvm.atomic.and.gen.i.cta.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.and.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_and_gen_l' needs target feature satom}} + __nvvm_atom_cta_and_gen_l(&dl, l); + // CHECK: call i64 @llvm.nvvm.atomic.and.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_and_gen_ll' needs target feature satom}} + __nvvm_atom_cta_and_gen_ll(&sll, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.and.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_and_gen_i' needs target feature satom}} + __nvvm_atom_sys_and_gen_i(ip, i); + // LP32: call i32 @llvm.nvvm.atomic.and.gen.i.sys.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.and.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_and_gen_l' needs target feature satom}} + __nvvm_atom_sys_and_gen_l(&dl, l); + // CHECK: call i64 @llvm.nvvm.atomic.and.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_and_gen_ll' needs target feature satom}} + __nvvm_atom_sys_and_gen_ll(&sll, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.or.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_or_gen_i' needs target feature satom}} + __nvvm_atom_cta_or_gen_i(ip, i); + // LP32: call i32 @llvm.nvvm.atomic.or.gen.i.cta.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.or.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_or_gen_l' needs target feature satom}} + __nvvm_atom_cta_or_gen_l(&dl, l); + // CHECK: call i64 @llvm.nvvm.atomic.or.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_or_gen_ll' needs target feature satom}} + __nvvm_atom_cta_or_gen_ll(&sll, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.or.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_or_gen_i' needs target feature satom}} + __nvvm_atom_sys_or_gen_i(ip, i); + // LP32: call i32 @llvm.nvvm.atomic.or.gen.i.sys.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.or.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_or_gen_l' needs target feature satom}} + __nvvm_atom_sys_or_gen_l(&dl, l); + // CHECK: call i64 @llvm.nvvm.atomic.or.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_or_gen_ll' needs target feature satom}} + __nvvm_atom_sys_or_gen_ll(&sll, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.xor.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_i' needs target feature satom}} + __nvvm_atom_cta_xor_gen_i(ip, i); + // LP32: call i32 @llvm.nvvm.atomic.xor.gen.i.cta.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.xor.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_l' needs target feature satom}} + __nvvm_atom_cta_xor_gen_l(&dl, l); + // CHECK: call i64 @llvm.nvvm.atomic.xor.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_xor_gen_ll' needs target feature satom}} + __nvvm_atom_cta_xor_gen_ll(&sll, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.xor.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_i' needs target feature satom}} + __nvvm_atom_sys_xor_gen_i(ip, i); + // LP32: call i32 @llvm.nvvm.atomic.xor.gen.i.sys.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.xor.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_l' needs target feature satom}} + __nvvm_atom_sys_xor_gen_l(&dl, l); + // CHECK: call i64 @llvm.nvvm.atomic.xor.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_xor_gen_ll' needs target feature satom}} + __nvvm_atom_sys_xor_gen_ll(&sll, ll); + + // CHECK: call i32 @llvm.nvvm.atomic.cas.gen.i.cta.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_i' needs target feature satom}} + __nvvm_atom_cta_cas_gen_i(ip, i, 0); + // LP32: call i32 @llvm.nvvm.atomic.cas.gen.i.cta.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.cas.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_l' needs target feature satom}} + __nvvm_atom_cta_cas_gen_l(&dl, l, 0); + // CHECK: call i64 @llvm.nvvm.atomic.cas.gen.i.cta.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_cta_cas_gen_ll' needs target feature satom}} + __nvvm_atom_cta_cas_gen_ll(&sll, ll, 0); + + // CHECK: call i32 @llvm.nvvm.atomic.cas.gen.i.sys.i32.p0i32 + // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_i' needs target feature satom}} + __nvvm_atom_sys_cas_gen_i(ip, i, 0); + // LP32: call i32 @llvm.nvvm.atomic.cas.gen.i.sys.i32.p0i32 + // LP64: call i64 @llvm.nvvm.atomic.cas.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_l' needs target feature satom}} + __nvvm_atom_sys_cas_gen_l(&dl, l, 0); + // CHECK: call i64 @llvm.nvvm.atomic.cas.gen.i.sys.i64.p0i64 + // expected-error@+1 {{'__nvvm_atom_sys_cas_gen_ll' needs target feature satom}} + __nvvm_atom_sys_cas_gen_ll(&sll, ll, 0); +#endif + // CHECK: ret } |
