clang/test/OpenMP/nvptx_va_arg_delayed_diags.c

   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-target-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-target-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-target-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 void 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
  28 void 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
  42 int 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 }