clang/test/OpenMP/nvptx_target_cuda_mode_messages.cpp

   1 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
   2 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-cuda-mode -fopenmp-host-ir-file-path %t-ppc-host.bc -o -
   3
   4 template <typename tx, typename ty>
   5 struct TT {
   6   tx X;
   7   ty Y;
   8 };
   9
  10 int foo(int n, double *ptr) {
  11   int a = 0;
  12   short aa = 0;
  13   float b[10];
  14   double c[5][10];
  15   TT<long long, char> d;
  16
  17 #pragma omp target firstprivate(a) map(tofrom: b) // expected-note 2 {{defined as threadprivate or thread local}}
  18   {
  19     int c;                               // expected-note {{defined as threadprivate or thread local}}
  20 #pragma omp parallel shared(a, b, c, aa) // expected-error 3 {{threadprivate or thread local variable cannot be shared}}
  21     b[a] = a;
  22 #pragma omp parallel for
  23     for (int i = 0; i < 10; ++i) // expected-note {{defined as threadprivate or thread local}}
  24 #pragma omp parallel shared(i) // expected-error {{threadprivate or thread local variable cannot be shared}}
  25     ++i;
  26   }
  27
  28 #pragma omp target map(aa, b, c, d)
  29   {
  30     int e;                         // expected-note {{defined as threadprivate or thread local}}
  31 #pragma omp parallel private(b, e) // expected-error {{threadprivate or thread local variable cannot be private}}
  32     {
  33       aa += 1;
  34       b[2] = 1.0;
  35       c[1][2] = 1.0;
  36       d.X = 1;
  37       d.Y = 1;
  38     }
  39   }
  40
  41 #pragma omp target private(ptr)
  42   {
  43     ptr[0]++;
  44   }
  45
  46   return a;
  47 }
  48
  49 static int fstatic(int n) {
  50   int a = 0;
  51   char aaa = 0;
  52   int b[10];
  53
  54 #pragma omp target firstprivate(a, aaa, b)
  55   {
  56     a += 1;
  57     aaa += 1;
  58     b[2] += 1;
  59   }
  60
  61   return a;
  62 }
  63
  64 struct S1 {
  65   double a;
  66
  67   int r1(int n) {
  68     int b = n + 1;
  69
  70 #pragma omp target firstprivate(b) // expected-note {{defined as threadprivate or thread local}}
  71     {
  72       int c;                      // expected-note {{defined as threadprivate or thread local}}
  73 #pragma omp parallel shared(b, c) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
  74       this->a = (double)b + 1.5;
  75     }
  76
  77     return (int)b;
  78   }
  79 };
  80
  81 int bar(int n, double *ptr) {
  82   int a = 0;
  83   a += foo(n, ptr);
  84   S1 S;
  85   a += S.r1(n);
  86   a += fstatic(n);
  87
  88   return a;
  89 }
  90