This is the mail archive of the
gcc@gcc.gnu.org
mailing list for the GCC project.
Re: [gomp4] Questions about "declare target" and "target update" pragmas
- From: Jakub Jelinek <jakub at redhat dot com>
- To: Ilya Verbin <iverbin at gmail dot com>
- Cc: gcc at gcc dot gnu dot org, Kirill Yukhin <kirill dot yukhin at gmail dot com>
- Date: Thu, 19 Mar 2015 14:47:44 +0100
- Subject: Re: [gomp4] Questions about "declare target" and "target update" pragmas
- Authentication-results: sourceware.org; auth=none
- References: <20140122155151 dot GA50489 at msticlxl57 dot ims dot intel dot com> <20150310165252 dot GC37666 at msticlxl57 dot ims dot intel dot com> <20150316184153 dot GA42550 at msticlxl57 dot ims dot intel dot com>
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
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