106 lines · plain
1// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_70 \2// RUN: -fcuda-is-device -target-feature +ptx60 \3// RUN: -emit-llvm -o - -x cuda %s \4// RUN: | FileCheck -check-prefix=CHECK %s5// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_80 \6// RUN: -fcuda-is-device -target-feature +ptx65 \7// RUN: -emit-llvm -o - -x cuda %s \8// RUN: | FileCheck -check-prefix=CHECK %s9// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_80 \10// RUN: -fcuda-is-device -target-feature +ptx70 \11// RUN: -emit-llvm -o - -x cuda %s \12// RUN: | FileCheck -check-prefix=CHECK %s13// RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_70 \14// RUN: -fcuda-is-device -S -o /dev/null -x cuda -verify %s15 16#define __device__ __attribute__((device))17#define __global__ __attribute__((global))18#define __shared__ __attribute__((shared))19#define __constant__ __attribute__((constant))20 21typedef unsigned long long uint64_t;22 23// We have to keep all builtins that depend on particular target feature in the24// same function, because the codegen will stop after the very first function25// that encounters an error, so -verify will not be able to find errors in26// subsequent functions.27 28// CHECK-LABEL: nvvm_sync29__device__ void nvvm_sync(unsigned mask, int i, float f, int a, int b,30 bool pred, uint64_t i64) {31 32 // CHECK: call void @llvm.nvvm.bar.warp.sync(i3233 // expected-error@+1 {{'__nvvm_bar_warp_sync' needs target feature ptx60}}34 __nvvm_bar_warp_sync(mask);35 // CHECK: call void @llvm.nvvm.barrier.cta.sync.all(i3236 // expected-error@+1 {{'__nvvm_barrier_sync' needs target feature ptx60}}37 __nvvm_barrier_sync(mask);38 // CHECK: call void @llvm.nvvm.barrier.cta.sync.count(i3239 // expected-error@+1 {{'__nvvm_barrier_sync_cnt' needs target feature ptx60}}40 __nvvm_barrier_sync_cnt(mask, i);41 42 //43 // SHFL.SYNC44 //45 // CHECK: call i32 @llvm.nvvm.shfl.sync.down.i32(i32 {{%[0-9]+}}, i3246 // expected-error@+1 {{'__nvvm_shfl_sync_down_i32' needs target feature ptx60}}47 __nvvm_shfl_sync_down_i32(mask, i, a, b);48 // CHECK: call contract float @llvm.nvvm.shfl.sync.down.f32(i32 {{%[0-9]+}}, float49 // expected-error@+1 {{'__nvvm_shfl_sync_down_f32' needs target feature ptx60}}50 __nvvm_shfl_sync_down_f32(mask, f, a, b);51 // CHECK: call i32 @llvm.nvvm.shfl.sync.up.i32(i32 {{%[0-9]+}}, i3252 // expected-error@+1 {{'__nvvm_shfl_sync_up_i32' needs target feature ptx60}}53 __nvvm_shfl_sync_up_i32(mask, i, a, b);54 // CHECK: call contract float @llvm.nvvm.shfl.sync.up.f32(i32 {{%[0-9]+}}, float55 // expected-error@+1 {{'__nvvm_shfl_sync_up_f32' needs target feature ptx60}}56 __nvvm_shfl_sync_up_f32(mask, f, a, b);57 // CHECK: call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 {{%[0-9]+}}, i3258 // expected-error@+1 {{'__nvvm_shfl_sync_bfly_i32' needs target feature ptx60}}59 __nvvm_shfl_sync_bfly_i32(mask, i, a, b);60 // CHECK: call contract float @llvm.nvvm.shfl.sync.bfly.f32(i32 {{%[0-9]+}}, float61 // expected-error@+1 {{'__nvvm_shfl_sync_bfly_f32' needs target feature ptx60}}62 __nvvm_shfl_sync_bfly_f32(mask, f, a, b);63 // CHECK: call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 {{%[0-9]+}}, i3264 // expected-error@+1 {{'__nvvm_shfl_sync_idx_i32' needs target feature ptx60}}65 __nvvm_shfl_sync_idx_i32(mask, i, a, b);66 // CHECK: call contract float @llvm.nvvm.shfl.sync.idx.f32(i32 {{%[0-9]+}}, float67 // expected-error@+1 {{'__nvvm_shfl_sync_idx_f32' needs target feature ptx60}}68 __nvvm_shfl_sync_idx_f32(mask, f, a, b);69 70 //71 // VOTE.SYNC72 //73 74 // CHECK: call i1 @llvm.nvvm.vote.all.sync(i3275 // expected-error@+1 {{'__nvvm_vote_all_sync' needs target feature ptx60}}76 __nvvm_vote_all_sync(mask, pred);77 // CHECK: call i1 @llvm.nvvm.vote.any.sync(i3278 // expected-error@+1 {{'__nvvm_vote_any_sync' needs target feature ptx60}}79 __nvvm_vote_any_sync(mask, pred);80 // CHECK: call i1 @llvm.nvvm.vote.uni.sync(i3281 // expected-error@+1 {{'__nvvm_vote_uni_sync' needs target feature ptx60}}82 __nvvm_vote_uni_sync(mask, pred);83 // CHECK: call i32 @llvm.nvvm.vote.ballot.sync(i3284 // expected-error@+1 {{'__nvvm_vote_ballot_sync' needs target feature ptx60}}85 __nvvm_vote_ballot_sync(mask, pred);86 87 //88 // MATCH.{ALL,ANY}.SYNC89 //90 91 // CHECK: call i32 @llvm.nvvm.match.any.sync.i32(i3292 // expected-error-re@+1 {{'__nvvm_match_any_sync_i32' needs target feature (sm_70{{.*}}),(ptx60{{.*}})}}93 __nvvm_match_any_sync_i32(mask, i);94 // CHECK: call i32 @llvm.nvvm.match.any.sync.i64(i3295 // expected-error-re@+1 {{'__nvvm_match_any_sync_i64' needs target feature (sm_70{{.*}}),(ptx60{{.*}})}}96 __nvvm_match_any_sync_i64(mask, i64);97 // CHECK: call { i32, i1 } @llvm.nvvm.match.all.sync.i32p(i3298 // expected-error-re@+1 {{'__nvvm_match_all_sync_i32p' needs target feature (sm_70{{.*}}),(ptx60{{.*}})}}99 __nvvm_match_all_sync_i32p(mask, i, &i);100 // CHECK: call { i32, i1 } @llvm.nvvm.match.all.sync.i64p(i32101 // expected-error-re@+1 {{'__nvvm_match_all_sync_i64p' needs target feature (sm_70{{.*}}),(ptx60{{.*}})}}102 __nvvm_match_all_sync_i64p(mask, i64, &i);103 104 // CHECK: ret void105}106