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 foo(int r,...)12void foo(int r, ...) { 13 #ifdef IMMEDIATE 14 // expected-error@+4 {{CUDA device code does not support va_arg}} 15 #endif // IMMEDIATE 16 __builtin_va_list list; 17 __builtin_va_start(list, r); 18 (void)__builtin_va_arg(list, int); 19 __builtin_va_end(list); 20 } 21 #ifdef IMMEDIATE 22 #pragma omp declare target to(foo) 23 #endif //IMMEDIATE 24 25 #ifdef IMMEDIATE 26 #pragma omp declare target 27 #endif //IMMEDIATE t1(int r,...)28void t1(int r, ...) { 29 #ifdef DIAGS 30 // expected-error@+4 {{CUDA device code does not support va_arg}} 31 #endif // DIAGS 32 __builtin_va_list list; 33 __builtin_va_start(list, r); 34 (void)__builtin_va_arg(list, int); 35 __builtin_va_end(list); 36 } 37 38 #ifdef IMMEDIATE 39 #pragma omp end declare target 40 #endif //IMMEDIATE 41 main()42int main() { 43 #ifdef DELAYED 44 #pragma omp target 45 #endif // DELAYED 46 { 47 #ifdef DELAYED 48 // expected-note@+2 {{called by 'main'}} 49 #endif // DELAYED 50 t1(0); 51 } 52 return 0; 53 } 54