File: nvptx_target_cuda_mode_messages.cpp

package info (click to toggle)
llvm-toolchain-19 1%3A19.1.7-3
  • links: PTS, VCS
  • area: main
  • in suites: forky, sid, trixie
  • size: 1,998,520 kB
  • sloc: cpp: 6,951,680; ansic: 1,486,157; asm: 913,598; python: 232,024; f90: 80,126; objc: 75,281; lisp: 37,276; pascal: 16,990; sh: 10,009; ml: 5,058; perl: 4,724; awk: 3,523; makefile: 3,167; javascript: 2,504; xml: 892; fortran: 664; cs: 573
file content (90 lines) | stat: -rw-r--r-- 2,314 bytes parent folder | download | duplicates (7)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// 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 -

template <typename tx, typename ty>
struct TT {
  tx X;
  ty Y;
};

int foo(int n, double *ptr) {
  int a = 0;
  short aa = 0;
  float b[10];
  double c[5][10];
  TT<long long, char> d;

#pragma omp target firstprivate(a) map(tofrom: b) // expected-note 2 {{defined as threadprivate or thread local}}
  {
    int c;                               // expected-note {{defined as threadprivate or thread local}}
#pragma omp parallel shared(a, b, c, aa) // expected-error 3 {{threadprivate or thread local variable cannot be shared}}
    b[a] = a;
#pragma omp parallel for
    for (int i = 0; i < 10; ++i) // expected-note {{defined as threadprivate or thread local}}
#pragma omp parallel shared(i) // expected-error {{threadprivate or thread local variable cannot be shared}}
    ++i;
  }

#pragma omp target map(aa, b, c, d)
  {
    int e;                         // expected-note {{defined as threadprivate or thread local}}
#pragma omp parallel private(b, e) // expected-error {{threadprivate or thread local variable cannot be private}}
    {
      aa += 1;
      b[2] = 1.0;
      c[1][2] = 1.0;
      d.X = 1;
      d.Y = 1;
    }
  }

#pragma omp target private(ptr)
  {
    ptr[0]++;
  }

  return a;
}

static int fstatic(int n) {
  int a = 0;
  char aaa = 0;
  int b[10];

#pragma omp target firstprivate(a, aaa, b)
  {
    a += 1;
    aaa += 1;
    b[2] += 1;
  }

  return a;
}

struct S1 {
  double a;

  int r1(int n) {
    int b = n + 1;

#pragma omp target firstprivate(b) // expected-note {{defined as threadprivate or thread local}}
    {
      int c;                      // expected-note {{defined as threadprivate or thread local}}
#pragma omp parallel shared(b, c) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
      this->a = (double)b + 1.5;
    }

    return (int)b;
  }
};

int bar(int n, double *ptr) {
  int a = 0;
  a += foo(n, ptr);
  S1 S;
  a += S.r1(n);
  a += fstatic(n);

  return a;
}