* [gomp4] Questions about "declare target" and "target update" pragmas @ 2014-01-22 16:02 Ilya Verbin 2014-01-22 18:02 ` Jakub Jelinek 2015-03-10 16:53 ` Ilya Verbin 0 siblings, 2 replies; 13+ messages in thread From: Ilya Verbin @ 2014-01-22 16:02 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc Hi Jakub, I have 2 questions concerning OpenMP 4.0 specification. 1. Do I understand correctly that every "declare target" directive should be closed with "end declare target"? E.g. in this example GCC marks both foo1 and foo2 with "omp declare target" attribute: #pragma omp declare target int foo1 () { return 1; } int foo2 () { return 2; } /* EOF */ Shouldn't the frontend issue an error message that there is "declare target" without corresponding "end declare target"? 2. Do I understand correctly that the "target update" directive can be used outside the "target" or "target data" regions? E.g. this example should print '2' (and it prints '2' while building with icc): #pragma omp declare target int G = 1; #pragma omp end declare target int main () { G = 2; #pragma omp target update to(G) G = 3; int x = 0; #pragma omp target { x = G; } printf ("%d\n", x); } If it is acceptable, then GOMP_target_update should also map variables that wasn't mapped. Thanks, -- Ilya ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2014-01-22 16:02 [gomp4] Questions about "declare target" and "target update" pragmas Ilya Verbin @ 2014-01-22 18:02 ` Jakub Jelinek 2014-01-28 18:14 ` Ilya Verbin 2015-03-10 16:53 ` Ilya Verbin 1 sibling, 1 reply; 13+ messages in thread From: Jakub Jelinek @ 2014-01-22 18:02 UTC (permalink / raw) To: Ilya Verbin; +Cc: gcc On Wed, Jan 22, 2014 at 07:51:51PM +0400, Ilya Verbin wrote: > I have 2 questions concerning OpenMP 4.0 specification. > > > 1. Do I understand correctly that every "declare target" directive should be > closed with "end declare target"? E.g. in this example GCC marks both foo1 and > foo2 with "omp declare target" attribute: > > #pragma omp declare target > int foo1 () { return 1; } > int foo2 () { return 2; } > /* EOF */ Yes, this is invalid. > Shouldn't the frontend issue an error message that there is "declare target" > without corresponding "end declare target"? I guess so, e.g. the C FE has current_omp_declare_target_attribute, supposedly it could check it after parsing the whole file and if it is non-zero at the end, complain. > 2. Do I understand correctly that the "target update" directive can be used > outside the "target" or "target data" regions? E.g. this example should > print '2' (and it prints '2' while building with icc): > > #pragma omp declare target > int G = 1; > #pragma omp end declare target > > int main () > { > G = 2; > #pragma omp target update to(G) > G = 3; > int x = 0; > #pragma omp target > { > x = G; > } > printf ("%d\n", x); > } This can print 3 (if doing host fallback or device shares address space with host), or 2 (otherwise). It shouldn't print 1 ever, and yes, the target update is then well defined. All variables from omp declare target are allocated on the device sometime before the first target data/target update/target region; given that they will be allocated in the data section of the target DSO, they actually just need to be registered with the mapping data structure when the DSO is loaded. > If it is acceptable, then GOMP_target_update should also map variables that > wasn't mapped. No, the target DSO initialization should use the tables we've talked about to initialize the mapping. Jakub ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2014-01-22 18:02 ` Jakub Jelinek @ 2014-01-28 18:14 ` Ilya Verbin 2014-01-28 21:39 ` Jakub Jelinek 0 siblings, 1 reply; 13+ messages in thread From: Ilya Verbin @ 2014-01-28 18:14 UTC (permalink / raw) To: Jakub Jelinek; +Cc: GCC Development 2014/1/22 Jakub Jelinek <jakub@redhat.com>: > This can print 3 (if doing host fallback or device shares address space > with host), or 2 (otherwise). It shouldn't print 1 ever, and yes, > the target update is then well defined. All variables from omp declare > target are allocated on the device sometime before > the first target data/target update/target region; given that they will > be allocated in the data section of the target DSO, they actually just need > to be registered with the mapping data structure when the DSO is loaded. > > No, the target DSO initialization should use the tables we've talked about > to initialize the mapping. > > Jakub Yes, when G is global variable marked with 'declare target', everything works fine. But this testcase crashes at runtime in GOMP_target_update: int main () { int G = 2; #pragma omp target update to(G) G = 3; int x = 0; #pragma omp target { x = G; } printf ("%d\n", x); } Is it right, that such usage of 'target update' is not allowed by omp specification? -- Ilya ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2014-01-28 18:14 ` Ilya Verbin @ 2014-01-28 21:39 ` Jakub Jelinek [not found] ` <CADG=Z0GzEZw8a8S_M+3SzEXKn++qf5UDRjcnWD5FdPNjN7u0Mw@mail.gmail.com> 2014-07-08 18:31 ` Ilya Verbin 0 siblings, 2 replies; 13+ messages in thread From: Jakub Jelinek @ 2014-01-28 21:39 UTC (permalink / raw) To: Ilya Verbin; +Cc: GCC Development On Tue, Jan 28, 2014 at 09:54:09PM +0400, Ilya Verbin wrote: > Yes, when G is global variable marked with 'declare target', everything works > fine. But this testcase crashes at runtime in GOMP_target_update: > > int main () > { > int G = 2; > #pragma omp target update to(G) > G = 3; > int x = 0; > #pragma omp target > { > x = G; > } > printf ("%d\n", x); > } > > Is it right, that such usage of 'target update' is not allowed by omp > specification? Yes, this testcase is invalid. "If the corresponding list item is not present in the device data environment, the behavior is unspecified." Perhaps we shouldn't crash but do nothing, or complain to stderr, as QoI. Jakub ^ permalink raw reply [flat|nested] 13+ messages in thread
[parent not found: <CADG=Z0GzEZw8a8S_M+3SzEXKn++qf5UDRjcnWD5FdPNjN7u0Mw@mail.gmail.com>]
* Fwd: [gomp4] Questions about "declare target" and "target update" pragmas [not found] ` <CADG=Z0GzEZw8a8S_M+3SzEXKn++qf5UDRjcnWD5FdPNjN7u0Mw@mail.gmail.com> @ 2014-01-30 19:17 ` Ilya Verbin 0 siblings, 0 replies; 13+ messages in thread From: Ilya Verbin @ 2014-01-30 19:17 UTC (permalink / raw) To: Jakub Jelinek; +Cc: GCC Development One more question. Is it valid to use arr[MAX/2..MAX] on target? #define MAX 20 void foo () { int arr[MAX]; #pragma omp target map(from: arr[0:MAX/2]) { int i; for (i = 0; i < MAX; i++) arr[i] = i; } } In this case GOMP_target gets sizes[0]==40 as input. Due to this, gomp_map_vars allocates 40 bytes of memory on target for 'arr', instead of 80 bytes. -- Ilya ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2014-01-28 21:39 ` Jakub Jelinek [not found] ` <CADG=Z0GzEZw8a8S_M+3SzEXKn++qf5UDRjcnWD5FdPNjN7u0Mw@mail.gmail.com> @ 2014-07-08 18:31 ` Ilya Verbin 1 sibling, 0 replies; 13+ messages in thread From: Ilya Verbin @ 2014-07-08 18:31 UTC (permalink / raw) To: Jakub Jelinek; +Cc: GCC Development Hi Jakub, I discovered an issue related to global variables. In this testcase the 'omp target' child fn uses the local copy of glob_var. But the 'omp parallel' child fn tries to use the glob_var directly and therefore crashes. int glob_var; void foo (void) { glob_var = 1; #pragma omp target map(to: glob_var) { glob_var = 2; #pragma omp parallel { glob_var = 3; } } } In the spec I found only that if a variable is not present in the enclosing device data environment, then it is allocated in the device data environment associated with the construct. Effectively this means glob_var become a non-global variable within the omp_target construct. Then it's not quite clear what type of glob_var should be in case target is not present and we fallback to host execution. Should we forbid 'omp target' to map global vars that are not defined as target? Or force 'omp parallel' to use local copies inside the target regions? (ICC reports an error about mapping glob_var for this testcase) -- Ilya ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2014-01-22 16:02 [gomp4] Questions about "declare target" and "target update" pragmas Ilya Verbin 2014-01-22 18:02 ` Jakub Jelinek @ 2015-03-10 16:53 ` Ilya Verbin 2015-03-16 18:42 ` Ilya Verbin 1 sibling, 1 reply; 13+ messages in thread From: Ilya Verbin @ 2015-03-10 16:53 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc, Kirill Yukhin Hi Jakub, I have one more question :) This testcase seems to be correct... or not? #pragma omp declare target extern int G; #pragma omp end declare target int G; int main () { #pragma omp target update to(G) return 0; } If yes, then we have a problem that the decl of G in varpool_node::get_create doesn't have "omp declare target" attribute. Thanks, -- Ilya ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2015-03-10 16:53 ` Ilya Verbin @ 2015-03-16 18:42 ` Ilya Verbin 2015-03-19 13:47 ` Jakub Jelinek 0 siblings, 1 reply; 13+ messages in thread From: Ilya Verbin @ 2015-03-16 18:42 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc, Kirill Yukhin On Tue, Mar 10, 2015 at 19:52:52 +0300, Ilya Verbin wrote: > Hi Jakub, > > I have one more question :) > This testcase seems to be correct... or not? > > #pragma omp declare target > extern int G; > #pragma omp end declare target > > int G; > > int main () > { > #pragma omp target update to(G) > > return 0; > } > > If yes, then we have a problem that the decl of G in varpool_node::get_create > doesn't have "omp declare target" attribute. Ping? I am investigating run-fails on some benchmark, and have found a second questionable place, where a function argument overrides a global array. Just to be sure, is this a bug in the test? #pragma omp declare target int a1[50], a2[50]; #pragma omp end declare target void foo (int a1[]) { #pragma omp target { a1[10]++; a2[10]++; } } int main () { a1[10] = a2[10] = 10; #pragma omp target update to(a1, a2) foo (a1); #pragma omp target update from(a1, a2) if (a1[10] != a2[10]) abort (); return 0; } Thanks, -- Ilya ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2015-03-16 18:42 ` Ilya Verbin @ 2015-03-19 13:47 ` Jakub Jelinek 2015-03-19 14:50 ` Ilya Verbin 0 siblings, 1 reply; 13+ messages in thread From: Jakub Jelinek @ 2015-03-19 13:47 UTC (permalink / raw) To: Ilya Verbin; +Cc: gcc, Kirill Yukhin On Mon, Mar 16, 2015 at 09:41:53PM +0300, Ilya Verbin wrote: > On Tue, Mar 10, 2015 at 19:52:52 +0300, Ilya Verbin wrote: > > Hi Jakub, > > > > I have one more question :) > > This testcase seems to be correct... or not? > > > > #pragma omp declare target > > extern int G; > > #pragma omp end declare target > > > > int G; > > > > int main () > > { > > #pragma omp target update to(G) > > > > return 0; > > } > > > > If yes, then we have a problem that the decl of G in varpool_node::get_create > > doesn't have "omp declare target" attribute. > > Ping? Here is untested patch. I'm going to check it in after bootstrap/regtest. > I am investigating run-fails on some benchmark, and have found a second > questionable place, where a function argument overrides a global array. > Just to be sure, is this a bug in the test? > > #pragma omp declare target > int a1[50], a2[50]; > #pragma omp end declare target > > void foo (int a1[]) > { > #pragma omp target > { > a1[10]++; > a2[10]++; > } > } That is a buggy test. int a1[] function argument is changed into int *a1, so it is actually #pragma omp target map(tofrom:a1, a2) { a1[10]++; a2[10]++; } which copies the a1 pointer to the device by value (no pointer transformation). Perhaps the testcase writer meant to use #pragma omp target map(a1[10]) instead (or map(a1[0:50])? 2015-03-19 Jakub Jelinek <jakub@redhat.com> * c-decl.c (c_decl_attributes): Also add "omp declare target" attribute for DECL_EXTERNAL VAR_DECLs. * decl2.c (cplus_decl_attributes): Also add "omp declare target" attribute for DECL_EXTERNAL VAR_DECLs. * testsuite/libgomp.c/target-10.c: New test. * testsuite/libgomp.c++/target-4.C: New test. --- gcc/c/c-decl.c.jj 2015-03-09 19:24:34.000000000 +0100 +++ gcc/c/c-decl.c 2015-03-19 13:01:15.423381262 +0100 @@ -4407,7 +4407,8 @@ c_decl_attributes (tree *node, tree attr { /* Add implicit "omp declare target" attribute if requested. */ if (current_omp_declare_target_attribute - && ((TREE_CODE (*node) == VAR_DECL && TREE_STATIC (*node)) + && ((TREE_CODE (*node) == VAR_DECL + && (TREE_STATIC (*node) || DECL_EXTERNAL (*node))) || TREE_CODE (*node) == FUNCTION_DECL)) { if (TREE_CODE (*node) == VAR_DECL --- gcc/cp/decl2.c.jj 2015-03-18 11:53:13.000000000 +0100 +++ gcc/cp/decl2.c 2015-03-19 13:04:30.739176009 +0100 @@ -1440,7 +1440,8 @@ cplus_decl_attributes (tree *decl, tree /* Add implicit "omp declare target" attribute if requested. */ if (scope_chain->omp_declare_target_attribute - && ((TREE_CODE (*decl) == VAR_DECL && TREE_STATIC (*decl)) + && ((TREE_CODE (*decl) == VAR_DECL + && (TREE_STATIC (*decl) || DECL_EXTERNAL (*decl))) || TREE_CODE (*decl) == FUNCTION_DECL)) { if (TREE_CODE (*decl) == VAR_DECL --- libgomp/testsuite/libgomp.c/target-10.c.jj 2015-03-19 13:06:56.812778618 +0100 +++ libgomp/testsuite/libgomp.c/target-10.c 2015-03-19 13:07:03.857662996 +0100 @@ -0,0 +1,14 @@ +/* { dg-do run } */ + +#pragma omp declare target +extern int v; +#pragma omp end declare target + +int v; + +int +main () +{ + #pragma omp target update to(v) + return 0; +} --- libgomp/testsuite/libgomp.c++/target-4.C.jj 2015-03-19 13:07:19.286409775 +0100 +++ libgomp/testsuite/libgomp.c++/target-4.C 2015-03-19 13:07:30.322228651 +0100 @@ -0,0 +1,3 @@ +// { dg-do run } + +#include "../libgomp.c/target-10.c" Jakub ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2015-03-19 13:47 ` Jakub Jelinek @ 2015-03-19 14:50 ` Ilya Verbin 2015-03-19 14:57 ` Jakub Jelinek 0 siblings, 1 reply; 13+ messages in thread From: Ilya Verbin @ 2015-03-19 14:50 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc, Kirill Yukhin On Thu, Mar 19, 2015 at 14:47:44 +0100, Jakub Jelinek wrote: > Here is untested patch. I'm going to check it in after bootstrap/regtest. Thanks. > > I am investigating run-fails on some benchmark, and have found a second > > questionable place, where a function argument overrides a global array. > > Just to be sure, is this a bug in the test? > > > > #pragma omp declare target > > int a1[50], a2[50]; > > #pragma omp end declare target > > > > void foo (int a1[]) > > { > > #pragma omp target > > { > > a1[10]++; > > a2[10]++; > > } > > } > > That is a buggy test. int a1[] function argument is changed > into int *a1, so it is actually > #pragma omp target map(tofrom:a1, a2) Actually, it copies only a1 pointer, since a2 points to the global array. > { > a1[10]++; > a2[10]++; > } > which copies the a1 pointer to the device by value (no pointer > transformation). > Perhaps the testcase writer meant to use #pragma omp target map(a1[10]) > instead (or map(a1[0:50])? If I understand correctly, it's not allowed to map global target arrays this way, since it's already present in the initial device data environment: 2.9.4 declare target Directive If a list item is a variable then the original variable is mapped to a corresponding variable in the initial device data environment for all devices. 2.14.5 map Clause If a corresponding list item of the original list item is in the enclosing device data environment, the new device data environment uses the corresponding list item from the enclosing device data environment. No additional storage is allocated in the new device data environment and neither initialization nor assignment is performed, regardless of the map-type that is specified. So, to fix this testcase I can just remove the "int a1[]" function argument, and add some "#pragma omp target update" where needed. -- Ilya ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2015-03-19 14:50 ` Ilya Verbin @ 2015-03-19 14:57 ` Jakub Jelinek 2015-03-19 18:43 ` Ilya Verbin 0 siblings, 1 reply; 13+ messages in thread From: Jakub Jelinek @ 2015-03-19 14:57 UTC (permalink / raw) To: Ilya Verbin; +Cc: gcc, Kirill Yukhin On Thu, Mar 19, 2015 at 05:49:47PM +0300, Ilya Verbin wrote: > > > void foo (int a1[]) > > > { > > > #pragma omp target > > > { > > > a1[10]++; > > > a2[10]++; > > > } > > > } > > > > That is a buggy test. int a1[] function argument is changed > > into int *a1, so it is actually > > #pragma omp target map(tofrom:a1, a2) > > Actually, it copies only a1 pointer, since a2 points to the global array. Sure, the implicit map(tofrom:a2) doesn't really do anything, since a2 is clearly already mapped and will stay to be mapped, and is a global var. > > > { > > a1[10]++; > > a2[10]++; > > } > > which copies the a1 pointer to the device by value (no pointer > > transformation). > > Perhaps the testcase writer meant to use #pragma omp target map(a1[10]) > > instead (or map(a1[0:50])? > > If I understand correctly, it's not allowed to map global target arrays this > way, since it's already present in the initial device data environment: It of course is allowed. It just means that it doesn't allocate new memory (sizeof(int) large in the map(a1[10]) case, sizeof(int)*50 large in the map(a1[0:50]) case), nor copy the bytes around, all it will do is allocate memory for the target copy of the a1 pointer, and do pointer transformation such that the a1 pointer on the target will point to the global target a1 array. Without the map(a1[10]) or map(a1[0:50]) clauses (i.e. implicit map(tofrom:a1)) it does similar thing, except it copies the pointer value to the target (and back at the end of the region) instead, which is not what you want... > So, to fix this testcase I can just remove the "int a1[]" function argument, and > add some "#pragma omp target update" where needed. Well, supposedly the test was meant to test it with a pointer parameter, otherwise why would there be both a1 and a2 when a2 would be enough? Jakub ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2015-03-19 14:57 ` Jakub Jelinek @ 2015-03-19 18:43 ` Ilya Verbin 2015-03-19 19:00 ` Jakub Jelinek 0 siblings, 1 reply; 13+ messages in thread From: Ilya Verbin @ 2015-03-19 18:43 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc, Kirill Yukhin On Thu, Mar 19, 2015 at 15:57:10 +0100, Jakub Jelinek wrote: > On Thu, Mar 19, 2015 at 05:49:47PM +0300, Ilya Verbin wrote: > > If I understand correctly, it's not allowed to map global target arrays this > > way, since it's already present in the initial device data environment: > > It of course is allowed. It just means that it doesn't allocate new memory > (sizeof(int) large in the map(a1[10]) case, sizeof(int)*50 large in the map(a1[0:50]) > case), nor copy the bytes around, all it will do is allocate memory for the > target copy of the a1 pointer, and do pointer transformation such that the > a1 pointer on the target will point to the global target a1 array. > Without the map(a1[10]) or map(a1[0:50]) clauses (i.e. implicit map(tofrom:a1)) > it does similar thing, except it copies the pointer value to the target (and > back at the end of the region) instead, which is not what you want... Ok, got it. And what about global allocatable fortran arrays? I didn't find any restrictions in the specification. Here is a reduced testcase: module test integer, allocatable, target :: x(:) !$omp declare target(x) end module test use test integer :: n = 1000 allocate (x(n)) !$omp target map(x(1:n)) x(123) = 456 !$omp end target deallocate (x) end It crashes on target with NULL-pointer access, however the memory for x(1:n) is allocated on target. Looks like there's something wrong with pointer transformation. Is this a wrong testcase or a bug in gcc? Thanks, -- Ilya ^ permalink raw reply [flat|nested] 13+ messages in thread
* Re: [gomp4] Questions about "declare target" and "target update" pragmas 2015-03-19 18:43 ` Ilya Verbin @ 2015-03-19 19:00 ` Jakub Jelinek 0 siblings, 0 replies; 13+ messages in thread From: Jakub Jelinek @ 2015-03-19 19:00 UTC (permalink / raw) To: Ilya Verbin; +Cc: gcc, Kirill Yukhin On Thu, Mar 19, 2015 at 09:42:58PM +0300, Ilya Verbin wrote: > Ok, got it. > > And what about global allocatable fortran arrays? I didn't find any > restrictions in the specification. Here is a reduced testcase: This really can't be supported. If you have global allocatables, you really should allocate them in the target IMNSHO. The allocation on the host doesn't make it allocated on the target. Otherwise, how you could e.g. deallocate it on the target or allocate again? > > module test > integer, allocatable, target :: x(:) > !$omp declare target(x) > end module test > use test > integer :: n = 1000 > allocate (x(n)) > !$omp target map(x(1:n)) > x(123) = 456 > !$omp end target > deallocate (x) > end > > It crashes on target with NULL-pointer access, however the memory for x(1:n) is > allocated on target. Looks like there's something wrong with pointer > transformation. Is this a wrong testcase or a bug in gcc? Jakub ^ permalink raw reply [flat|nested] 13+ messages in thread
end of thread, other threads:[~2015-03-19 19:00 UTC | newest] Thread overview: 13+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2014-01-22 16:02 [gomp4] Questions about "declare target" and "target update" pragmas Ilya Verbin 2014-01-22 18:02 ` Jakub Jelinek 2014-01-28 18:14 ` Ilya Verbin 2014-01-28 21:39 ` Jakub Jelinek [not found] ` <CADG=Z0GzEZw8a8S_M+3SzEXKn++qf5UDRjcnWD5FdPNjN7u0Mw@mail.gmail.com> 2014-01-30 19:17 ` Fwd: " Ilya Verbin 2014-07-08 18:31 ` Ilya Verbin 2015-03-10 16:53 ` Ilya Verbin 2015-03-16 18:42 ` Ilya Verbin 2015-03-19 13:47 ` Jakub Jelinek 2015-03-19 14:50 ` Ilya Verbin 2015-03-19 14:57 ` Jakub Jelinek 2015-03-19 18:43 ` Ilya Verbin 2015-03-19 19:00 ` Jakub Jelinek
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).