| Message ID | 763f86ad-c8b0-de34-e752-f591a87ea851@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | [libgomp] Add tests for print from offload target | expand |
On Thu, Nov 14, 2019 at 04:47:49PM +0000, Andrew Stubbs wrote: > This patch adds new libgomp tests to ensure that C "printf" and Fortran > "write" work correctly within offload kernels. Both should work for amdgcn, > but nvptx uses the libgfortran "minimal" mode which lacks "write" support. So, do those *.f90 testcases now FAIL with nvptx offloading? If yes, perhaps there should be effective target check that it is not nvptx offloading. Once the declare variant support is finished, at least in OpenMP it could be handled through that, but Fortran support for that will definitely not make GCC 10. Jakub
On 11/14/19 5:47 PM, Andrew Stubbs wrote: > This patch adds new libgomp tests to ensure that C "printf" and > Fortran "write" work correctly within offload kernels. Both should > work for amdgcn, but nvptx uses the libgfortran "minimal" mode which > lacks "write" support. Can't you add something like: ! { dg-do run { target { ! { openacc_nvidia_accel_selected } } } } ! For openacc_nvidia_accel_selected, there is no I/O support. To avoid FAILs? Cheers, Tobias
On 14/11/2019 17:05, Jakub Jelinek wrote: > On Thu, Nov 14, 2019 at 04:47:49PM +0000, Andrew Stubbs wrote: >> This patch adds new libgomp tests to ensure that C "printf" and Fortran >> "write" work correctly within offload kernels. Both should work for amdgcn, >> but nvptx uses the libgfortran "minimal" mode which lacks "write" support. > > So, do those *.f90 testcases now FAIL with nvptx offloading? > If yes, perhaps there should be effective target check that it is not nvptx > offloading. > Once the declare variant support is finished, at least in OpenMP it could be > handled through that, but Fortran support for that will definitely not make > GCC 10. Oops, I forgot to regenerate the patch before posting it. Here's the version with the nvptx xfails. OK? Andrew
On Thu, Nov 14, 2019 at 05:18:41PM +0000, Andrew Stubbs wrote: > On 14/11/2019 17:05, Jakub Jelinek wrote: > > On Thu, Nov 14, 2019 at 04:47:49PM +0000, Andrew Stubbs wrote: > > > This patch adds new libgomp tests to ensure that C "printf" and Fortran > > > "write" work correctly within offload kernels. Both should work for amdgcn, > > > but nvptx uses the libgfortran "minimal" mode which lacks "write" support. > > > > So, do those *.f90 testcases now FAIL with nvptx offloading? > > If yes, perhaps there should be effective target check that it is not nvptx > > offloading. > > Once the declare variant support is finished, at least in OpenMP it could be > > handled through that, but Fortran support for that will definitely not make > > GCC 10. > > Oops, I forgot to regenerate the patch before posting it. > > Here's the version with the nvptx xfails. > > OK? Ok. > Add tests for print from offload target. > > 2019-11-14 Andrew Stubbs <ams@codesourcery.com> > > libgomp/ > * testsuite/libgomp.c/target-print-1.c: New file. > * testsuite/libgomp.fortran/target-print-1.f90: New file. > * testsuite/libgomp.oacc-c/print-1.c: New file. > * testsuite/libgomp.oacc-fortran/print-1.f90: New file. > +int > +main () > +{ > +#pragma omp target > + { > + printf ("The answer is %d\n", var); > + } Just a nit, #pragma omp target printf ("The answer is %d\n", var); would be valid too, but no need to change the testcase. Jakub
Hi! On 2019-11-14T18:22:39+0100, Jakub Jelinek <jakub@redhat.com> wrote: > On Thu, Nov 14, 2019 at 05:18:41PM +0000, Andrew Stubbs wrote: >> On 14/11/2019 17:05, Jakub Jelinek wrote: >> > On Thu, Nov 14, 2019 at 04:47:49PM +0000, Andrew Stubbs wrote: >> > > This patch adds new libgomp tests to ensure that C "printf" and Fortran >> > > "write" work correctly within offload kernels. Thanks. >> > > Both should work for amdgcn, >> > > but nvptx uses the libgfortran "minimal" mode which lacks "write" support. >> > >> > So, do those *.f90 testcases now FAIL with nvptx offloading? Note that for libgomp OpenMP testing that means FAIL as soon as nvptx plus possibly any other offload targets are configured, as we will always compile for all offload targets, as no specific '-foffload' gets passed -- in contrast to the libgomp OpenACC testing. (Changing that would be a separate discussion.) >> > If yes, perhaps there should be effective target check that it is not nvptx >> > offloading. >> Here's the version with the nvptx xfails. >> >> OK? > > Ok. ... but that doesn't work as expected: {+ERROR: libgomp.fortran/target-print-1.f90 -O0 : can't read "openacc_device_type": no such variable for " dg-xfail-if 5 "no write for nvidia" { openacc_nvidia_accel_selected } "+} [Etc.] ..., and: {+XFAIL: libgomp.oacc-fortran/print-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 (test for excess errors)+} {+UNRESOLVED: libgomp.oacc-fortran/print-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 compilation failed to produce executable+} [Etc.] To fix the ERRORs, I'm adding 'check_effective_target_offload_target_nvptx' to generally test whether "compiling for offload target nvptx", and to fix the UNRESOLVEDs, I'm adding separate "nvptx" variants, as I couldn't find a way to express the 'dg-*' directives in one combined file. I want to actually see the nvptx XFAILs, so that we'll notice XPASSes once/if nvptx switches away from the minimal libgfortran. (Changing that would be a separate discussion.) Committed "Fix 'libgomp.fortran/target-print-1.f90', 'libgomp.oacc-fortran/print-1.f90' for offload target nvptx" to trunk in r278779, see attached. Grüße Thomas
Add tests for print from offload target. 2019-11-14 Andrew Stubbs <ams@codesourcery.com> libgomp/ * testsuite/libgomp.c/target-print-1.c: New file. * testsuite/libgomp.fortran/target-print-1.f90: New file. * testsuite/libgomp.oacc-c/print-1.c: New file. * testsuite/libgomp.oacc-fortran/print-1.f90: New file. diff --git a/libgomp/testsuite/libgomp.c/target-print-1.c b/libgomp/testsuite/libgomp.c/target-print-1.c new file mode 100644 index 00000000000..5857b875ced --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-print-1.c @@ -0,0 +1,17 @@ +/* Ensure that printf on the offload device works. */ + +/* { dg-do run } */ +/* { dg-output "The answer is 42(\n|\r\n|\r)+" } */ + +#include <stdio.h> + +int var = 42; + +int +main () +{ +#pragma omp target + { + printf ("The answer is %d\n", var); + } +} diff --git a/libgomp/testsuite/libgomp.fortran/target-print-1.f90 b/libgomp/testsuite/libgomp.fortran/target-print-1.f90 new file mode 100644 index 00000000000..73ee09a2b79 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-print-1.f90 @@ -0,0 +1,14 @@ +! Ensure that printf on the offload device works. + +! { dg-do run } +! { dg-output "The answer is 42(\n|\r\n|\r)+" } + +program main + implicit none + integer :: var = 42 + +!$omp target + write (0, '("The answer is ", I2)') var +!$omp end target + +end program main diff --git a/libgomp/testsuite/libgomp.oacc-c/print-1.c b/libgomp/testsuite/libgomp.oacc-c/print-1.c new file mode 100644 index 00000000000..593885b5c2c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c/print-1.c @@ -0,0 +1,17 @@ +/* Ensure that printf on the offload device works. */ + +/* { dg-do run } */ +/* { dg-output "The answer is 42(\n|\r\n|\r)+" } */ + +#include <stdio.h> + +int var = 42; + +int +main () +{ +#pragma acc parallel + { + printf ("The answer is %d\n", var); + } +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 new file mode 100644 index 00000000000..bef664df4fa --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/print-1.f90 @@ -0,0 +1,14 @@ +! Ensure that printf on the offload device works. + +! { dg-do run } +! { dg-output "The answer is 42(\n|\r\n|\r)+" } + +program main + implicit none + integer :: var = 42 + +!$acc parallel + write (0, '("The answer is ", I2)') var +!$acc end parallel + +end program main