Message ID | 20150717130559.GI1780@tucnak.redhat.com |
---|---|
State | New |
Headers | show |
Jakub, On 07/17/2015 08:05 AM, Jakub Jelinek wrote: > Hi! > > ... > > I believe OpenACC has something similar, but no idea if it is already > implemented. Yes, it is implemented in gomp-4_0-branch. While the purpose for 'omp declare target' and 'acc declare' are similar, the data movement, via the clauses, provided with the latter make it very different than the former. The data movement requires that data be moved at the entry and exit of an 'associated region'. Associated region to mean either a function, subroutine, entire program or Fortran module. I choose to implement this in the front-ends. For discussion purposes, I'll use the C front-end: c_parser_oacc_declare and finish_oacc_declare. As far as the syntax, OpenMP is alot easier to deal with than OpenACC. The handling of said is reflected in c_parser_oacc_declare. Here also is the handling of the numerous data movement clauses. One in particular requires special handling: create. This can be seen toward the end of the function, There is a libgomp component GOACC_register_static (oacc-parallel.c) that is used in conjunction with the create clause. The creation and deletion of the 'associated region' is done in finish_oacc_declare. Depending upon where the directive was found requires different handling, i.e., global variable scope versus local variable scope. In addition, if there is data movement from target -> host, this must be handled appropriately. > >... > > Ilya, Thomas, thoughts on this? > Jim answering at the behest of Thomas.... If the above explanation is not sufficient please yell. It may make more sense to carve out the code in question and document it more thoroughly for discussion purposes. Also the implementation approach in the front-ends may be entirely wrong. There may be an approach to do it in the 'middle'. However, my lack of experience in the middle may have caused me to go down the wrong path. Jim
On Fri, Jul 17, 2015 at 15:05:59 +0200, Jakub Jelinek wrote: > As the testcases show, #pragma omp declare target has now a new form (well, > two; with some issues on it pending), where it is used just as a single > declarative directive rather than a pair of them and allows marking > vars and functions by name as "omp declare target" vars/functions (which the > middle-end etc. already handles), but also "omp declare target link", which > is a deferred var, that is not initially mapped (on devices without shared > memory with host), but has to be mapped explicitly. I don't quite understand how link should work. OpenMP 4.5 says: "The list items of a link clause are not mapped by the declare target directive. Instead, their mapping is deferred until they are mapped by target data or target constructs. They are mapped only for such regions." But doesn't this mean that the example bellow should work identically with/without USE_LINK defined? Or is there some difference on other testcases? int a = 1; #ifdef USE_LINK #pragma omp declare target link(a) #endif int main () { a = 2; int res; #pragma omp target map(to: a) map(from: res) res = a; return res; } > This patch only marks them with the new attribute, the actual middle-end > implementation needs to be implemented. > > I believe OpenACC has something similar, but no idea if it is already > implemented. > > Anyway, I think the implementation should be that in some pass running on > the ACCEL_COMPILER side (guarded by separate address space aka non-HSA) HSA does not define ACCEL_COMPILER, because it uses only one compiler. > we actually replace the variables with pointers to variables, then need > to somehow also mark those in the offloading tables, so that the library I see 2 possible options: use the MSB of the size, or introduce the third field for flags. > registers them (the locations of the pointers to the vars), but also marks > them for special treatment, and then when actually trying to map them > (or their parts, guess that needs to be discussed) we allocate them or > whatever is requested and store the device pointer into the corresponding > variable. > > Ilya, Thomas, thoughts on this? -- Ilya
On Mon, Oct 26, 2015 at 09:35:52PM +0300, Ilya Verbin wrote: > On Fri, Jul 17, 2015 at 15:05:59 +0200, Jakub Jelinek wrote: > > As the testcases show, #pragma omp declare target has now a new form (well, > > two; with some issues on it pending), where it is used just as a single > > declarative directive rather than a pair of them and allows marking > > vars and functions by name as "omp declare target" vars/functions (which the > > middle-end etc. already handles), but also "omp declare target link", which > > is a deferred var, that is not initially mapped (on devices without shared > > memory with host), but has to be mapped explicitly. > > I don't quite understand how link should work. OpenMP 4.5 says: > > "The list items of a link clause are not mapped by the declare target directive. > Instead, their mapping is deferred until they are mapped by target data or > target constructs. They are mapped only for such regions." > > But doesn't this mean that the example bellow should work identically > with/without USE_LINK defined? Or is there some difference on other testcases? On your testcase, the end result is pretty much the same, the variable is not mapped initially to the device, and at the beginning of omp target it is mapped to device, at the end of the region it is unmapped from the device (without copying back). But consider: int a = 1, b = 1; #pragma omp declare target link (a) to (b) int foo (void) { return a++ + b++; } #pragma omp declare target to (foo) int main () { a = 2; b = 2; int res; #pragma omp target map (to: a, b) map (from: res) { res = foo () + foo (); } // This assumes only non-shared address space, so would need to be guarded // for that. if (res != (2 + 1) + (3 + 2)) __builtin_abort (); return 0; } Without declare target link or to, you can't use the global variables in orphaned accelerated routines (unless you e.g. take the address of the mapped variable in the region and pass it around). The to variables (non-deferred) are always mapped and are initialized with the original initializer, refcount is infinity. link (deferred) work more like the normal mapping, referencing those vars when they aren't explicitly (or implicitly) mapped is unspecified behavior, if it is e.g. mapped freshly with to kind, it gets the current value of the host var rather than the original one. But, beyond the mapping the compiler needs to ensure that all uses of the link global var (or perhaps just all uses of the link global var outside of the target construct body where it is mapped, because you could use there the pointer you got from GOMP_target) are replaced by dereference of some artificial pointer, so a becomes *a_tmp and &a becomes &*a_tmp, and that the runtime library during registration of the tables is told about the address of this artificial pointer. During registration, I'd expect it would stick an entry for this range into the table, with some special flag or something similar, indicating that it is deferred mapping and where the offloading device pointer is. During mapping, it would map it as any other not yet mapped object, but additionally would also set this device pointer to the device address of the mapped object. We also need to ensure that when we drop the refcount of that mapping back to 0, we get it back to the state where it is described as a range with registered deferred mapping and where the device pointer is. > > This patch only marks them with the new attribute, the actual middle-end > > implementation needs to be implemented. > > > > I believe OpenACC has something similar, but no idea if it is already > > implemented. > > > > Anyway, I think the implementation should be that in some pass running on > > the ACCEL_COMPILER side (guarded by separate address space aka non-HSA) > > HSA does not define ACCEL_COMPILER, because it uses only one compiler. HSA is a non-issue here, as it has shared address space, therefore map clause does nothing, declare target to or link clauses also don't do anything. > > we actually replace the variables with pointers to variables, then need > > to somehow also mark those in the offloading tables, so that the library > > I see 2 possible options: use the MSB of the size, or introduce the third field > for flags. Well, it can be either recorded in the host variable tables (which contain address and size pair, right), or in corresponding offloading device table (which contains the pointer, something else?). Jakub
On Mon, Oct 26, 2015 at 20:05:39 +0100, Jakub Jelinek wrote: > On Mon, Oct 26, 2015 at 09:35:52PM +0300, Ilya Verbin wrote: > > On Fri, Jul 17, 2015 at 15:05:59 +0200, Jakub Jelinek wrote: > > > As the testcases show, #pragma omp declare target has now a new form (well, > > > two; with some issues on it pending), where it is used just as a single > > > declarative directive rather than a pair of them and allows marking > > > vars and functions by name as "omp declare target" vars/functions (which the > > > middle-end etc. already handles), but also "omp declare target link", which > > > is a deferred var, that is not initially mapped (on devices without shared > > > memory with host), but has to be mapped explicitly. > > > > I don't quite understand how link should work. OpenMP 4.5 says: > > > > "The list items of a link clause are not mapped by the declare target directive. > > Instead, their mapping is deferred until they are mapped by target data or > > target constructs. They are mapped only for such regions." > > > > But doesn't this mean that the example bellow should work identically > > with/without USE_LINK defined? Or is there some difference on other testcases? > > On your testcase, the end result is pretty much the same, the variable is > not mapped initially to the device, and at the beginning of omp target it is > mapped to device, at the end of the region it is unmapped from the device > (without copying back). > > But consider: > > int a = 1, b = 1; > #pragma omp declare target link (a) to (b) > int > foo (void) > { > return a++ + b++; > } > #pragma omp declare target to (foo) > int > main () > { > a = 2; > b = 2; > int res; > #pragma omp target map (to: a, b) map (from: res) > { > res = foo () + foo (); > } > // This assumes only non-shared address space, so would need to be guarded > // for that. > if (res != (2 + 1) + (3 + 2)) > __builtin_abort (); > return 0; > } > > Without declare target link or to, you can't use the global variables > in orphaned accelerated routines (unless you e.g. take the address of the > mapped variable in the region and pass it around). > The to variables (non-deferred) are always mapped and are initialized with > the original initializer, refcount is infinity. link (deferred) work more > like the normal mapping, referencing those vars when they aren't explicitly > (or implicitly) mapped is unspecified behavior, if it is e.g. mapped freshly > with to kind, it gets the current value of the host var rather than the > original one. But, beyond the mapping the compiler needs to ensure that > all uses of the link global var (or perhaps just all uses of the link global > var outside of the target construct body where it is mapped, because you > could use there the pointer you got from GOMP_target) are replaced by > dereference of some artificial pointer, so a becomes *a_tmp and &a becomes > &*a_tmp, and that the runtime library during registration of the tables is > told about the address of this artificial pointer. During registration, > I'd expect it would stick an entry for this range into the table, with some > special flag or something similar, indicating that it is deferred mapping > and where the offloading device pointer is. During mapping, it would map it > as any other not yet mapped object, but additionally would also set this > device pointer to the device address of the mapped object. We also need to > ensure that when we drop the refcount of that mapping back to 0, we get it > back to the state where it is described as a range with registered deferred > mapping and where the device pointer is. Ok, got it, I'll try implement this... > > > we actually replace the variables with pointers to variables, then need > > > to somehow also mark those in the offloading tables, so that the library > > > > I see 2 possible options: use the MSB of the size, or introduce the third field > > for flags. > > Well, it can be either recorded in the host variable tables (which contain > address and size pair, right), or in corresponding offloading device table > (which contains the pointer, something else?). It contains a size too, which is checked in libgomp: gomp_fatal ("Can't map target variables (size mismatch)"); Yes, we can remove this check, and use second field in device table for flags. -- Ilya
On Mon, Oct 26, 2015 at 10:39:04PM +0300, Ilya Verbin wrote: > > Without declare target link or to, you can't use the global variables > > in orphaned accelerated routines (unless you e.g. take the address of the > > mapped variable in the region and pass it around). > > The to variables (non-deferred) are always mapped and are initialized with > > the original initializer, refcount is infinity. link (deferred) work more > > like the normal mapping, referencing those vars when they aren't explicitly > > (or implicitly) mapped is unspecified behavior, if it is e.g. mapped freshly > > with to kind, it gets the current value of the host var rather than the > > original one. But, beyond the mapping the compiler needs to ensure that > > all uses of the link global var (or perhaps just all uses of the link global > > var outside of the target construct body where it is mapped, because you > > could use there the pointer you got from GOMP_target) are replaced by > > dereference of some artificial pointer, so a becomes *a_tmp and &a becomes > > &*a_tmp, and that the runtime library during registration of the tables is > > told about the address of this artificial pointer. During registration, > > I'd expect it would stick an entry for this range into the table, with some > > special flag or something similar, indicating that it is deferred mapping > > and where the offloading device pointer is. During mapping, it would map it > > as any other not yet mapped object, but additionally would also set this > > device pointer to the device address of the mapped object. We also need to > > ensure that when we drop the refcount of that mapping back to 0, we get it > > back to the state where it is described as a range with registered deferred > > mapping and where the device pointer is. > > Ok, got it, I'll try implement this... Thanks. > > > > we actually replace the variables with pointers to variables, then need > > > > to somehow also mark those in the offloading tables, so that the library > > > > > > I see 2 possible options: use the MSB of the size, or introduce the third field > > > for flags. > > > > Well, it can be either recorded in the host variable tables (which contain > > address and size pair, right), or in corresponding offloading device table > > (which contains the pointer, something else?). > > It contains a size too, which is checked in libgomp: > gomp_fatal ("Can't map target variables (size mismatch)"); > Yes, we can remove this check, and use second field in device table for flags. Yeah, or e.g. just use MSB of that size (so check that either the size is the same (then it is target to) or it is MSB | size (then it is target link). Objects larger than half of the address space aren't really supportable anyway. Jakub
Hi Jakub! On Fri, 17 Jul 2015 15:05:59 +0200, Jakub Jelinek <jakub@redhat.com> wrote: > [...] "omp declare target link" [...] > This patch only marks them with the new attribute, [...] > --- gcc/c/c-parser.c.jj 2015-07-16 18:09:25.000000000 +0200 > +++ gcc/c/c-parser.c 2015-07-17 14:11:08.553694975 +0200 > static void > c_parser_omp_declare_target (c_parser *parser) > { > [...] > + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) > + { > + tree t = OMP_CLAUSE_DECL (c), id; > + tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t)); > + tree at2 = lookup_attribute ("omp declare target link", > + DECL_ATTRIBUTES (t)); > + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK) > + { > + id = get_identifier ("omp declare target link"); > + std::swap (at1, at2); > + } > + else > + id = get_identifier ("omp declare target"); Is it intentional that you didn't add "omp declare target link" to gcc/c-family/c-common.c:c_common_attribute_table, next to the existing "omp declare target"? Grüße Thomas
On Mon, Nov 23, 2015 at 12:31:24PM +0100, Thomas Schwinge wrote: > Hi Jakub! > > On Fri, 17 Jul 2015 15:05:59 +0200, Jakub Jelinek <jakub@redhat.com> wrote: > > [...] "omp declare target link" [...] > > > This patch only marks them with the new attribute, [...] > > > --- gcc/c/c-parser.c.jj 2015-07-16 18:09:25.000000000 +0200 > > +++ gcc/c/c-parser.c 2015-07-17 14:11:08.553694975 +0200 > > > static void > > c_parser_omp_declare_target (c_parser *parser) > > { > > [...] > > + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) > > + { > > + tree t = OMP_CLAUSE_DECL (c), id; > > + tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t)); > > + tree at2 = lookup_attribute ("omp declare target link", > > + DECL_ATTRIBUTES (t)); > > + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK) > > + { > > + id = get_identifier ("omp declare target link"); > > + std::swap (at1, at2); > > + } > > + else > > + id = get_identifier ("omp declare target"); > > Is it intentional that you didn't add "omp declare target link" to > gcc/c-family/c-common.c:c_common_attribute_table, next to the existing > "omp declare target"? No. But the link attribute support is still unfinished, Ilya is working on the support. Jakub
--- gcc/tree-core.h.jj 2015-07-15 13:02:31.000000000 +0200 +++ gcc/tree-core.h 2015-07-17 09:30:44.944431669 +0200 @@ -256,6 +256,13 @@ enum omp_clause_code { /* OpenMP clause: uniform (argument-list). */ OMP_CLAUSE_UNIFORM, + /* OpenMP clause: to (extended-list). + Only when it appears in declare target. */ + OMP_CLAUSE_TO_DECLARE, + + /* OpenMP clause: link (variable-list). */ + OMP_CLAUSE_LINK, + /* OpenMP clause: from (variable-list). */ OMP_CLAUSE_FROM, --- gcc/tree.c.jj 2015-07-14 14:49:57.000000000 +0200 +++ gcc/tree.c 2015-07-17 09:33:51.270692623 +0200 @@ -288,6 +288,8 @@ unsigned const char omp_clause_num_ops[] 2, /* OMP_CLAUSE_ALIGNED */ 1, /* OMP_CLAUSE_DEPEND */ 1, /* OMP_CLAUSE_UNIFORM */ + 1, /* OMP_CLAUSE_TO_DECLARE */ + 1, /* OMP_CLAUSE_LINK */ 2, /* OMP_CLAUSE_FROM */ 2, /* OMP_CLAUSE_TO */ 2, /* OMP_CLAUSE_MAP */ @@ -357,6 +359,8 @@ const char * const omp_clause_code_name[ "aligned", "depend", "uniform", + "to", + "link", "from", "to", "map", @@ -11392,6 +11396,8 @@ walk_tree_1 (tree *tp, walk_tree_fn func case OMP_CLAUSE_GRAINSIZE: case OMP_CLAUSE_NUM_TASKS: case OMP_CLAUSE_HINT: + case OMP_CLAUSE_TO_DECLARE: + case OMP_CLAUSE_LINK: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE__LOOPTEMP_: --- gcc/tree-nested.c.jj 2015-07-14 14:49:57.000000000 +0200 +++ gcc/tree-nested.c 2015-07-17 09:35:11.905507270 +0200 @@ -1098,6 +1098,8 @@ convert_nonlocal_omp_clauses (tree *pcla case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_SHARED: + case OMP_CLAUSE_TO_DECLARE: + case OMP_CLAUSE_LINK: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: @@ -1745,6 +1747,8 @@ convert_local_omp_clauses (tree *pclause case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_SHARED: + case OMP_CLAUSE_TO_DECLARE: + case OMP_CLAUSE_LINK: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: --- gcc/tree-pretty-print.c.jj 2015-07-15 13:02:31.000000000 +0200 +++ gcc/tree-pretty-print.c 2015-07-17 09:36:30.822347172 +0200 @@ -344,6 +344,12 @@ dump_omp_clause (pretty_printer *pp, tre case OMP_CLAUSE_USE_DEVICE: name = "use_device"; goto print_remap; + case OMP_CLAUSE_TO_DECLARE: + name = "to"; + goto print_remap; + case OMP_CLAUSE_LINK: + name = "link"; + goto print_remap; print_remap: pp_string (pp, name); pp_left_paren (pp); --- gcc/c-family/c-pragma.h.jj 2015-07-14 14:49:57.000000000 +0200 +++ gcc/c-family/c-pragma.h 2015-07-17 09:21:03.190983600 +0200 @@ -101,6 +101,7 @@ typedef enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR, PRAGMA_OMP_CLAUSE_LASTPRIVATE, PRAGMA_OMP_CLAUSE_LINEAR, + PRAGMA_OMP_CLAUSE_LINK, PRAGMA_OMP_CLAUSE_MAP, PRAGMA_OMP_CLAUSE_MERGEABLE, PRAGMA_OMP_CLAUSE_NOGROUP, --- gcc/c/c-parser.c.jj 2015-07-16 18:09:25.000000000 +0200 +++ gcc/c/c-parser.c 2015-07-17 14:11:08.553694975 +0200 @@ -9953,6 +9953,8 @@ c_parser_omp_clause_name (c_parser *pars result = PRAGMA_OMP_CLAUSE_LASTPRIVATE; else if (!strcmp ("linear", p)) result = PRAGMA_OMP_CLAUSE_LINEAR; + else if (!strcmp ("link", p)) + result = PRAGMA_OMP_CLAUSE_LINK; break; case 'm': if (!strcmp ("map", p)) @@ -10235,7 +10237,7 @@ c_parser_omp_variable_list (c_parser *pa && !TREE_READONLY (low_bound)) { error_at (clause_loc, - "%qD is not a constant", low_bound); + "%qD is not a constant", low_bound); t = error_mark_node; } @@ -10243,7 +10245,7 @@ c_parser_omp_variable_list (c_parser *pa && !TREE_READONLY (length)) { error_at (clause_loc, - "%qD is not a constant", length); + "%qD is not a constant", length); t = error_mark_node; } } @@ -12600,8 +12602,18 @@ c_parser_omp_all_clauses (c_parser *pars if (!first) goto clause_not_first; break; + case PRAGMA_OMP_CLAUSE_LINK: + clauses + = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_LINK, clauses); + c_name = "link"; + break; case PRAGMA_OMP_CLAUSE_TO: - clauses = c_parser_omp_clause_to (parser, clauses); + if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LINK)) != 0) + clauses + = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO_DECLARE, + clauses); + else + clauses = c_parser_omp_clause_to (parser, clauses); c_name = "to"; break; case PRAGMA_OMP_CLAUSE_FROM: @@ -15313,13 +15325,64 @@ c_finish_omp_declare_simd (c_parser *par /* OpenMP 4.0: # pragma omp declare target new-line declarations and definitions - # pragma omp end declare target new-line */ + # pragma omp end declare target new-line + + OpenMP 4.1: + # pragma omp declare target ( extended-list ) new-line + + # pragma omp declare target declare-target-clauses[seq] new-line */ + +#define OMP_DECLARE_TARGET_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_TO) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LINK)) static void c_parser_omp_declare_target (c_parser *parser) { - c_parser_skip_to_pragma_eol (parser); - current_omp_declare_target_attribute++; + location_t loc = c_parser_peek_token (parser)->location; + tree clauses = NULL_TREE; + if (c_parser_next_token_is (parser, CPP_NAME)) + clauses = c_parser_omp_all_clauses (parser, OMP_DECLARE_TARGET_CLAUSE_MASK, + "#pragma omp declare target"); + else if (c_parser_next_token_is (parser, CPP_OPEN_PAREN)) + { + clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO_DECLARE, + clauses); + c_parser_skip_to_pragma_eol (parser); + } + else + { + c_parser_skip_to_pragma_eol (parser); + current_omp_declare_target_attribute++; + return; + } + if (current_omp_declare_target_attribute) + error_at (loc, "%<#pragma omp declare target%> with clauses in between " + "%<#pragma omp declare target%> without clauses and " + "%<#pragma omp end declare target%>"); + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + tree t = OMP_CLAUSE_DECL (c), id; + tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t)); + tree at2 = lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (t)); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK) + { + id = get_identifier ("omp declare target link"); + std::swap (at1, at2); + } + else + id = get_identifier ("omp declare target"); + if (at2) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD specified both in declare target %<link%> and %<to%>" + " clauses", t); + continue; + } + if (!at1) + DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); + } } static void --- gcc/c/c-typeck.c.jj 2015-07-15 13:00:32.000000000 +0200 +++ gcc/c/c-typeck.c 2015-07-17 13:06:58.297769199 +0200 @@ -12576,6 +12576,36 @@ c_finish_omp_clauses (tree clauses, bool bitmap_set_bit (&map_head, DECL_UID (t)); break; + case OMP_CLAUSE_TO_DECLARE: + t = OMP_CLAUSE_DECL (c); + if (TREE_CODE (t) == FUNCTION_DECL) + break; + /* FALLTHRU */ + case OMP_CLAUSE_LINK: + t = OMP_CLAUSE_DECL (c); + if (!VAR_P (t)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE is not a variable in clause %qs", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + else if (DECL_THREAD_LOCAL_P (t)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD is threadprivate variable in %qs clause", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + else if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD does not have a mappable type in %qs clause", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + break; + case OMP_CLAUSE_UNIFORM: t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) != PARM_DECL) --- gcc/cp/parser.c.jj 2015-07-16 18:09:25.000000000 +0200 +++ gcc/cp/parser.c 2015-07-17 14:04:34.945101113 +0200 @@ -27748,6 +27748,8 @@ cp_parser_omp_clause_name (cp_parser *pa result = PRAGMA_OMP_CLAUSE_LASTPRIVATE; else if (!strcmp ("linear", p)) result = PRAGMA_OMP_CLAUSE_LINEAR; + else if (!strcmp ("link", p)) + result = PRAGMA_OMP_CLAUSE_LINK; break; case 'm': if (!strcmp ("map", p)) @@ -27987,7 +27989,7 @@ cp_parser_omp_var_list_no_open (cp_parse && !TREE_READONLY (low_bound)) { error_at (token->location, - "%qD is not a constant", low_bound); + "%qD is not a constant", low_bound); decl = error_mark_node; } @@ -27995,7 +27997,7 @@ cp_parser_omp_var_list_no_open (cp_parse && !TREE_READONLY (length)) { error_at (token->location, - "%qD is not a constant", length); + "%qD is not a constant", length); decl = error_mark_node; } } @@ -30198,14 +30200,20 @@ cp_parser_omp_all_clauses (cp_parser *pa if (!first) goto clause_not_first; break; + case PRAGMA_OMP_CLAUSE_LINK: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_LINK, clauses); + c_name = "to"; + break; case PRAGMA_OMP_CLAUSE_TO: - clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, - clauses); + if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LINK)) != 0) + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO_DECLARE, + clauses); + else + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses); c_name = "to"; break; case PRAGMA_OMP_CLAUSE_FROM: - clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, - clauses); + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses); c_name = "from"; break; case PRAGMA_OMP_CLAUSE_UNIFORM: @@ -33168,13 +33176,65 @@ cp_parser_late_parsing_omp_declare_simd /* OpenMP 4.0: # pragma omp declare target new-line declarations and definitions - # pragma omp end declare target new-line */ + # pragma omp end declare target new-line + + OpenMP 4.1: + # pragma omp declare target ( extended-list ) new-line + + # pragma omp declare target declare-target-clauses[seq] new-line */ + +#define OMP_DECLARE_TARGET_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_TO) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LINK)) static void cp_parser_omp_declare_target (cp_parser *parser, cp_token *pragma_tok) { - cp_parser_skip_to_pragma_eol (parser, pragma_tok); - scope_chain->omp_declare_target_attribute++; + tree clauses = NULL_TREE; + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + clauses + = cp_parser_omp_all_clauses (parser, OMP_DECLARE_TARGET_CLAUSE_MASK, + "#pragma omp declare target", pragma_tok); + else if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN)) + { + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO_DECLARE, + clauses); + cp_parser_skip_to_pragma_eol (parser, pragma_tok); + } + else + { + cp_parser_skip_to_pragma_eol (parser, pragma_tok); + scope_chain->omp_declare_target_attribute++; + return; + } + if (scope_chain->omp_declare_target_attribute) + error_at (pragma_tok->location, + "%<#pragma omp declare target%> with clauses in between " + "%<#pragma omp declare target%> without clauses and " + "%<#pragma omp end declare target%>"); + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + tree t = OMP_CLAUSE_DECL (c), id; + tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t)); + tree at2 = lookup_attribute ("omp declare target link", + DECL_ATTRIBUTES (t)); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK) + { + id = get_identifier ("omp declare target link"); + std::swap (at1, at2); + } + else + id = get_identifier ("omp declare target"); + if (at2) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD specified both in declare target %<link%> and %<to%>" + " clauses", t); + continue; + } + if (!at1) + DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t)); + } } static void --- gcc/cp/semantics.c.jj 2015-07-16 17:56:41.000000000 +0200 +++ gcc/cp/semantics.c 2015-07-17 13:59:27.177346223 +0200 @@ -6266,6 +6266,36 @@ finish_omp_clauses (tree clauses, bool a bitmap_set_bit (&map_head, DECL_UID (t)); break; + case OMP_CLAUSE_TO_DECLARE: + t = OMP_CLAUSE_DECL (c); + if (TREE_CODE (t) == FUNCTION_DECL) + break; + /* FALLTHRU */ + case OMP_CLAUSE_LINK: + t = OMP_CLAUSE_DECL (c); + if (!VAR_P (t)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE is not a variable in clause %qs", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + else if (DECL_THREAD_LOCAL_P (t)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD is threadprivate variable in %qs clause", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + else if (!cp_omp_mappable_type (TREE_TYPE (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD does not have a mappable type in %qs clause", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + break; + case OMP_CLAUSE_UNIFORM: t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) != PARM_DECL) --- gcc/testsuite/c-c++-common/gomp/declare-target-1.c.jj 2015-07-17 14:07:10.523953776 +0200 +++ gcc/testsuite/c-c++-common/gomp/declare-target-1.c 2015-07-17 14:07:30.472678409 +0200 @@ -0,0 +1,13 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +int foo (void), bar (void); +extern int a; +int b; +char d; +#pragma omp declare target +long c; +#pragma omp end declare target + +#pragma omp declare target (bar, a) +#pragma omp declare target to (b) link (d) to (foo) --- gcc/testsuite/c-c++-common/gomp/declare-target-2.c.jj 2015-07-17 14:23:16.246720738 +0200 +++ gcc/testsuite/c-c++-common/gomp/declare-target-2.c 2015-07-17 14:21:32.000000000 +0200 @@ -0,0 +1,27 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +extern int a; +#pragma omp declare target +#pragma omp declare target to (a) /* { dg-error "with clauses in between" } */ +#pragma omp end declare target +int b; +#pragma omp declare target to (b) link (b) /* { dg-error "specified both in declare target" } */ +int c; +#pragma omp declare target (c) +#pragma omp declare target link (c) /* { dg-error "specified both in declare target" } */ +int foo (void); +#pragma omp declare target link (foo) /* { dg-error "is not a variable in clause" } */ +struct S; +extern struct S d[]; /* { dg-error "array type has incomplete element type" "" { target c } } */ +#pragma omp declare target to (d) /* { dg-error "does not have a mappable type in" } */ +extern struct S e; +#pragma omp declare target link (e) /* { dg-error "does not have a mappable type in" } */ +extern int f[]; +#pragma omp declare target to (f) /* { dg-error "does not have a mappable type in" } */ +int g, h; +#pragma omp threadprivate (g, h) +#pragma omp declare target to (g) /* { dg-error "is threadprivate variable in" } */ +#pragma omp declare target link (h) /* { dg-error "is threadprivate variable in" } */ +int j[10]; +#pragma omp declare target to (j[0:4]) /* { dg-error "expected" } */