Skip to content

Commit 184b0ae

Browse files
kevinsalakevinsala
committed
[OpenMP][Offload] Add tests for dyn_groupprivate
1 parent 1a54d76 commit 184b0ae

File tree

3 files changed

+314
-0
lines changed

3 files changed

+314
-0
lines changed

‎clang/test/OpenMP/target_dyn_groupprivate_messages.cpp

Lines changed: 87 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,87 @@
1+
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 %s -Wuninitialized
2+
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 %s -Wuninitialized
3+
// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
4+
// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
5+
6+
void foo() {
7+
}
8+
9+
bool foobool(int argc) {
10+
return argc;
11+
}
12+
13+
struct S1; // expected-note {{declared here}}
14+
15+
template <class T, class S> // expected-note {{declared here}}
16+
int tmain(T argc, S **argv) {
17+
T z;
18+
#pragma omp target dyn_groupprivate // expected-error {{expected '(' after 'dyn_groupprivate'}}
19+
foo();
20+
#pragma omp target dyn_groupprivate ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
21+
foo();
22+
#pragma omp target dyn_groupprivate () // expected-error {{expected expression}}
23+
foo();
24+
#pragma omp target dyn_groupprivate (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
25+
foo();
26+
#pragma omp target dyn_groupprivate (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
27+
foo();
28+
#pragma omp target dyn_groupprivate (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
29+
foo();
30+
#pragma omp target dyn_groupprivate (foobool(argc)), dyn_groupprivate (true) // expected-error {{directive '#pragma omp target' cannot contain more than one 'dyn_groupprivate' clause}}
31+
foo();
32+
#pragma omp target dyn_groupprivate (S) // expected-error {{'S' does not refer to a value}}
33+
foo();
34+
#pragma omp target dyn_groupprivate (argv[1]=2) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
35+
foo();
36+
#pragma omp target dyn_groupprivate (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
37+
foo();
38+
#pragma omp target dyn_groupprivate(argc+z)
39+
foo();
40+
return 0;
41+
}
42+
43+
int main(int argc, char **argv) {
44+
constexpr int n = -1;
45+
int z;
46+
#pragma omp target dyn_groupprivate // expected-error {{expected '(' after 'dyn_groupprivate'}}
47+
foo();
48+
#pragma omp target dyn_groupprivate ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
49+
foo();
50+
#pragma omp target dyn_groupprivate () // expected-error {{expected expression}}
51+
foo();
52+
#pragma omp target dyn_groupprivate (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
53+
foo();
54+
#pragma omp target dyn_groupprivate (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
55+
foo();
56+
#pragma omp target dyn_groupprivate (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
57+
foo();
58+
#pragma omp target dyn_groupprivate (foobool(argc)), dyn_groupprivate (true) // expected-error {{directive '#pragma omp target' cannot contain more than one 'dyn_groupprivate' clause}}
59+
foo();
60+
#pragma omp target dyn_groupprivate (S1) // expected-error {{'S1' does not refer to a value}}
61+
foo();
62+
#pragma omp target dyn_groupprivate (argv[1]=2) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
63+
foo();
64+
#pragma omp target dyn_groupprivate (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
65+
foo();
66+
#pragma omp target dyn_groupprivate (1 0) // expected-error {{expected ')'}} expected-note {{to match this '('}}
67+
foo();
68+
#pragma omp target dyn_groupprivate(dyn_groupprivate(tmain(argc, argv) // expected-error2 {{expected ')'}} expected-note2 {{to match this '('}} expected-note {{in instantiation of function template specialization 'tmain<int, char>' requested here}}
69+
foo();
70+
#pragma omp target dyn_groupprivate(-1) // expected-error {{argument to 'dyn_groupprivate' clause must be a non-negative integer value}}
71+
foo();
72+
#pragma omp target dyn_groupprivate(cgrou) // expected-error {{use of undeclared identifier 'cgrou'}}
73+
foo();
74+
#pragma omp target dyn_groupprivate(cgrou: argc) // expected-error {{use of undeclared identifier 'cgrou'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
75+
foo();
76+
#pragma omp target dyn_groupprivate(cgroup,cgroup: argc) // expected-error {{modifier 'cgroup' cannot be used along with modifier 'cgroup' in dyn_groupprivate}}
77+
foo();
78+
#pragma omp target dyn_groupprivate(fallback,strict: argc) // expected-error {{modifier 'strict' cannot be used along with modifier 'fallback' in dyn_groupprivate}}
79+
foo();
80+
#pragma omp target dyn_groupprivate(strict,fallback: argc) // expected-error {{modifier 'fallback' cannot be used along with modifier 'strict' in dyn_groupprivate}}
81+
foo();
82+
#pragma omp target dyn_groupprivate(: argc) // expected-error {{expected ')'}} expected-error {{expected expression}} expected-note {{to match this '('}}
83+
foo();
84+
85+
return tmain(argc, argv);
86+
}
87+

‎clang/test/OpenMP/target_teams_dyn_groupprivate_messages.cpp

Lines changed: 87 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,87 @@
1+
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 %s -Wuninitialized
2+
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 %s -Wuninitialized
3+
// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
4+
// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
5+
6+
void foo() {
7+
}
8+
9+
bool foobool(int argc) {
10+
return argc;
11+
}
12+
13+
struct S1; // expected-note {{declared here}}
14+
15+
template <class T, class S> // expected-note {{declared here}}
16+
int tmain(T argc, S **argv) {
17+
T z;
18+
#pragma omp target teams dyn_groupprivate // expected-error {{expected '(' after 'dyn_groupprivate'}}
19+
foo();
20+
#pragma omp target teams dyn_groupprivate ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
21+
foo();
22+
#pragma omp target teams dyn_groupprivate () // expected-error {{expected expression}}
23+
foo();
24+
#pragma omp target teams dyn_groupprivate (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
25+
foo();
26+
#pragma omp target teams dyn_groupprivate (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
27+
foo();
28+
#pragma omp target teams dyn_groupprivate (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
29+
foo();
30+
#pragma omp target teams dyn_groupprivate (foobool(argc)), dyn_groupprivate (true) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'dyn_groupprivate' clause}}
31+
foo();
32+
#pragma omp target teams dyn_groupprivate (S) // expected-error {{'S' does not refer to a value}}
33+
foo();
34+
#pragma omp target teams dyn_groupprivate (argv[1]=2) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
35+
foo();
36+
#pragma omp target teams dyn_groupprivate (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
37+
foo();
38+
#pragma omp target teams dyn_groupprivate(argc+z)
39+
foo();
40+
return 0;
41+
}
42+
43+
int main(int argc, char **argv) {
44+
constexpr int n = -1;
45+
int z;
46+
#pragma omp target teams dyn_groupprivate // expected-error {{expected '(' after 'dyn_groupprivate'}}
47+
foo();
48+
#pragma omp target teams dyn_groupprivate ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
49+
foo();
50+
#pragma omp target teams dyn_groupprivate () // expected-error {{expected expression}}
51+
foo();
52+
#pragma omp target teams dyn_groupprivate (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
53+
foo();
54+
#pragma omp target teams dyn_groupprivate (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
55+
foo();
56+
#pragma omp target teams dyn_groupprivate (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
57+
foo();
58+
#pragma omp target teams dyn_groupprivate (foobool(argc)), dyn_groupprivate (true) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'dyn_groupprivate' clause}}
59+
foo();
60+
#pragma omp target teams dyn_groupprivate (S1) // expected-error {{'S1' does not refer to a value}}
61+
foo();
62+
#pragma omp target teams dyn_groupprivate (argv[1]=2) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
63+
foo();
64+
#pragma omp target teams dyn_groupprivate (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
65+
foo();
66+
#pragma omp target teams dyn_groupprivate (1 0) // expected-error {{expected ')'}} expected-note {{to match this '('}}
67+
foo();
68+
#pragma omp target teams dyn_groupprivate(dyn_groupprivate(tmain(argc, argv) // expected-error2 {{expected ')'}} expected-note2 {{to match this '('}} expected-note {{in instantiation of function template specialization 'tmain<int, char>' requested here}}
69+
foo();
70+
#pragma omp target teams dyn_groupprivate(-1) // expected-error {{argument to 'dyn_groupprivate' clause must be a non-negative integer value}}
71+
foo();
72+
#pragma omp target teams dyn_groupprivate(cgrou) // expected-error {{use of undeclared identifier 'cgrou'}}
73+
foo();
74+
#pragma omp target teams dyn_groupprivate(cgrou: argc) // expected-error {{use of undeclared identifier 'cgrou'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
75+
foo();
76+
#pragma omp target teams dyn_groupprivate(cgroup,cgroup: argc) // expected-error {{modifier 'cgroup' cannot be used along with modifier 'cgroup' in dyn_groupprivate}}
77+
foo();
78+
#pragma omp target teams dyn_groupprivate(fallback,strict: argc) // expected-error {{modifier 'strict' cannot be used along with modifier 'fallback' in dyn_groupprivate}}
79+
foo();
80+
#pragma omp target teams dyn_groupprivate(strict,fallback: argc) // expected-error {{modifier 'fallback' cannot be used along with modifier 'strict' in dyn_groupprivate}}
81+
foo();
82+
#pragma omp target teams dyn_groupprivate(: argc) // expected-error {{expected ')'}} expected-error {{expected expression}} expected-note {{to match this '('}}
83+
foo();
84+
85+
return tmain(argc, argv);
86+
}
87+

‎offload/test/offloading/dyn_groupprivate_strict.cpp

Lines changed: 140 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,140 @@
1+
// RUN: %libomptarget-compilexx-run-and-check-generic
2+
// REQUIRES: gpu
3+
4+
#include <omp.h>
5+
#include <stdio.h>
6+
7+
#define N 512
8+
9+
int main() {
10+
int Result[N], NumThreads;
11+
12+
#pragma omp target teams num_teams(1) thread_limit(N) \
13+
dyn_groupprivate(strict: N * sizeof(Result[0])) \
14+
map(from : Result, NumThreads)
15+
{
16+
int Buffer[N];
17+
#pragma omp parallel
18+
{
19+
int *DynBuffer = (int *)omp_get_dyn_groupprivate_ptr();
20+
int TId = omp_get_thread_num();
21+
if (TId == 0)
22+
NumThreads = omp_get_num_threads();
23+
Buffer[TId] = 7;
24+
DynBuffer[TId] = 3;
25+
#pragma omp barrier
26+
int WrappedTId = (TId + 37) % NumThreads;
27+
Result[TId] = Buffer[WrappedTId] + DynBuffer[WrappedTId];
28+
}
29+
}
30+
31+
if (NumThreads < N / 2 || NumThreads > N) {
32+
printf("Expected number of threads to be in [%i:%i], but got: %i", N / 2, N,
33+
NumThreads);
34+
return -1;
35+
}
36+
37+
int Failed = 0;
38+
for (int i = 0; i < NumThreads; ++i) {
39+
if (Result[i] != 7 + 3) {
40+
printf("Result[%i] is %i, expected %i\n", i, Result[i], 7 + 3);
41+
++Failed;
42+
}
43+
}
44+
45+
// Verify that the routines in the host returns NULL and zero.
46+
if (omp_get_dyn_groupprivate_ptr())
47+
++Failed;
48+
if (omp_get_dyn_groupprivate_size())
49+
++Failed;
50+
51+
size_t MaxSize = omp_get_groupprivate_limit(0, omp_access_cgroup);
52+
size_t ExceededSize = MaxSize + 10;
53+
54+
// Verify that the fallback modifier works.
55+
#pragma omp target dyn_groupprivate(fallback: ExceededSize) map(tofrom: Failed)
56+
{
57+
int IsFallback;
58+
if (!omp_get_dyn_groupprivate_ptr(0, &IsFallback))
59+
++Failed;
60+
if (!omp_get_dyn_groupprivate_size())
61+
++Failed;
62+
if (omp_get_dyn_groupprivate_size() != ExceededSize)
63+
++Failed;
64+
if (!IsFallback)
65+
++Failed;
66+
}
67+
68+
// Verify that the default modifier is fallback.
69+
#pragma omp target dyn_groupprivate(ExceededSize)
70+
{
71+
}
72+
73+
// Verify that the strict modifier works.
74+
#pragma omp target dyn_groupprivate(strict: N) map(tofrom: Failed)
75+
{
76+
int IsFallback;
77+
if (!omp_get_dyn_groupprivate_ptr(0, &IsFallback))
78+
++Failed;
79+
if (!omp_get_dyn_groupprivate_size())
80+
++Failed;
81+
if (omp_get_dyn_groupprivate_size() != N)
82+
++Failed;
83+
if (IsFallback)
84+
++Failed;
85+
}
86+
87+
// Verify that the fallback does not trigger when not needed.
88+
#pragma omp target dyn_groupprivate(fallback: N) map(tofrom: Failed)
89+
{
90+
int IsFallback;
91+
if (!omp_get_dyn_groupprivate_ptr(0, &IsFallback))
92+
++Failed;
93+
if (!omp_get_dyn_groupprivate_size())
94+
++Failed;
95+
if (omp_get_dyn_groupprivate_size() != N)
96+
++Failed;
97+
if (IsFallback)
98+
++Failed;
99+
}
100+
101+
// Verify that the clause works when passing a zero size.
102+
#pragma omp target dyn_groupprivate(strict: 0) map(tofrom: Failed)
103+
{
104+
int IsFallback;
105+
if (omp_get_dyn_groupprivate_ptr(0, &IsFallback))
106+
++Failed;
107+
if (omp_get_dyn_groupprivate_size())
108+
++Failed;
109+
if (IsFallback)
110+
++Failed;
111+
}
112+
113+
// Verify that the clause works when passing a zero size.
114+
#pragma omp target dyn_groupprivate(fallback: 0) map(tofrom: Failed)
115+
{
116+
int IsFallback;
117+
if (omp_get_dyn_groupprivate_ptr(0, &IsFallback))
118+
++Failed;
119+
if (omp_get_dyn_groupprivate_size())
120+
++Failed;
121+
if (IsFallback)
122+
++Failed;
123+
}
124+
125+
// Verify that omitting the clause is the same as setting zero size.
126+
#pragma omp target map(tofrom: Failed)
127+
{
128+
int IsFallback;
129+
if (omp_get_dyn_groupprivate_ptr(0, &IsFallback))
130+
++Failed;
131+
if (omp_get_dyn_groupprivate_size())
132+
++Failed;
133+
if (IsFallback)
134+
++Failed;
135+
}
136+
137+
// CHECK: PASS
138+
if (!Failed)
139+
printf("PASS\n");
140+
}

0 commit comments

Comments
 (0)