| Message ID | 564642E0.6090303@acm.org |
|---|---|
| State | New |
| Headers | show |
On 11/13/2015 09:06 PM, Nathan Sidwell wrote: > On 11/11/15 09:19, Bernd Schmidt wrote: >> On 11/11/2015 02:59 PM, Nathan Sidwell wrote: >>> That's not the problem. How to conditionally enable the test is the >>> difficulty. I suspect porting something concerning accel_compiler from >>> the libgomp testsuite is needed? >> >> Maybe a check_effective_target_offload_nvptx which tries to see if >> -foffload=nvptx gives an error (I would hope it does if it's >> unsupported). > > This patch seems to do the trick. tested on an offload-aware build > (passes), and a regular x86_64-linux build (skipped as unsupported). I think this is fine. Bernd
Hi! On Fri, 13 Nov 2015 21:22:11 +0100, Bernd Schmidt <bschmidt@redhat.com> wrote: > On 11/13/2015 09:06 PM, Nathan Sidwell wrote: > > On 11/11/15 09:19, Bernd Schmidt wrote: > >> On 11/11/2015 02:59 PM, Nathan Sidwell wrote: > >>> That's not the problem. How to conditionally enable the test is the > >>> difficulty. I suspect porting something concerning accel_compiler from > >>> the libgomp testsuite is needed? > >> > >> Maybe a check_effective_target_offload_nvptx which tries to see if > >> -foffload=nvptx gives an error (I would hope it does if it's > >> unsupported). > > > > This patch seems to do the trick. tested on an offload-aware build > > (passes), and a regular x86_64-linux build (skipped as unsupported). Thanks for adding such tests! > I think this is fine. I may have pointed this out before a few times... ;-) -- <http://news.gmane.org/find-root.php?message_id=%3C8737wcv4lg.fsf%40kepler.schwinge.homeip.net%3E> -- this doesn't work for build-tree testing ("make check"), if GCC has not yet been installed ("make install"), which is still the default testing procedure as far as I know. spawn [...]/build-gcc/gcc/xgcc -B[...]/build-gcc/gcc/ [...]/source-gcc/gcc/testsuite/gcc.dg/goacc/nvptx-merged-loop.c -fno-diagnostics-show-caret -fdiagnostics-color=never -fopenacc -O2 -foffload=-fdump-rtl-mach -dumpbase nvptx-merged-loop.c -Wa,--no-verify -ffat-lto-objects -lm -o nvptx-merged-loop.exe xgcc: error: libgomp.spec: No such file or directory compiler exited with status 1 [...] FAIL: gcc.dg/goacc/nvptx-merged-loop.c (test for excess errors) [...] UNRESOLVED: gcc.dg/goacc/nvptx-merged-loop.c scan-rtl-dump mach "Merging loop .* into " Here, -fopenacc induces -lgomp. So, we'll either need a (dummy?) libgomp available to link against in gcc/testsuite/, or come up with a way to do LTO/offloading compilation without actually linking (libgomp into) the final executable, or move such tests into libgomp/testsuite/. (Jakub?) Grüße Thomas
On Mon, Nov 23, 2015 at 08:46:30AM +0100, Thomas Schwinge wrote: > Here, -fopenacc induces -lgomp. So, we'll either need a (dummy?) libgomp > available to link against in gcc/testsuite/, or come up with a way to do > LTO/offloading compilation without actually linking (libgomp into) the > final executable, or move such tests into libgomp/testsuite/. (Jakub?) Link/run tests that link against libgomp belong to libgomp/testsuite/. Jakub
2015-11-13 Nathan Sidwell <nathan@codesourcery.com> * lib/target-supports.exp (check_effective_target_offload_nvptx): New. * gcc.dg/goacc/nvptx-merged-loop.c: New. Index: testsuite/gcc.dg/goacc/nvptx-merged-loop.c =================================================================== --- testsuite/gcc.dg/goacc/nvptx-merged-loop.c (revision 0) +++ testsuite/gcc.dg/goacc/nvptx-merged-loop.c (working copy) @@ -0,0 +1,30 @@ +/* { dg-do link } */ +/* { dg-require-effective-target offload_nvptx } */ +/* { dg-options "-fopenacc -O2 -foffload=-fdump-rtl-mach\\ -dumpbase\\ nvptx-merged-loop.c\\ -Wa,--no-verify" } */ + +#define N (32*32*32+17) +void __attribute__ ((noinline)) Foo (int *ary) +{ + int ix; + +#pragma acc parallel num_workers(32) vector_length(32) copyout(ary[0:N]) + { + /* Loop partitioning should be merged. */ +#pragma acc loop worker vector + for (unsigned ix = 0; ix < N; ix++) + { + ary[ix] = ix; + } + } +} + +int main () +{ + int ary[N]; + + Foo (ary); + + return 0; +} + +/* { dg-final { scan-rtl-dump "Merging loop .* into " "mach" } } */ Index: testsuite/lib/target-supports.exp =================================================================== --- testsuite/lib/target-supports.exp (revision 230324) +++ testsuite/lib/target-supports.exp (working copy) @@ -6716,3 +6716,11 @@ proc check_effective_target_vect_max_red } return 0 } + +# Return 1 if there is an nvptx offload compiler. + +proc check_effective_target_offload_nvptx { } { + return [check_no_compiler_messages offload_nvptx object { + int main () {return 0;} + } "-foffload=nvptx-none" ] +}