| 1 | // RUN: %clang_cc1 -triple nvptx64-unknown-unknown -target-cpu sm_60 \ |
|---|
| 2 | // RUN: -fcuda-is-device -target-feature +ptx60 \ |
|---|
| 3 | // RUN: -S -emit-llvm -o - -x cuda %s \ |
|---|
| 4 | // RUN: | FileCheck -check-prefix=CHECK %s |
|---|
| 5 | // RUN: %clang_cc1 -triple nvptx-unknown-unknown -target-cpu sm_60 \ |
|---|
| 6 | // RUN: -fcuda-is-device -S -o /dev/null -x cuda -verify %s |
|---|
| 7 | |
|---|
| 8 | #define __device__ __attribute__((device)) |
|---|
| 9 | #define __global__ __attribute__((global)) |
|---|
| 10 | #define __shared__ __attribute__((shared)) |
|---|
| 11 | #define __constant__ __attribute__((constant)) |
|---|
| 12 | |
|---|
| 13 | typedef unsigned long long uint64_t; |
|---|
| 14 | |
|---|
| 15 | // We have to keep all builtins that depend on particular target feature in the |
|---|
| 16 | // same function, because the codegen will stop after the very first function |
|---|
| 17 | // that encounters an error, so -verify will not be able to find errors in |
|---|
| 18 | // subsequent functions. |
|---|
| 19 | |
|---|
| 20 | // CHECK-LABEL: nvvm_sync |
|---|
| 21 | __device__ void nvvm_sync(unsigned mask, int i, float f, int a, int b, |
|---|
| 22 | bool pred, uint64_t i64) { |
|---|
| 23 | |
|---|
| 24 | // CHECK: call void @llvm.nvvm.bar.warp.sync(i32 |
|---|
| 25 | // expected-error@+1 {{'__nvvm_bar_warp_sync' needs target feature ptx60}} |
|---|
| 26 | __nvvm_bar_warp_sync(mask); |
|---|
| 27 | // CHECK: call void @llvm.nvvm.barrier.sync(i32 |
|---|
| 28 | // expected-error@+1 {{'__nvvm_barrier_sync' needs target feature ptx60}} |
|---|
| 29 | __nvvm_barrier_sync(mask); |
|---|
| 30 | // CHECK: call void @llvm.nvvm.barrier.sync.cnt(i32 |
|---|
| 31 | // expected-error@+1 {{'__nvvm_barrier_sync_cnt' needs target feature ptx60}} |
|---|
| 32 | __nvvm_barrier_sync_cnt(mask, i); |
|---|
| 33 | |
|---|
| 34 | // |
|---|
| 35 | // SHFL.SYNC |
|---|
| 36 | // |
|---|
| 37 | // CHECK: call i32 @llvm.nvvm.shfl.sync.down.i32(i32 {{%[0-9]+}}, i32 |
|---|
| 38 | // expected-error@+1 {{'__nvvm_shfl_sync_down_i32' needs target feature ptx60}} |
|---|
| 39 | __nvvm_shfl_sync_down_i32(mask, i, a, b); |
|---|
| 40 | // CHECK: call float @llvm.nvvm.shfl.sync.down.f32(i32 {{%[0-9]+}}, float |
|---|
| 41 | // expected-error@+1 {{'__nvvm_shfl_sync_down_f32' needs target feature ptx60}} |
|---|
| 42 | __nvvm_shfl_sync_down_f32(mask, f, a, b); |
|---|
| 43 | // CHECK: call i32 @llvm.nvvm.shfl.sync.up.i32(i32 {{%[0-9]+}}, i32 |
|---|
| 44 | // expected-error@+1 {{'__nvvm_shfl_sync_up_i32' needs target feature ptx60}} |
|---|
| 45 | __nvvm_shfl_sync_up_i32(mask, i, a, b); |
|---|
| 46 | // CHECK: call float @llvm.nvvm.shfl.sync.up.f32(i32 {{%[0-9]+}}, float |
|---|
| 47 | // expected-error@+1 {{'__nvvm_shfl_sync_up_f32' needs target feature ptx60}} |
|---|
| 48 | __nvvm_shfl_sync_up_f32(mask, f, a, b); |
|---|
| 49 | // CHECK: call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 {{%[0-9]+}}, i32 |
|---|
| 50 | // expected-error@+1 {{'__nvvm_shfl_sync_bfly_i32' needs target feature ptx60}} |
|---|
| 51 | __nvvm_shfl_sync_bfly_i32(mask, i, a, b); |
|---|
| 52 | // CHECK: call float @llvm.nvvm.shfl.sync.bfly.f32(i32 {{%[0-9]+}}, float |
|---|
| 53 | // expected-error@+1 {{'__nvvm_shfl_sync_bfly_f32' needs target feature ptx60}} |
|---|
| 54 | __nvvm_shfl_sync_bfly_f32(mask, f, a, b); |
|---|
| 55 | // CHECK: call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 {{%[0-9]+}}, i32 |
|---|
| 56 | // expected-error@+1 {{'__nvvm_shfl_sync_idx_i32' needs target feature ptx60}} |
|---|
| 57 | __nvvm_shfl_sync_idx_i32(mask, i, a, b); |
|---|
| 58 | // CHECK: call float @llvm.nvvm.shfl.sync.idx.f32(i32 {{%[0-9]+}}, float |
|---|
| 59 | // expected-error@+1 {{'__nvvm_shfl_sync_idx_f32' needs target feature ptx60}} |
|---|
| 60 | __nvvm_shfl_sync_idx_f32(mask, f, a, b); |
|---|
| 61 | |
|---|
| 62 | // |
|---|
| 63 | // VOTE.SYNC |
|---|
| 64 | // |
|---|
| 65 | |
|---|
| 66 | // CHECK: call i1 @llvm.nvvm.vote.all.sync(i32 |
|---|
| 67 | // expected-error@+1 {{'__nvvm_vote_all_sync' needs target feature ptx60}} |
|---|
| 68 | __nvvm_vote_all_sync(mask, pred); |
|---|
| 69 | // CHECK: call i1 @llvm.nvvm.vote.any.sync(i32 |
|---|
| 70 | // expected-error@+1 {{'__nvvm_vote_any_sync' needs target feature ptx60}} |
|---|
| 71 | __nvvm_vote_any_sync(mask, pred); |
|---|
| 72 | // CHECK: call i1 @llvm.nvvm.vote.uni.sync(i32 |
|---|
| 73 | // expected-error@+1 {{'__nvvm_vote_uni_sync' needs target feature ptx60}} |
|---|
| 74 | __nvvm_vote_uni_sync(mask, pred); |
|---|
| 75 | // CHECK: call i32 @llvm.nvvm.vote.ballot.sync(i32 |
|---|
| 76 | // expected-error@+1 {{'__nvvm_vote_ballot_sync' needs target feature ptx60}} |
|---|
| 77 | __nvvm_vote_ballot_sync(mask, pred); |
|---|
| 78 | |
|---|
| 79 | // |
|---|
| 80 | // MATCH.{ALL,ANY}.SYNC |
|---|
| 81 | // |
|---|
| 82 | |
|---|
| 83 | // CHECK: call i32 @llvm.nvvm.match.any.sync.i32(i32 |
|---|
| 84 | // expected-error@+1 {{'__nvvm_match_any_sync_i32' needs target feature ptx60}} |
|---|
| 85 | __nvvm_match_any_sync_i32(mask, i); |
|---|
| 86 | // CHECK: call i64 @llvm.nvvm.match.any.sync.i64(i32 |
|---|
| 87 | // expected-error@+1 {{'__nvvm_match_any_sync_i64' needs target feature ptx60}} |
|---|
| 88 | __nvvm_match_any_sync_i64(mask, i64); |
|---|
| 89 | // CHECK: call { i32, i1 } @llvm.nvvm.match.all.sync.i32p(i32 |
|---|
| 90 | // expected-error@+1 {{'__nvvm_match_all_sync_i32p' needs target feature ptx60}} |
|---|
| 91 | __nvvm_match_all_sync_i32p(mask, i, &i); |
|---|
| 92 | // CHECK: call { i64, i1 } @llvm.nvvm.match.all.sync.i64p(i32 |
|---|
| 93 | // expected-error@+1 {{'__nvvm_match_all_sync_i64p' needs target feature ptx60}} |
|---|
| 94 | __nvvm_match_all_sync_i64p(mask, i64, &i); |
|---|
| 95 | |
|---|
| 96 | // CHECK: ret void |
|---|
| 97 | } |
|---|
| 98 | |
|---|