1 /* Test of worker-private variables declared in a local scope, broadcasting
2 to vector-partitioned mode. Back-to-back worker loops. */
3
4 /* { dg-additional-options "--param=openacc-kernels=decompose" } */
5
6 /* { dg-additional-options "-fopt-info-omp-all" }
7 { dg-additional-options "-foffload=-fopt-info-omp-all" } */
8
9 /* { dg-additional-options "--param=openacc-privatization=noisy" }
10 { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
11 Prune a few: uninteresting:
12 { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
13
14 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
15 passed to 'incr' may be unset, and in that case, it will be set to [...]",
16 so to maintain compatibility with earlier Tcl releases, we manually
17 initialize counter variables:
18 { dg-line l_dummy[variable c_compute 0 c_loop_i 0 c_loop_j 0 c_loop_k 0] }
19 { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
20 "WARNING: dg-line var l_dummy defined, but not used". */
21
22 #include <assert.h>
23
24 int
25 main (int argc, char* argv[])
26 {
27 int i, arr[32 * 32 * 32];
28
29 for (i = 0; i < 32 * 32 * 32; i++)
30 arr[i] = i;
31
32 #pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
33 /* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
34 decides that a data region is needed for 'j', and subsequently requests it
35 to be made addressable.
36 { dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
37 { dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
38 { dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
39 {
40 int j;
41
42 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
43 #pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
44 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
45 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
46 for (i = 0; i < 32; i++)
47 {
48 #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
49 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
50 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
51 /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
52 for (j = 0; j < 32; j++)
53 {
54 int k;
55 int x = i ^ j * 3;
56
57 #pragma acc loop vector(length:32) /* { dg-line l_loop_k[incr c_loop_k] } */
58 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k$c_loop_k } */
59 for (k = 0; k < 32; k++)
60 arr[i * 1024 + j * 32 + k] += x * k;
61 }
62
63 #pragma acc loop worker(num:32) /* { dg-line l_loop_j[incr c_loop_j] } */
64 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
65 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
66 /* { dg-note {variable 'x' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
67 for (j = 0; j < 32; j++)
68 {
69 int k;
70 int x = i | j * 5;
71
72 #pragma acc loop vector(length:32) /* { dg-line l_loop_k[incr c_loop_k] } */
73 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k$c_loop_k } */
74 for (k = 0; k < 32; k++)
75 arr[i * 1024 + j * 32 + k] += x * k;
76 }
77 }
78 /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
79 }
80
81 for (i = 0; i < 32; i++)
82 for (int j = 0; j < 32; j++)
83 for (int k = 0; k < 32; k++)
84 {
85 int idx = i * 1024 + j * 32 + k;
86 assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
87 }
88
89 return 0;
90 }