[OpenMP][OpenACC] Implement `ompx_hold` map type modifier extension in Clang (1/2)
This patch implements Clang support for an original OpenMP extension
we have developed to support OpenACC: the `ompx_hold` map type
modifier. The next patch in this series, D106510, implements OpenMP
runtime support.
Consider the following example:
```
#pragma omp target data map(ompx_hold, tofrom: x) // holds onto mapping of x
{
foo(); // might have map(delete: x)
#pragma omp target map(present, alloc: x) // x is guaranteed to be present
printf("%d\n", x);
}
```
The `ompx_hold` map type modifier above specifies that the `target
data` directive holds onto the mapping for `x` throughout the
associated region regardless of any `target exit data` directives
executed during the call to `foo`. Thus, the presence assertion for
`x` at the enclosed `target` construct cannot fail. (As usual, the
standard OpenMP reference count for `x` must also reach zero before
the data is unmapped.)
Justification for inclusion in Clang and LLVM's OpenMP runtime:
* The `ompx_hold` modifier supports OpenACC functionality (structured
reference count) that cannot be achieved in standard OpenMP, as of
5.1.
* The runtime implementation for `ompx_hold` (next patch) will thus be
used by Flang's OpenACC support.
* The Clang implementation for `ompx_hold` (this patch) as well as the
runtime implementation are required for the Clang OpenACC support
being developed as part of the ECP Clacc project, which translates
OpenACC to OpenMP at the directive AST level. These patches are the
first step in upstreaming OpenACC functionality from Clacc.
* The Clang implementation for `ompx_hold` is also used by the tests
in the runtime implementation. That syntactic support makes the
tests more readable than low-level runtime calls can. Moreover,
upstream Flang and Clang do not yet support OpenACC syntax
sufficiently for writing the tests.
* More generally, the Clang implementation enables a clean separation
of concerns between OpenACC and OpenMP development in LLVM. That
is, LLVM's OpenMP developers can discuss, modify, and debug LLVM's
extended OpenMP implementation and test suite without directly
considering OpenACC's language and execution model, which can be
handled by LLVM's OpenACC developers.
* OpenMP users might find the `ompx_hold` modifier useful, as in the
above example.
See new documentation introduced by this patch in `openmp/docs` for
more detail on the functionality of this extension and its
relationship with OpenACC. For example, it explains how the runtime
must support two reference counts, as specified by OpenACC.
Clang recognizes `ompx_hold` unless `-fno-openmp-extensions`, a new
command-line option introduced by this patch, is specified.
Reviewed By: ABataev, jdoerfert, protze.joachim, grokos
Differential Revision: https://reviews.llvm.org/D106509
2021-09-01 03:17:07 +08:00
// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp -fno-openmp-extensions -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
// RUN: %clang_cc1 -verify=expected,lt45,lt50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
// RUN: %clang_cc1 -verify=expected,ge45,lt50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp -fno-openmp-extensions -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
// RUN: %clang_cc1 -verify=expected,ge45,ge50,ge51 -fopenmp -fno-openmp-extensions -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp-simd -fno-openmp-extensions -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
// RUN: %clang_cc1 -DCCODE -verify=expected,ge45,ge50,lt51 -fopenmp -fno-openmp-extensions -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized
2016-12-17 13:48:59 +08:00
# ifdef CCODE
void foo ( int arg ) {
const int n = 0 ;
double marr [ 10 ] [ 10 ] [ 10 ] ;
# pragma omp target teams map(marr[2][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr[:][0:][:])
{ }
# pragma omp target teams map(marr[:][1:][:]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr[:][n:][:])
{ }
}
# else
2019-08-28 22:55:08 +08:00
void xxx ( int argc ) {
int map ; // expected-note {{initialize the variable 'map' to silence this warning}}
# pragma omp target teams map(tofrom: map) // expected-warning {{variable 'map' is uninitialized when used here}}
for ( int i = 0 ; i < 10 ; + + i )
;
}
2016-12-17 13:48:59 +08:00
template < typename T , int I >
struct SA {
static int ss ;
# pragma omp threadprivate(ss) // expected-note {{defined as threadprivate or thread local}}
float a ;
int b [ 12 ] ;
float * c ;
T d ;
float e [ I ] ;
T * f ;
void func ( int arg ) {
# pragma omp target teams map(arg,a,d)
{ }
# pragma omp target teams map(arg[2:2],a,d) // expected-error {{subscripted value is not an array or pointer}}
{ }
2020-07-22 22:14:00 +08:00
// ge50-error@+2 {{expected addressable lvalue in 'map' clause}}
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(arg,a*2)
2016-12-17 13:48:59 +08:00
{ }
2020-07-22 22:14:00 +08:00
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(arg,(c+1)[2])
2016-12-17 13:48:59 +08:00
{ }
# pragma omp target teams map(arg,a[:2],d) // expected-error {{subscripted value is not an array or pointer}}
{ }
# pragma omp target teams map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}}
{ }
# pragma omp target teams map(to:ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
{ }
# pragma omp target teams map(to:b,e)
{ }
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(to:b,e) map(to:b) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
2016-12-17 13:48:59 +08:00
{ }
# pragma omp target teams map(to:b[:2],e)
{ }
# pragma omp target teams map(to:b,e[:])
{ }
# pragma omp target teams map(b[-1:]) // expected-error {{array section must be a subset of the original array}}
{ }
# pragma omp target teams map(b[:-1]) // expected-error {{section length is evaluated to a negative value -1}}
{ }
[APSInt][OpenMP] Fix isNegative, etc. for unsigned types
Without this patch, APSInt inherits APInt::isNegative, which merely
checks the sign bit without regard to whether the type is actually
signed. isNonNegative and isStrictlyPositive call isNegative and so
are also affected.
This patch adjusts APSInt to override isNegative, isNonNegative, and
isStrictlyPositive with implementations that consider whether the type
is signed.
A large set of Clang OpenMP tests are affected. Without this patch,
these tests assume that `true` is not a valid argument for clauses
like `collapse`. Indeed, `true` fails APInt::isStrictlyPositive but
not APSInt::isStrictlyPositive. This patch adjusts those tests to
assume `true` should be accepted.
This patch also adds tests revealing various other similar fixes due
to APSInt::isNegative calls in Clang's ExprConstant.cpp and
SemaExpr.cpp: `++` and `--` overflow in `constexpr`, evaluated object
size based on `alloc_size`, `<<` and `>>` shift count validation, and
OpenMP array section validation.
Reviewed By: lebedev.ri, ABataev, hfinkel
Differential Revision: https://reviews.llvm.org/D59712
llvm-svn: 359012
2019-04-24 01:04:15 +08:00
# pragma omp target teams map(b[true:true])
{ }
2016-12-17 13:48:59 +08:00
# pragma omp target teams map(always, tofrom: c,f)
{ }
# pragma omp target teams map(always, tofrom: c[1:2],f)
{ }
# pragma omp target teams map(always, tofrom: c,f[1:2])
{ }
# pragma omp target teams map(always, tofrom: c[:],f) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
{ }
# pragma omp target teams map(always, tofrom: c,f[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
{ }
return ;
}
} ;
struct SB {
unsigned A ;
unsigned B ;
float Arr [ 100 ] ;
float * Ptr ;
float * foo ( ) {
return & Arr [ 0 ] ;
}
} ;
struct SC {
unsigned A : 2 ;
unsigned B : 3 ;
unsigned C ;
unsigned D ;
float Arr [ 100 ] ;
SB S ;
SB ArrS [ 100 ] ;
SB * PtrS ;
SB * & RPtrS ;
float * Ptr ;
SC ( SB * & _RPtrS ) : RPtrS ( _RPtrS ) { }
} ;
union SD {
unsigned A ;
float B ;
} ;
void SAclient ( int arg ) {
SA < int , 123 > s ;
s . func ( arg ) ; // expected-note {{in instantiation of member function}}
double marr [ 10 ] [ 10 ] [ 10 ] ;
double marr2 [ 5 ] [ 10 ] [ 1 ] ;
double mvla [ 5 ] [ arg ] [ 10 ] ;
double * * * mptr ;
const int n = 0 ;
const int m = 1 ;
double mvla2 [ 5 ] [ arg ] [ m + n + 10 ] ;
SB * p ;
SD u ;
SC r ( p ) , t ( p ) ;
# pragma omp target teams map(r)
{ }
# pragma omp target teams map(marr[2][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr[:][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr[2][3][0:2])
{ }
# pragma omp target teams map(marr[:][:][:])
{ }
# pragma omp target teams map(marr[:2][:][:])
{ }
# pragma omp target teams map(marr[arg:][:][:])
{ }
# pragma omp target teams map(marr[arg:])
{ }
# pragma omp target teams map(marr[arg:][:arg][:]) // correct if arg is the size of dimension 2
{ }
# pragma omp target teams map(marr[:arg][:])
{ }
# pragma omp target teams map(marr[:arg][n:])
{ }
# pragma omp target teams map(marr[:][:arg][n:]) // correct if arg is the size of dimension 2
{ }
# pragma omp target teams map(marr[:][:m][n:]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr[n:m][:arg][n:])
{ }
# pragma omp target teams map(marr[:2][:1][:]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr[:2][1:][:]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr[:2][:][:1]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr[:2][:][1:]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr[:1][:2][:])
{ }
# pragma omp target teams map(marr[:1][0][:])
{ }
# pragma omp target teams map(marr[:arg][:2][:]) // correct if arg is 1
{ }
# pragma omp target teams map(marr[:1][3:1][:2])
{ }
# pragma omp target teams map(marr[:1][3:arg][:2]) // correct if arg is 1
{ }
# pragma omp target teams map(marr[:1][3:2][:2]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr[:2][:10][:])
{ }
# pragma omp target teams map(marr[:2][:][:5+5])
{ }
# pragma omp target teams map(marr[:2][2+2-4:][0:5+5])
{ }
# pragma omp target teams map(marr[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(marr2[:1][:2][0])
{ }
# pragma omp target teams map(mvla[:1][:][0]) // correct if the size of dimension 2 is 1.
{ }
# pragma omp target teams map(mvla[:2][:arg][:]) // correct if arg is the size of dimension 2.
{ }
# pragma omp target teams map(mvla[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(mvla[1][2:arg][:])
{ }
# pragma omp target teams map(mvla[:1][:][:])
{ }
# pragma omp target teams map(mvla2[:1][:2][:11])
{ }
# pragma omp target teams map(mvla2[:1][:2][:10]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(mptr[:1][:2-1][2:4-3])
{ }
# pragma omp target teams map(mptr[:1][:arg][2:4-3]) // correct if arg is 1.
{ }
# pragma omp target teams map(mptr[:1][:2-1][0:2])
{ }
# pragma omp target teams map(mptr[:1][:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(mptr[:1][:][0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
{ }
# pragma omp target teams map(mptr[:2][:1][0:2]) // expected-error {{array section does not specify contiguous storage}}
{ }
# pragma omp target teams map(r.ArrS[0].B)
{ }
# pragma omp target teams map(r.ArrS[:1].B) // expected-error {{OpenMP array section is not allowed here}}
{ }
# pragma omp target teams map(r.ArrS[:arg].B) // expected-error {{OpenMP array section is not allowed here}}
{ }
# pragma omp target teams map(r.ArrS[0].Arr[1:23])
{ }
# pragma omp target teams map(r.ArrS[0].Arr[1:arg])
{ }
# pragma omp target teams map(r.ArrS[0].Arr[arg:23])
{ }
# pragma omp target teams map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}}
{ }
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(r.ArrS[0].A, r.ArrS[1].A) // lt50-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} lt50-note {{used here}}
2016-12-17 13:48:59 +08:00
{ }
# pragma omp target teams map(r.ArrS[0].A, t.ArrS[1].A)
{ }
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(r.PtrS[0], r.PtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}}
2016-12-17 13:48:59 +08:00
{ }
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(r.PtrS, r.PtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
2018-02-28 01:42:00 +08:00
{ }
# pragma omp target teams map(r.PtrS->A, r.PtrS->B)
{ }
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}}
2018-02-28 01:42:00 +08:00
{ }
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(r.RPtrS, r.RPtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
2018-02-28 01:42:00 +08:00
{ }
# pragma omp target teams map(r.RPtrS->A, r.RPtrS->B)
2016-12-17 13:48:59 +08:00
{ }
# pragma omp target teams map(r.S.Arr[:12])
{ }
2020-07-22 22:14:00 +08:00
// ge50-error@+2 {{expected addressable lvalue in 'map' clause}}
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(r.S.foo()[:12])
2016-12-17 13:48:59 +08:00
{ }
# pragma omp target teams map(r.C, r.D)
{ }
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(r.C, r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
2016-12-17 13:48:59 +08:00
{ }
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(r.C) map(r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
2016-12-17 13:48:59 +08:00
{ }
# pragma omp target teams map(r.C, r.S) // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
{ }
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(r, r.S) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
2016-12-17 13:48:59 +08:00
{ }
# pragma omp target teams map(r.C, t.C)
{ }
# pragma omp target teams map(r.A) // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}}
{ }
# pragma omp target teams map(r.Arr)
{ }
# pragma omp target teams map(r.Arr[3:5])
{ }
# pragma omp target teams map(r.Ptr[3:5])
{ }
# pragma omp target teams map(r.ArrS[3:5].A) // expected-error {{OpenMP array section is not allowed here}}
{ }
# pragma omp target teams map(r.ArrS[3:5].Arr[6:7]) // expected-error {{OpenMP array section is not allowed here}}
{ }
# pragma omp target teams map(r.ArrS[3].Arr[6:7])
{ }
# pragma omp target teams map(r.S.Arr[4:5])
{ }
# pragma omp target teams map(r.S.Ptr[4:5])
{ }
# pragma omp target teams map(r.S.Ptr[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
{ }
2020-07-22 22:14:00 +08:00
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map((p+1)->A)
2016-12-17 13:48:59 +08:00
{ }
2017-12-06 23:04:36 +08:00
# pragma omp target teams map(u.B) // expected-error {{mapping of union members is not allowed}}
2016-12-17 13:48:59 +08:00
{ }
2021-02-11 21:10:54 +08:00
# pragma omp target data map(to: r.C) // lt50-note {{used here}}
2016-12-17 13:48:59 +08:00
{
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(r.D) // lt50-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
2016-12-17 13:48:59 +08:00
{ }
}
2021-02-11 21:10:54 +08:00
# pragma omp target data map(to: t.Ptr) // lt50-note {{used here}}
2016-12-17 13:48:59 +08:00
{
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(t.Ptr[:23]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
2016-12-17 13:48:59 +08:00
{ }
}
# pragma omp target data map(to: t.C, t.D)
{
# pragma omp target data map(to: t.C)
{
# pragma omp target teams map(t.D)
{ }
}
}
# pragma omp target data map(to: t)
{
# pragma omp target data map(to: t.C)
{
# pragma omp target teams map(t.D)
{ }
}
}
}
void foo ( ) {
}
bool foobool ( int argc ) {
return argc ;
}
2021-08-04 01:35:04 +08:00
struct S1 ; // expected-note 2 {{declared here}} // expected-note 3 {{forward declaration of 'S1'}}
2016-12-17 13:48:59 +08:00
extern S1 a ;
class S2 {
mutable int a ;
public :
S2 ( ) : a ( 0 ) { }
S2 ( S2 & s2 ) : a ( s2 . a ) { }
2017-09-13 19:12:35 +08:00
static float S2s ;
static const float S2sc ;
2016-12-17 13:48:59 +08:00
} ;
const float S2 : : S2sc = 0 ;
const S2 b ;
const S2 ba [ 5 ] ;
class S3 {
int a ;
public :
S3 ( ) : a ( 0 ) { }
S3 ( S3 & s3 ) : a ( s3 . a ) { }
} ;
const S3 c ;
const S3 ca [ 5 ] ;
extern const int f ;
class S4 {
int a ;
S4 ( ) ;
S4 ( const S4 & s4 ) ;
public :
S4 ( int v ) : a ( v ) { }
} ;
class S5 {
int a ;
S5 ( ) : a ( 0 ) { }
S5 ( const S5 & s5 ) : a ( s5 . a ) { }
public :
S5 ( int v ) : a ( v ) { }
} ;
template < class T >
struct S6 ;
template < >
2017-09-13 19:12:35 +08:00
struct S6 < int >
2016-12-17 13:48:59 +08:00
{
virtual void foo ( ) ;
} ;
S3 h ;
# pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
typedef int from ;
template < typename T , int I > // expected-note {{declared here}}
T tmain ( T argc ) {
const T d = 5 ;
const T da [ 5 ] = { 0 } ;
S4 e ( 4 ) ;
S5 g ( 5 ) ;
T i , t [ 20 ] ;
T & j = i ;
T * k = & j ;
T x ;
T y ;
T to , tofrom , always ;
const T ( & l ) [ 5 ] = da ;
# pragma omp target teams map // expected-error {{expected '(' after 'map'}}
{ }
# pragma omp target teams map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
{ }
# pragma omp target teams map() // expected-error {{expected expression}}
{ }
# pragma omp target teams map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
{ }
# pragma omp target teams map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
{ }
# pragma omp target teams map(to:) // expected-error {{expected expression}}
{ }
# pragma omp target teams map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
{ }
# pragma omp target teams map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
{ }
# pragma omp target teams map(x)
foo ( ) ;
# pragma omp target teams map(tofrom: t[:I])
foo ( ) ;
# pragma omp target teams map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
foo ( ) ;
# pragma omp target teams map(T) // expected-error {{'T' does not refer to a value}}
foo ( ) ;
2020-07-22 22:14:00 +08:00
// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(I)
2016-12-17 13:48:59 +08:00
foo ( ) ;
# pragma omp target teams map(S2::S2s)
foo ( ) ;
# pragma omp target teams map(S2::S2sc)
foo ( ) ;
# pragma omp target teams map(x)
foo ( ) ;
# pragma omp target teams map(to: x)
foo ( ) ;
# pragma omp target teams map(to: to)
foo ( ) ;
# pragma omp target teams map(to)
foo ( ) ;
# pragma omp target teams map(to, x)
foo ( ) ;
# pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
2020-07-22 22:14:00 +08:00
// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target data map(tofrom: argc > 0 ? x : y)
[OpenMP5.0] Allow pointer arithmetic in motion/map clause, by Chi Chun
Chen
Summary:
Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.
Reviewers: ABataev, jdoerfert
Reviewed By: ABataev
Subscribers: guansong, cfe-commits, sandoval, dreachem
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75077
2020-02-29 03:37:14 +08:00
2016-12-17 13:48:59 +08:00
# pragma omp target data map(argc)
# pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
2017-09-13 19:12:35 +08:00
# pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
# pragma omp target data map(ba)
2016-12-17 13:48:59 +08:00
# pragma omp target data map(ca)
# pragma omp target data map(da)
# pragma omp target data map(S2::S2s)
# pragma omp target data map(S2::S2sc)
# pragma omp target data map(e, g)
# pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
2021-02-11 21:10:54 +08:00
# pragma omp target data map(k) map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
# pragma omp target teams map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
2016-12-17 13:48:59 +08:00
foo ( ) ;
# pragma omp target data map(da)
# pragma omp target teams map(da[:4])
foo ( ) ;
2021-02-11 21:10:54 +08:00
# pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
# pragma omp target data map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
2016-12-17 13:48:59 +08:00
# pragma omp target data map(j)
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
2016-12-17 13:48:59 +08:00
foo ( ) ;
2021-02-11 21:10:54 +08:00
# pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
# pragma omp target data map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
2016-12-17 13:48:59 +08:00
# pragma omp target data map(j)
# pragma omp target teams map(l)
foo ( ) ;
# pragma omp target data map(always, tofrom: x)
# pragma omp target data map(always: x) // expected-error {{missing map type}}
[OpenMP][OpenACC] Implement `ompx_hold` map type modifier extension in Clang (1/2)
This patch implements Clang support for an original OpenMP extension
we have developed to support OpenACC: the `ompx_hold` map type
modifier. The next patch in this series, D106510, implements OpenMP
runtime support.
Consider the following example:
```
#pragma omp target data map(ompx_hold, tofrom: x) // holds onto mapping of x
{
foo(); // might have map(delete: x)
#pragma omp target map(present, alloc: x) // x is guaranteed to be present
printf("%d\n", x);
}
```
The `ompx_hold` map type modifier above specifies that the `target
data` directive holds onto the mapping for `x` throughout the
associated region regardless of any `target exit data` directives
executed during the call to `foo`. Thus, the presence assertion for
`x` at the enclosed `target` construct cannot fail. (As usual, the
standard OpenMP reference count for `x` must also reach zero before
the data is unmapped.)
Justification for inclusion in Clang and LLVM's OpenMP runtime:
* The `ompx_hold` modifier supports OpenACC functionality (structured
reference count) that cannot be achieved in standard OpenMP, as of
5.1.
* The runtime implementation for `ompx_hold` (next patch) will thus be
used by Flang's OpenACC support.
* The Clang implementation for `ompx_hold` (this patch) as well as the
runtime implementation are required for the Clang OpenACC support
being developed as part of the ECP Clacc project, which translates
OpenACC to OpenMP at the directive AST level. These patches are the
first step in upstreaming OpenACC functionality from Clacc.
* The Clang implementation for `ompx_hold` is also used by the tests
in the runtime implementation. That syntactic support makes the
tests more readable than low-level runtime calls can. Moreover,
upstream Flang and Clang do not yet support OpenACC syntax
sufficiently for writing the tests.
* More generally, the Clang implementation enables a clean separation
of concerns between OpenACC and OpenMP development in LLVM. That
is, LLVM's OpenMP developers can discuss, modify, and debug LLVM's
extended OpenMP implementation and test suite without directly
considering OpenACC's language and execution model, which can be
handled by LLVM's OpenACC developers.
* OpenMP users might find the `ompx_hold` modifier useful, as in the
above example.
See new documentation introduced by this patch in `openmp/docs` for
more detail on the functionality of this extension and its
relationship with OpenACC. For example, it explains how the runtime
must support two reference counts, as specified by OpenACC.
Clang recognizes `ompx_hold` unless `-fno-openmp-extensions`, a new
command-line option introduced by this patch, is specified.
Reviewed By: ABataev, jdoerfert, protze.joachim, grokos
Differential Revision: https://reviews.llvm.org/D106509
2021-09-01 03:17:07 +08:00
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
2020-07-22 22:14:00 +08:00
// expected-error@+1 {{missing map type}}
# pragma omp target data map(tofrom, always: x)
2016-12-17 13:48:59 +08:00
# pragma omp target data map(always, tofrom: always, tofrom, x)
# pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
foo ( ) ;
return 0 ;
}
int main ( int argc , char * * argv ) {
const int d = 5 ;
const int da [ 5 ] = { 0 } ;
S4 e ( 4 ) ;
S5 g ( 5 ) ;
int i ;
int & j = i ;
int * k = & j ;
S6 < int > m ;
int x ;
int y ;
int to , tofrom , always ;
const int ( & l ) [ 5 ] = da ;
2020-07-22 22:14:00 +08:00
// expected-error@+3 {{expected '(' after 'map'}}
// ge50-error@+2 {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
// lt50-error@+1 {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}}
# pragma omp target data map
2016-12-17 13:48:59 +08:00
# pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
# pragma omp target data map() // expected-error {{expected expression}}
# pragma omp target data map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
# pragma omp target data map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
# pragma omp target data map(to:) // expected-error {{expected expression}}
# pragma omp target data map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
# pragma omp target data map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
# pragma omp target teams map(x)
foo ( ) ;
# pragma omp target teams map(to: x)
foo ( ) ;
# pragma omp target teams map(to: to)
foo ( ) ;
# pragma omp target teams map(to)
foo ( ) ;
# pragma omp target teams map(to, x)
foo ( ) ;
# pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
2020-07-22 22:14:00 +08:00
// ge50-error@+2 {{expected addressable lvalue in 'map' clause}}
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target data map(tofrom: argc > 0 ? argv[1] : argv[2])
2016-12-17 13:48:59 +08:00
# pragma omp target data map(argc)
# pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
2017-09-13 19:12:35 +08:00
# pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
2016-12-17 13:48:59 +08:00
# pragma omp target data map(argv[1])
2017-09-13 19:12:35 +08:00
# pragma omp target data map(ba)
2016-12-17 13:48:59 +08:00
# pragma omp target data map(ca)
# pragma omp target data map(da)
# pragma omp target data map(S2::S2s)
# pragma omp target data map(S2::S2sc)
# pragma omp target data map(e, g)
# pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
2021-02-11 21:10:54 +08:00
# pragma omp target data map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
# pragma omp target teams map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
2016-12-17 13:48:59 +08:00
foo ( ) ;
# pragma omp target data map(da)
# pragma omp target teams map(da[:4])
foo ( ) ;
2021-02-11 21:10:54 +08:00
# pragma omp target data map(k, j, l) // lt50-note {{used here}}
# pragma omp target data map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
2016-12-17 13:48:59 +08:00
# pragma omp target data map(j)
2021-02-11 21:10:54 +08:00
# pragma omp target teams map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
2016-12-17 13:48:59 +08:00
foo ( ) ;
2021-02-11 21:10:54 +08:00
# pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
# pragma omp target data map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
2016-12-17 13:48:59 +08:00
# pragma omp target data map(j)
# pragma omp target teams map(l)
foo ( ) ;
# pragma omp target data map(always, tofrom: x)
# pragma omp target data map(always: x) // expected-error {{missing map type}}
[OpenMP][OpenACC] Implement `ompx_hold` map type modifier extension in Clang (1/2)
This patch implements Clang support for an original OpenMP extension
we have developed to support OpenACC: the `ompx_hold` map type
modifier. The next patch in this series, D106510, implements OpenMP
runtime support.
Consider the following example:
```
#pragma omp target data map(ompx_hold, tofrom: x) // holds onto mapping of x
{
foo(); // might have map(delete: x)
#pragma omp target map(present, alloc: x) // x is guaranteed to be present
printf("%d\n", x);
}
```
The `ompx_hold` map type modifier above specifies that the `target
data` directive holds onto the mapping for `x` throughout the
associated region regardless of any `target exit data` directives
executed during the call to `foo`. Thus, the presence assertion for
`x` at the enclosed `target` construct cannot fail. (As usual, the
standard OpenMP reference count for `x` must also reach zero before
the data is unmapped.)
Justification for inclusion in Clang and LLVM's OpenMP runtime:
* The `ompx_hold` modifier supports OpenACC functionality (structured
reference count) that cannot be achieved in standard OpenMP, as of
5.1.
* The runtime implementation for `ompx_hold` (next patch) will thus be
used by Flang's OpenACC support.
* The Clang implementation for `ompx_hold` (this patch) as well as the
runtime implementation are required for the Clang OpenACC support
being developed as part of the ECP Clacc project, which translates
OpenACC to OpenMP at the directive AST level. These patches are the
first step in upstreaming OpenACC functionality from Clacc.
* The Clang implementation for `ompx_hold` is also used by the tests
in the runtime implementation. That syntactic support makes the
tests more readable than low-level runtime calls can. Moreover,
upstream Flang and Clang do not yet support OpenACC syntax
sufficiently for writing the tests.
* More generally, the Clang implementation enables a clean separation
of concerns between OpenACC and OpenMP development in LLVM. That
is, LLVM's OpenMP developers can discuss, modify, and debug LLVM's
extended OpenMP implementation and test suite without directly
considering OpenACC's language and execution model, which can be
handled by LLVM's OpenACC developers.
* OpenMP users might find the `ompx_hold` modifier useful, as in the
above example.
See new documentation introduced by this patch in `openmp/docs` for
more detail on the functionality of this extension and its
relationship with OpenACC. For example, it explains how the runtime
must support two reference counts, as specified by OpenACC.
Clang recognizes `ompx_hold` unless `-fno-openmp-extensions`, a new
command-line option introduced by this patch, is specified.
Reviewed By: ABataev, jdoerfert, protze.joachim, grokos
Differential Revision: https://reviews.llvm.org/D106509
2021-09-01 03:17:07 +08:00
// ge51-error@+3 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper', 'present'}}
// lt51-error@+2 {{incorrect map type modifier, expected one of: 'always', 'close', 'mapper'}}
2020-07-22 22:14:00 +08:00
// expected-error@+1 {{missing map type}}
# pragma omp target data map(tofrom, always: x)
2016-12-17 13:48:59 +08:00
# pragma omp target data map(always, tofrom: always, tofrom, x)
# pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
foo ( ) ;
2020-07-22 22:14:00 +08:00
// lt50-error@+2 {{private variable cannot be in a map clause in '#pragma omp target teams' directive}}
// lt50-note@+1 {{defined as private}}
# pragma omp target teams private(j) map(j)
2016-12-17 13:48:59 +08:00
{ }
2020-07-22 22:14:00 +08:00
// lt50-error@+2 {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}}
// lt50-note@+1 {{defined as firstprivate}}
# pragma omp target teams firstprivate(j) map(j)
2016-12-17 13:48:59 +08:00
{ }
2017-09-13 19:12:35 +08:00
# pragma omp target teams map(m)
2016-12-17 13:48:59 +08:00
{ }
[OpenMP5.0] Allow pointer arithmetic in motion/map clause, by Chi Chun
Chen
Summary:
Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.
Reviewers: ABataev, jdoerfert
Reviewed By: ABataev
Subscribers: guansong, cfe-commits, sandoval, dreachem
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75077
2020-02-29 03:37:14 +08:00
int * * BB , * offset , * a ;
2020-07-22 22:14:00 +08:00
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(**(BB+*offset))
[OpenMP5.0] Allow pointer arithmetic in motion/map clause, by Chi Chun
Chen
Summary:
Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.
Reviewers: ABataev, jdoerfert
Reviewed By: ABataev
Subscribers: guansong, cfe-commits, sandoval, dreachem
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75077
2020-02-29 03:37:14 +08:00
{ }
2020-07-22 22:14:00 +08:00
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(**(BB+y))
[OpenMP5.0] Allow pointer arithmetic in motion/map clause, by Chi Chun
Chen
Summary:
Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.
Reviewers: ABataev, jdoerfert
Reviewed By: ABataev
Subscribers: guansong, cfe-commits, sandoval, dreachem
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75077
2020-02-29 03:37:14 +08:00
{ }
2020-07-22 22:14:00 +08:00
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(*(a+*offset))
[OpenMP5.0] Allow pointer arithmetic in motion/map clause, by Chi Chun
Chen
Summary:
Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.
Reviewers: ABataev, jdoerfert
Reviewed By: ABataev
Subscribers: guansong, cfe-commits, sandoval, dreachem
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75077
2020-02-29 03:37:14 +08:00
{ }
2020-07-22 22:14:00 +08:00
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(**(*offset+BB))
[OpenMP5.0] Allow pointer arithmetic in motion/map clause, by Chi Chun
Chen
Summary:
Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.
Reviewers: ABataev, jdoerfert
Reviewed By: ABataev
Subscribers: guansong, cfe-commits, sandoval, dreachem
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75077
2020-02-29 03:37:14 +08:00
{ }
2020-07-22 22:14:00 +08:00
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(**(y+BB))
[OpenMP5.0] Allow pointer arithmetic in motion/map clause, by Chi Chun
Chen
Summary:
Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.
Reviewers: ABataev, jdoerfert
Reviewed By: ABataev
Subscribers: guansong, cfe-commits, sandoval, dreachem
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75077
2020-02-29 03:37:14 +08:00
{ }
2020-07-22 22:14:00 +08:00
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(*(*offset+a))
[OpenMP5.0] Allow pointer arithmetic in motion/map clause, by Chi Chun
Chen
Summary:
Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.
Reviewers: ABataev, jdoerfert
Reviewed By: ABataev
Subscribers: guansong, cfe-commits, sandoval, dreachem
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75077
2020-02-29 03:37:14 +08:00
{ }
2020-07-22 22:14:00 +08:00
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(**(*offset+BB+*a))
[OpenMP5.0] Allow pointer arithmetic in motion/map clause, by Chi Chun
Chen
Summary:
Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.
Reviewers: ABataev, jdoerfert
Reviewed By: ABataev
Subscribers: guansong, cfe-commits, sandoval, dreachem
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75077
2020-02-29 03:37:14 +08:00
{ }
2020-07-22 22:14:00 +08:00
// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
# pragma omp target teams map(**(*(*(&offset))+BB+*a))
[OpenMP5.0] Allow pointer arithmetic in motion/map clause, by Chi Chun
Chen
Summary:
Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.
Reviewers: ABataev, jdoerfert
Reviewed By: ABataev
Subscribers: guansong, cfe-commits, sandoval, dreachem
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75077
2020-02-29 03:37:14 +08:00
{ }
# pragma omp target teams map(*(a+(a))) // expected-error {{invalid operands to binary expression ('int *' and 'int *')}}
{ }
# pragma omp target teams map(*(1+*a+*a)) // expected-error {{indirection requires pointer operand ('int' invalid)}}
{ }
2020-04-07 17:18:44 +08:00
# pragma omp target teams map(delete: j) // expected-error {{map type 'delete' is not allowed for '#pragma omp target teams'}}
{ }
# pragma omp target teams map(release: j) // expected-error {{map type 'release' is not allowed for '#pragma omp target teams'}}
{ }
2016-12-17 13:48:59 +08:00
return tmain < int , 3 > ( argc ) + tmain < from , 4 > ( argc ) ; // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
}
# endif