| Message ID | 87mvpqrets.fsf@kepler.schwinge.homeip.net |
|---|---|
| State | New |
| Headers | show |
On 03/22/2016 11:23 AM, Thomas Schwinge wrote: > diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c > index 01d1dc8..5806cb3 100644 > --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c > +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c > @@ -1,5 +1,6 @@ > -/* { dg-do run } */ > -/* { dg-additional-options "-O2" } */ > +/* Dead code elimination for blocks guarded by acc_on_device () only works with > + optimizations enabled. > + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ What exactly is going on with these? Do these tests fail with -O0, and is that likely to be a problem in practice? Also, why remove the dg-do run? Bernd
Hi! On Tue, 22 Mar 2016 23:52:11 +0100, Bernd Schmidt <bschmidt@redhat.com> wrote: > On 03/22/2016 11:23 AM, Thomas Schwinge wrote: > > --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c > > +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c > > @@ -1,5 +1,6 @@ > > -/* { dg-do run } */ > > -/* { dg-additional-options "-O2" } */ > > +/* Dead code elimination for blocks guarded by acc_on_device () only works with > > + optimizations enabled. > > + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ > > What exactly is going on with these? Do these tests fail with -O0, and > is that likely to be a problem in practice? Want me to re-word that? :-| I thought it would be obvious from looking at the test case code; will not be a problem in practice. It's because of constructs used in the test cases, like the following, for example: #pragma acc loop worker for (unsigned ix = 0; ix < N; ix++) { if (__builtin_acc_on_device (5)) { int g = 0, w = 0, v = 0; __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); ary[ix] = (g << 16) | (w << 8) | v; } else ary[ix] = ix; } Without optimizations, the target (x86_64) assembler will bail out seeing the device (nvptx) inline assembly code, even if it's dead code always because of the acc_on_device () conditional. Long ago, my suggestion has been to have GCC provide builtin functions for users to retrieve the number of gangs, workers, vectors, and the current thread's IDs of these; not sure why Nathan didn't implement that? (Should be easy to do -- want me to have a look at that, as a separate patch?) > Also, why remove the dg-do run? Because that's the default anyway. Grüße Thomas
On 03/23/2016 07:47 AM, Thomas Schwinge wrote: > Want me to re-word that? :-| I thought it would be obvious from looking > at the test case code; will not be a problem in practice. It's because > of constructs used in the test cases, like the following, for example: > if (__builtin_acc_on_device (5)) > { > int g = 0, w = 0, v = 0; > > __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); > __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); > __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); > ary[ix] = (g << 16) | (w << 8) | v; Ok, maybe write "This code uses nvptx assembly guarded with __builtin_acc_on_device, which fails to be optimized away at -O0." Ok with something like that. Bernd
On Tue, Mar 22, 2016 at 11:23:43AM +0100, Thomas Schwinge wrote: > As discussed in > <http://news.gmane.org/find-root.php?message_id=%3C87oaj3z280.fsf%40schwinge.name%3E> > (and similar to what we're already doing for Fortran, and similar to what > recently got committed to libgomp/testsuite/libgomp.hsa.c/c.exp), it has > been helpful to also run C, C++ offloading test cases with -O0 in > addition to the -O2 default. Making my earlier gomp-4_0-branch patch > conceptually simpler, I came up with the following; OK for trunk? How big difference in make check-target-libgomp time is that? Without PTX offloading I bet zero, but with PTX offloading configured, is it 10% or 50% slower? Jakub
diff --git libgomp/testsuite/libgomp.oacc-c++/c++.exp libgomp/testsuite/libgomp.oacc-c++/c++.exp index 88b0269..bbdbe2f 100644 --- libgomp/testsuite/libgomp.oacc-c++/c++.exp +++ libgomp/testsuite/libgomp.oacc-c++/c++.exp @@ -2,6 +2,7 @@ load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp +load_gcc_lib torture-options.exp global shlib_ext @@ -13,13 +14,9 @@ if [info exists lang_include_flags] then { unset lang_include_flags } -# If a testcase doesn't have special options, use these. -if ![info exists DEFAULT_CFLAGS] then { - set DEFAULT_CFLAGS "-O2" -} - # Initialize dg. dg-init +torture-init # Turn on OpenACC. lappend ALWAYS_CFLAGS "additional_flags=-fopenacc" @@ -104,7 +101,26 @@ if { $lang_test_file_found } { setenv ACC_DEVICE_TYPE $offload_target_openacc - dg-runtest $tests "$tagopt" "$libstdcxx_includes $DEFAULT_CFLAGS" + # To get better test coverage for device-specific code that is only + # ever used in offloading configurations, we'd like more thorough + # testing for test cases that deal with offloading, which most of all + # OpenACC test cases are. We enable torture testing, but limit it to + # -O0 and -O2 only, to avoid testing times exploding too much, under + # the assumption that between -O0 and -O[something] there is the + # biggest difference in the overall structure of the generated code. + switch $offload_target_openacc { + host { + set-torture-options [list \ + { -O2 } ] + } + default { + set-torture-options [list \ + { -O0 } \ + { -O2 } ] + } + } + + gcc-dg-runtest $tests "$tagopt" "$libstdcxx_includes" } } @@ -112,4 +128,5 @@ if { $lang_test_file_found } { set GCC_UNDER_TEST "$SAVE_GCC_UNDER_TEST" # All done. +torture-finish dg-finish diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c index e5d9c36..bfcb67d 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c @@ -1,9 +1,8 @@ -/* { dg-additional-options "-O0" } */ +/* Test the acc_on_device library function. */ +/* { dg-additional-options "-fno-builtin-acc_on_device" } */ #include <openacc.h> -/* acc_on_device might not be folded at -O0, but it should work. */ - int main () { int dev; diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device.c libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device.c index 88c000e..e0d8710 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/acc-on-device.c @@ -1,5 +1,6 @@ /* { dg-do compile } */ -/* { dg-additional-options "-O2" } */ +/* We don't expect this to work with optimizations disabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <openacc.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c index 0f323c8..e8d65df 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c @@ -1,4 +1,4 @@ -/* { dg-additional-options "-O2 -fipa-pta" } */ +/* { dg-additional-options "-fipa-pta" } */ #include <stdlib.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c index 654e750..dd8ca87 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c @@ -1,4 +1,4 @@ -/* { dg-additional-options "-O2 -fipa-pta" } */ +/* { dg-additional-options "-fipa-pta" } */ #include <stdlib.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c index 44d4fd2..50e7dc1 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c @@ -1,4 +1,4 @@ -/* { dg-additional-options "-O2 -fipa-pta" } */ +/* { dg-additional-options "-fipa-pta" } */ #include <stdlib.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c index 4922c9c..e44d26e 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c @@ -1,5 +1,7 @@ /* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> #include <openacc.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c index 36b882f..80e4363 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c @@ -1,5 +1,7 @@ - -/* { dg-additional-options "-O2 -fopenacc-dim=16:16" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ +/* { dg-additional-options "-fopenacc-dim=16:16" } */ #include <openacc.h> #include <alloca.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c index 23c2a75..0c8e9ff 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c index 1a93db3..2316ce3 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c index 26b2df9..e05573f 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c index c14bddd..17a27b1 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c index 3de8b09..4566568 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c index bae5c66..94c752f 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c index ada92e3..3c8b4b8 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c index 706d0d8..7616b19 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c index a073ac8..5f0a439 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c index b9ec95b..cdba658 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c index 539e41d..a65d535 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c index fcb9079..df1a9fc 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c index 201dc72..b30424f 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c index 8fc4cb5..b1cdfd0 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c index fefeb93..c49f643 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c index 01d1dc8..5806cb3 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c index bd9f1df..49761ce 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c @@ -1,5 +1,6 @@ -/* { dg-do run } */ -/* { dg-additional-options "-O2" } */ +/* Dead code elimination for blocks guarded by acc_on_device () only works with + optimizations enabled. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ #include <stdio.h> diff --git libgomp/testsuite/libgomp.oacc-c/c.exp libgomp/testsuite/libgomp.oacc-c/c.exp index 5020e6a..b509a10 100644 --- libgomp/testsuite/libgomp.oacc-c/c.exp +++ libgomp/testsuite/libgomp.oacc-c/c.exp @@ -13,14 +13,11 @@ if [info exists lang_include_flags] then { load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp - -# If a testcase doesn't have special options, use these. -if ![info exists DEFAULT_CFLAGS] then { - set DEFAULT_CFLAGS "-O2" -} +load_gcc_lib torture-options.exp # Initialize dg. dg-init +torture-init # Turn on OpenACC. lappend ALWAYS_CFLAGS "additional_flags=-fopenacc" @@ -67,8 +64,28 @@ foreach offload_target_openacc $offload_targets_s_openacc { setenv ACC_DEVICE_TYPE $offload_target_openacc - dg-runtest $tests "$tagopt" $DEFAULT_CFLAGS + # To get better test coverage for device-specific code that is only + # ever used in offloading configurations, we'd like more thorough + # testing for test cases that deal with offloading, which most of all + # OpenACC test cases are. We enable torture testing, but limit it to + # -O0 and -O2 only, to avoid testing times exploding too much, under + # the assumption that between -O0 and -O[something] there is the + # biggest difference in the overall structure of the generated code. + switch $offload_target_openacc { + host { + set-torture-options [list \ + { -O2 } ] + } + default { + set-torture-options [list \ + { -O0 } \ + { -O2 } ] + } + } + + gcc-dg-runtest $tests "$tagopt" "" } # All done. +torture-finish dg-finish
Hi! As discussed in <http://news.gmane.org/find-root.php?message_id=%3C87oaj3z280.fsf%40schwinge.name%3E> (and similar to what we're already doing for Fortran, and similar to what recently got committed to libgomp/testsuite/libgomp.hsa.c/c.exp), it has been helpful to also run C, C++ offloading test cases with -O0 in addition to the -O2 default. Making my earlier gomp-4_0-branch patch conceptually simpler, I came up with the following; OK for trunk? commit 879c8f6dcb9dad514fb3bf11c721fed37b6574be Author: Thomas Schwinge <thomas@codesourcery.com> Date: Tue Mar 22 10:26:19 2016 +0100 Also test -O0 for OpenACC C, C++ offloading test cases libgomp/ * testsuite/libgomp.oacc-c++/c++.exp: Set up torture testing, use gcc-dg-runtest. * testsuite/libgomp.oacc-c/c.exp: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc-on-device-2.c: Specify -fno-builtin-acc_on_device instead of -O0. * testsuite/libgomp.oacc-c-c++-common/acc-on-device.c: Skip for -O0. * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-v-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-v-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-2.c: Don't specify -O2. * testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-alias-ipa-pta.c: Likewise. --- libgomp/testsuite/libgomp.oacc-c++/c++.exp | 29 +++++++++++++++++----- .../libgomp.oacc-c-c++-common/acc-on-device-2.c | 5 ++-- .../libgomp.oacc-c-c++-common/acc-on-device.c | 3 ++- .../kernels-alias-ipa-pta-2.c | 2 +- .../kernels-alias-ipa-pta-3.c | 2 +- .../kernels-alias-ipa-pta.c | 2 +- .../libgomp.oacc-c-c++-common/loop-auto-1.c | 4 ++- .../libgomp.oacc-c-c++-common/loop-dim-default.c | 6 +++-- .../testsuite/libgomp.oacc-c-c++-common/loop-g-1.c | 5 ++-- .../testsuite/libgomp.oacc-c-c++-common/loop-g-2.c | 5 ++-- .../libgomp.oacc-c-c++-common/loop-gwv-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/loop-red-g-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/loop-red-gwv-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/loop-red-v-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/loop-red-v-2.c | 5 ++-- .../libgomp.oacc-c-c++-common/loop-red-w-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/loop-red-w-2.c | 5 ++-- .../testsuite/libgomp.oacc-c-c++-common/loop-v-1.c | 5 ++-- .../testsuite/libgomp.oacc-c-c++-common/loop-w-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/loop-wv-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/routine-g-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/routine-gwv-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/routine-v-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/routine-w-1.c | 5 ++-- .../libgomp.oacc-c-c++-common/routine-wv-1.c | 5 ++-- libgomp/testsuite/libgomp.oacc-c/c.exp | 29 +++++++++++++++++----- 26 files changed, 111 insertions(+), 56 deletions(-) Grüße Thomas