[C++,OpenACC/OpenMP] Allow static constexpr fields in mappable types
diff mbox series

Message ID 8b6d4f26-d11b-f564-3712-a6fdda4e4c2a@mentor.com
State New
Headers show
Series
  • [C++,OpenACC/OpenMP] Allow static constexpr fields in mappable types
Related show

Commit Message

Chung-Lin Tang Jan. 20, 2020, 4:32 p.m. UTC
Hi Jakub, Thomas,
We had a customer with a C++ program using GPU offloading failing to compile
due to the code's extensive use of 'static constexpr' in its many template
classes (code was using OpenMP, but OpenACC is no different)

While the FE should ensure that no static members should exist for struct/class
types that are being mapped to the GPU, 'static constexpr' are completely
resolved and folded statically during compile time, so they really shouldn't
count.

This is a small patch to cp/decl2.c:cp_omp_mappable_type_1() to allow the
DECL_DECLARED_CONSTEXPR_P == true case to be mapped, and a g++ testcase. Patch
has been tested with no regressions in g++ and libgomp testsuites.

Probably not okay for trunk now, okay for stage1?

Thanks,
Chung-Lin

	cp/
	* decl2.c (cp_omp_mappable_type_1): Allow fields with
	DECL_DECLARED_CONSTEXPR_P to be mapped.

	testsuite/
	* g++.dg/goacc/static-constexpr-1.C: New test.

Comments

Jakub Jelinek Jan. 20, 2020, 4:49 p.m. UTC | #1
On Tue, Jan 21, 2020 at 12:32:00AM +0800, Chung-Lin Tang wrote:
> Hi Jakub, Thomas,
> We had a customer with a C++ program using GPU offloading failing to compile
> due to the code's extensive use of 'static constexpr' in its many template
> classes (code was using OpenMP, but OpenACC is no different)
> 
> While the FE should ensure that no static members should exist for struct/class
> types that are being mapped to the GPU, 'static constexpr' are completely
> resolved and folded statically during compile time, so they really shouldn't
> count.
> 
> This is a small patch to cp/decl2.c:cp_omp_mappable_type_1() to allow the
> DECL_DECLARED_CONSTEXPR_P == true case to be mapped, and a g++ testcase. Patch
> has been tested with no regressions in g++ and libgomp testsuites.
> 
> Probably not okay for trunk now, okay for stage1?

This at least in OpenMP is something covered by the standard, so it can't be
changed just because somebody uses something in their code.
The OpenMP 4.5 definition of mappable type for C++ is that
  - All data members must be non-static.
among other requirements.  In OpenMP 5.0 that has been removed.
So, if we follow the 4.5 definition, it shouldn't change, if we follow 5.0
definition, the whole loop should be dropped, but in no case shall static
constexpr data members be treated any differently from any other static data
members.  There are many changes that need to be done for the 5.0 mapping and
just cherry-picking one random rule from there is not a good idea.

I have no idea what OpenACC says here.

If OpenACC doesn't have this restriction, then perhaps you need to pass down
whether it is OpenMP or OpenACC and decide based on that.

	Jakub

Patch
diff mbox series

diff --git a/gcc/cp/decl2.c b/gcc/cp/decl2.c
index 042d6fa12df..4f7d9b0ebd4 100644
--- a/gcc/cp/decl2.c
+++ b/gcc/cp/decl2.c
@@ -1461,7 +1461,10 @@  cp_omp_mappable_type_1 (tree type, bool notes)
     {
       tree field;
       for (field = TYPE_FIELDS (type); field; field = DECL_CHAIN (field))
-	if (VAR_P (field))
+	if (VAR_P (field)
+	    /* Fields that are 'static constexpr' can be folded away at compile
+	       time, thus does not interfere with mapping.  */
+	    && !DECL_DECLARED_CONSTEXPR_P (field))
 	  {
 	    if (notes)
 	      inform (DECL_SOURCE_LOCATION (field),
diff --git a/gcc/testsuite/g++.dg/goacc/static-constexpr-1.C b/gcc/testsuite/g++.dg/goacc/static-constexpr-1.C
new file mode 100644
index 00000000000..2bf69209de4
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/static-constexpr-1.C
@@ -0,0 +1,16 @@ 
+// { dg-do compile }
+
+/* Test that static constexpr members do not interfere with offloading.  */
+struct rec
+{
+  static constexpr int x = 1;
+  int y, z;
+};
+
+void foo (rec& r)
+{
+  #pragma acc parallel copy(r)
+  {
+    r.y = r.y = r.x;
+  }
+}