| Message ID | 20141209123233.GA32616@msticlxl57.ims.intel.com |
|---|---|
| State | New |
| Headers | show |
On Tue, Dec 09, 2014 at 03:32:33PM +0300, Ilya Verbin wrote: > On 04 Dec 20:52, Jakub Jelinek wrote: > > On Thu, Dec 04, 2014 at 10:35:19PM +0300, Ilya Verbin wrote: > > > This issue can be resolved by forcing output of such variables. > > > Is this fix ok? Should I add a testcase? > > > > Yes, with proper ChangeLog. Yes. > > Here is updated patch, ok to commit? > > However, I don't see -flto option in the build log. It seems that > check_effective_target_lto isn't working inside libgomp/ directory. > Maybe because ENABLE_LTO is defined only in gcc/configure.ac ? > > > gcc/ > * varpool.c (varpool_node::get_create): Force output of vars with > "omp declare target" attribute. > libgomp/ > * testsuite/libgomp.c/target-9.c: New test. Ok, though please try to find out why effective target lto check doesn't work in libgomp. Perhaps you just need to include some further *.exp file? Jakub
On 10 Dec 09:22, Jakub Jelinek wrote: > On Tue, Dec 09, 2014 at 03:32:33PM +0300, Ilya Verbin wrote: > > However, I don't see -flto option in the build log. It seems that > > check_effective_target_lto isn't working inside libgomp/ directory. > > Maybe because ENABLE_LTO is defined only in gcc/configure.ac ? > > > > gcc/ > > * varpool.c (varpool_node::get_create): Force output of vars with > > "omp declare target" attribute. > > libgomp/ > > * testsuite/libgomp.c/target-9.c: New test. > > Ok, though please try to find out why effective target lto check doesn't > work in libgomp. Perhaps you just need to include some further *.exp > file? It lives in gcc/testsuite/lib/target-supports.exp, which is already included into libgomp/testsuite/lib/libgomp.exp proc check_effective_target_lto { } { global ENABLE_LTO if { [istarget nvptx-*-*] } { return 0; } return [info exists ENABLE_LTO] } I'm not sure how it works, but ENABLE_LTO is defined only in gcc/configure.ac . Maybe it's possible to move it to top-level configure, or to check for "-flto" support instead. However, I will be able to fix this only in late Dec, I'm going on vacation without access to the computer :) -- Ilya
diff --git a/gcc/varpool.c b/gcc/varpool.c index 0526b7f..db28c2a 100644 --- a/gcc/varpool.c +++ b/gcc/varpool.c @@ -175,6 +175,7 @@ varpool_node::get_create (tree decl) g->have_offload = true; if (!in_lto_p) vec_safe_push (offload_vars, decl); + node->force_output = 1; #endif } diff --git a/libgomp/testsuite/libgomp.c/target-9.c b/libgomp/testsuite/libgomp.c/target-9.c new file mode 100644 index 0000000..00fe0cb --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-9.c @@ -0,0 +1,37 @@ +/* { dg-do run } */ +/* { dg-options "-O1" } */ +/* { dg-additional-options "-flto" { target lto } } */ + +#include <stdlib.h> + +#define N 123456 + +#pragma omp declare target +int X, Y; +#pragma omp end declare target + +void +foo () +{ + #pragma omp target map(alloc: X) + X = N; +} + +int +main () +{ + int res; + + foo (); + + #pragma omp target map(alloc: X, Y) map(from: res) + { + Y = N; + res = X + Y; + } + + if (res != N + N) + abort (); + + return 0; +}