Re: [PATCH][libgomp, testsuite, nvptx] Limit recursion in declare_target-{1,2}.f90

[Tom de Vries <tdevries@suse.de>]

[-- Attachment #1: Type: text/plain, Size: 1062 bytes --]

On 4/1/22 17:38, Jakub Jelinek wrote:
> On Fri, Apr 01, 2022 at 05:34:50PM +0200, Tom de Vries wrote:
>> Do you perhaps have an idea why it's failing?
> 
> Because you call on_device_arch_nvptx () outside of
> !$omp target region, so unless the host device is NVPTX,
> it will not be true.
> 

That bit does works because on_device_arch_nvptx calls on_device_arch 
which contains the omp target bit:
...
static int
on_device_arch (int d)
{
   int d_cur;
   #pragma omp target map(from:d_cur)
   d_cur = device_arch ();

   return d_cur == d;
}

int
on_device_arch_nvptx ()
{
   return on_device_arch (GOMP_DEVICE_NVIDIA_PTX);
}
...

So I realized that I didn't do a good job of specifying the problem I 
encountered, and went looking at it, at which point I realized the error 
message had changed, and knew how to fix it ... So, my apologies, some 
confusion on my part.

Anyway, attached patch avoids any nvptx-related tcl directives (just for 
once test-case for now).  To me, this seems the most robust solution.

It this approach acceptable?

Thanks,
- Tom

[-- Attachment #2: 0001-libgomp-testsuite-Fix-libgomp.fortran-examples-4-declare_target-1.f90.patch --]
[-- Type: text/x-patch, Size: 2648 bytes --]

[libgomp/testsuite] Fix libgomp.fortran/examples-4/declare_target-1.f90

---
 .../examples-4/declare_target-1.f90                | 31 +++++++++++++---------
 .../libgomp.fortran/examples-4/on_device_arch.c    |  3 +++
 2 files changed, 22 insertions(+), 12 deletions(-)

diff --git a/libgomp/testsuite/libgomp.fortran/examples-4/declare_target-1.f90 b/libgomp/testsuite/libgomp.fortran/examples-4/declare_target-1.f90
index 03c5c53ed67..acded20f756 100644
--- a/libgomp/testsuite/libgomp.fortran/examples-4/declare_target-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/examples-4/declare_target-1.f90
@@ -1,16 +1,6 @@
 ! { dg-do run }
-! { dg-additional-options "-cpp" }
-! Reduced from 25 to 23, otherwise execution runs out of thread stack on
-! Nvidia Titan V.
-! Reduced from 23 to 22, otherwise execution runs out of thread stack on
-! Nvidia T400 (2GB variant), when run with GOMP_NVPTX_JIT=-O0.
-! Reduced from 22 to 20, otherwise execution runs out of thread stack on
-! Nvidia RTX A2000 (6GB variant), when run with GOMP_NVPTX_JIT=-O0.
-! { dg-additional-options "-DREC_DEPTH=20" { target { offload_target_nvptx } } } */
-
-#ifndef REC_DEPTH
-#define REC_DEPTH 25
-#endif
+! { dg-additional-sources on_device_arch.c }
+! { dg-prune-output "command-line option '-fintrinsic-modules-path=.*' is valid for Fortran but not for C" }
 
 module e_53_1_mod
   integer :: THRESHOLD = 20
@@ -38,6 +28,23 @@ end module
 
 program e_53_1
   use e_53_1_mod, only : fib, fib_wrapper
+  integer :: REC_DEPTH = 25
+
+  interface
+    integer function on_device_arch_nvptx() bind(C)
+    end function on_device_arch_nvptx
+  end interface
+
+  if (on_device_arch_nvptx () /= 0) then
+     ! Reduced from 25 to 23, otherwise execution runs out of thread stack on
+     ! Nvidia Titan V.
+     ! Reduced from 23 to 22, otherwise execution runs out of thread stack on
+     ! Nvidia T400 (2GB variant), when run with GOMP_NVPTX_JIT=-O0.
+     ! Reduced from 22 to 20, otherwise execution runs out of thread stack on
+     ! Nvidia RTX A2000 (6GB variant), when run with GOMP_NVPTX_JIT=-O0.
+     REC_DEPTH = 20
+  end if
+
   if (fib (15) /= fib_wrapper (15)) stop 1
   if (fib (REC_DEPTH) /= fib_wrapper (REC_DEPTH)) stop 2
 end program
diff --git a/libgomp/testsuite/libgomp.fortran/examples-4/on_device_arch.c b/libgomp/testsuite/libgomp.fortran/examples-4/on_device_arch.c
new file mode 100644
index 00000000000..f8bef19e021
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/examples-4/on_device_arch.c
@@ -0,0 +1,3 @@
+/* Auxiliar file.  */
+/* { dg-do compile  { target skip-all-targets } } */
+#include "../../libgomp.c-c++-common/on_device_arch.h"