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 template <typename tx>
  50 tx ftemplate(int n) {
  51   tx a = 0;
  52   tx b[10];
  53
  54 #pragma omp target reduction(+ \
  55                              : a, b) // expected-note {{defined as threadprivate or thread local}}
  56   {
  57     int e;                        // expected-note {{defined as threadprivate or thread local}}
  58 #pragma omp parallel shared(a, e) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
  59     a += 1;
  60     b[2] += 1;
  61   }
  62
  63   return a;
  64 }
  65
  66 static int fstatic(int n) {
  67   int a = 0;
  68   char aaa = 0;
  69   int b[10];
  70
  71 #pragma omp target firstprivate(a, aaa, b)
  72   {
  73     a += 1;
  74     aaa += 1;
  75     b[2] += 1;
  76   }
  77
  78   return a;
  79 }
  80
  81 struct S1 {
  82   double a;
  83
  84   int r1(int n) {
  85     int b = n + 1;
  86
  87 #pragma omp target firstprivate(b) // expected-note {{defined as threadprivate or thread local}}
  88     {
  89       int c;                      // expected-note {{defined as threadprivate or thread local}}
  90 #pragma omp parallel shared(b, c) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
  91       this->a = (double)b + 1.5;
  92     }
  93
  94     return (int)b;
  95   }
  96 };
  97
  98 int bar(int n, double *ptr) {
  99   int a = 0;
 100   a += foo(n, ptr);
 101   S1 S;
 102   a += S.r1(n);
 103   a += fstatic(n);
 104   a += ftemplate<int>(n); // expected-note {{in instantiation of function template specialization 'ftemplate<int>' requested here}}
 105
 106   return a;
 107 }
 108