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.
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; + } +}