| Message ID | 8b6d4f26-d11b-f564-3712-a6fdda4e4c2a@mentor.com |
|---|---|
| State | New |
| Headers | show |
| Series | [C++,OpenACC/OpenMP] Allow static constexpr fields in mappable types | expand |
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
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; + } +}