| 1 | // RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc |
|---|
| 2 | // RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only |
|---|
| 3 | // RUN: %clang_cc1 -verify -DDIAGS -DIMMEDIATE -fopenmp -x c -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only |
|---|
| 4 | // RUN: %clang_cc1 -verify -DDIAGS -DDELAYED -fopenmp -x c -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only |
|---|
| 5 | // REQUIRES: x86-registered-target |
|---|
| 6 | // REQUIRES: nvptx-registered-target |
|---|
| 7 | |
|---|
| 8 | #ifndef DIAGS |
|---|
| 9 | // expected-no-diagnostics |
|---|
| 10 | #endif // DIAGS |
|---|
| 11 | |
|---|
| 12 | #ifdef IMMEDIATE |
|---|
| 13 | #pragma omp declare target |
|---|
| 14 | #endif //IMMEDIATE |
|---|
| 15 | void t1(int r, ...) { |
|---|
| 16 | #ifdef DIAGS |
|---|
| 17 | // expected-error@+4 {{CUDA device code does not support va_arg}} |
|---|
| 18 | #endif // DIAGS |
|---|
| 19 | __builtin_va_list list; |
|---|
| 20 | __builtin_va_start(list, r); |
|---|
| 21 | (void)__builtin_va_arg(list, int); |
|---|
| 22 | __builtin_va_end(list); |
|---|
| 23 | } |
|---|
| 24 | |
|---|
| 25 | #ifdef IMMEDIATE |
|---|
| 26 | #pragma omp end declare target |
|---|
| 27 | #endif //IMMEDIATE |
|---|
| 28 | |
|---|
| 29 | int main() { |
|---|
| 30 | #ifdef DELAYED |
|---|
| 31 | #pragma omp target |
|---|
| 32 | #endif // DELAYED |
|---|
| 33 | { |
|---|
| 34 | #ifdef DELAYED |
|---|
| 35 | // expected-note@+2 {{called by 'main'}} |
|---|
| 36 | #endif // DELAYED |
|---|
| 37 | t1(0); |
|---|
| 38 | } |
|---|
| 39 | return 0; |
|---|
| 40 | } |
|---|
| 41 | |
|---|