Add OpenACC 'acc_map_data' variant to 'libgomp.oacc-c-c++-common/deep-copy-8.c' (was: [PATCH 11/13] OpenACC 2.6 deep copy: C and C++ execution tests)

Message ID 87r0lauc2k.fsf@euler.schwinge.homeip.net
State Accepted
Headers
Series Add OpenACC 'acc_map_data' variant to 'libgomp.oacc-c-c++-common/deep-copy-8.c' (was: [PATCH 11/13] OpenACC 2.6 deep copy: C and C++ execution tests) |

Checks

Context Check Description
snail/gcc-patch-check success Github commit url

Commit Message

Thomas Schwinge Oct. 31, 2023, 2 p.m. UTC
  Hi!

On 2019-12-17T22:04:54-0800, Julian Brown <julian@codesourcery.com> wrote:
> This patch has been broken out of the "OpenACC 2.6 manual deep copy
> support" patch, [...]

> This part adds C and C++ execution tests to libgomp.

Pushed to master branch commit 3e888f94624294d2b9b34ebfee0916768e5d9c3f
"Add OpenACC 'acc_map_data' variant to 'libgomp.oacc-c-c++-common/deep-copy-8.c'",
see attached.

This will be helpful to detect (and then avoid) a regression due to an
libgomp patch elsewhere.


Grüße
 Thomas


>  .../testsuite/libgomp.oacc-c++/deep-copy-12.C | 72 +++++++++++++++
>  .../testsuite/libgomp.oacc-c++/deep-copy-13.C | 72 +++++++++++++++
>  .../libgomp.oacc-c-c++-common/deep-copy-1.c   | 24 +++++
>  .../libgomp.oacc-c-c++-common/deep-copy-10.c  | 53 +++++++++++
>  .../libgomp.oacc-c-c++-common/deep-copy-11.c  | 72 +++++++++++++++
>  .../libgomp.oacc-c-c++-common/deep-copy-14.c  | 63 ++++++++++++++
>  .../libgomp.oacc-c-c++-common/deep-copy-2.c   | 29 +++++++
>  .../libgomp.oacc-c-c++-common/deep-copy-4.c   | 87 +++++++++++++++++++
>  .../libgomp.oacc-c-c++-common/deep-copy-6.c   | 59 +++++++++++++
>  .../libgomp.oacc-c-c++-common/deep-copy-7.c   | 45 ++++++++++
>  .../libgomp.oacc-c-c++-common/deep-copy-8.c   | 54 ++++++++++++
>  .../libgomp.oacc-c-c++-common/deep-copy-9.c   | 53 +++++++++++


-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  

Patch

From 3e888f94624294d2b9b34ebfee0916768e5d9c3f Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Mon, 30 Oct 2023 17:11:40 +0100
Subject: [PATCH] Add OpenACC 'acc_map_data' variant to
 'libgomp.oacc-c-c++-common/deep-copy-8.c'

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: Add OpenACC
	'acc_map_data' variant.
---
 .../libgomp.oacc-c-c++-common/deep-copy-8.c   | 29 +++++++++++++++++--
 1 file changed, 27 insertions(+), 2 deletions(-)

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
index 1b4cf2fb684..e705f78c311 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
@@ -12,8 +12,8 @@  struct dc
   int *d;
 };
 
-int
-main ()
+static void
+test (unsigned variant)
 {
   int n = 100, i, j, k;
   struct dc v = { .a = 3 };
@@ -22,7 +22,16 @@  main ()
   v.c = (int *) malloc (sizeof (int) * n);
   v.d = (int *) malloc (sizeof (int) * n);
 
+  if (variant & 1)
+    {
 #pragma acc enter data copyin(v)
+    }
+  else
+    {
+      void *v_d = acc_malloc (sizeof v);
+      acc_map_data (&v, v_d, sizeof v);
+      acc_memcpy_to_device (v_d, &v, sizeof v);
+    }
 
   for (k = 0; k < 16; k++)
     {
@@ -46,9 +55,25 @@  main ()
       assert (!acc_is_present (v.d, sizeof (int) * n));
     }
 
+  if (variant & 1)
+    {
 #pragma acc exit data copyout(v)
+    }
+  else
+    {
+      void *v_d = acc_deviceptr (&v);
+      acc_unmap_data (&v);
+      acc_free (v_d);
+    }
 
   assert (!acc_is_present (&v, sizeof (v)));
+}
+
+int
+main ()
+{
+  for (unsigned variant = 0; variant < 2; ++variant)
+    test (variant);
 
   return 0;
 }
-- 
2.34.1