[Bug c++/110347] New: [OpenMP] private/firstprivate of a C++ member variable mishandled

[Bug c++/110347] New: [OpenMP] private/firstprivate of a C++ member variable mishandled

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=110347

            Bug ID: 110347
           Summary: [OpenMP] private/firstprivate of a C++ member variable
                    mishandled
           Product: gcc
           Version: 14.0
            Status: UNCONFIRMED
          Keywords: ice-on-valid-code, openmp, wrong-code
          Severity: normal
          Priority: P3
         Component: c++
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---

On the GCC side, this came up at
https://gcc.gnu.org/pipermail/gcc-patches/2022-June/596238.html

It was also discussed at length but with a different focus in a couple of
base-language OpenMP spec meetings, albeit Tom was only in one. I did not
quickly find an associated issue, maybe it was only shown on screen without an
associated issue by Kelvin or Deepak. ... I found the OpenMP spec issue → #3343
(see last comment for the result when Tom was there; I filed the issue based on
a previous discussion.)

Currently, we have for (assume a prior enter-data mapping of this[:1] if it
helps):
------------
struct C {
   int a, b, c, d;
   void foo (void) {
     #pragma omp target private (a) firstprivate(b)
      { a = 1; b = 2; c = 3; this->d = 4; }
   }
};
int main() { C c; c.foo(); }
------------

the original dump
------------
{
    int D.2925 [value-expr: ((struct C *) this)->a];
    int D.2924 [value-expr: ((struct C *) this)->b];
  #pragma omp target map(tofrom:*(struct C *) this [len: 16])
map(firstprivate:(struct C *) this [pointer assign, bias: 0])
firstprivate(D.2924) private(D.2925)
    {
      {
        <<cleanup_point <<< Unknown tree: expr_stmt
          (void) (D.2925 = 1) >>>>>;
        <<cleanup_point <<< Unknown tree: expr_stmt
          (void) (D.2924 = 2) >>>>>;
        <<cleanup_point <<< Unknown tree: expr_stmt
          (void) (((struct C *) this)->c = 3) >>>>>;
        <<cleanup_point <<< Unknown tree: expr_stmt
          (void) (((struct C *) this)->d = 4) >>>>>;
      }
    }
}
------------

the gimple dump
------------
void C::foo (struct C * const this)
{
  int a [value-expr: this->a];
  int b [value-expr: this->b];

  #pragma omp target num_teams(-2) thread_limit(0) map(tofrom:this [len:
8][implicit]) map(tofrom:*this [len: 16]) map(firstprivate:this [pointer
assign, bias: 0]) firstprivate(b) private(a)
    {
      this->a = 1;
      this->b = 2;
      this->c = 3;
      this->d = 4;
    }
}
------------


If I have understood the discussions and rules correctly, the expected result
would be:

{
    int D.2925;
    int D.2924 = ((struct C *) this)->b;
  #pragma omp target map(tofrom:*(struct C *) this [len: 16])
map(firstprivate:(struct C *) this [pointer assign, bias: 0])
firstprivate(D.2924) private(D.2925)
    {
      {
          D.2925 = 1;
          D.2924 = 2;
          ((struct C *) this)->c = 3;
          ((struct C *) this)->d = 4;
      }
    }
}

* * *

In terms of the spec, the following takes care of the
'firstprivate(cpp_member_var)' as it wouldn't be otherwise permitted in a
firstprivate:

"--- C++ ---
Unless otherwise specified, a variable that is part of another variable (as an
array element or a structure element) cannot be a variable list item, an
extended list item or locator list item except if the list appears on a clause
that is associated with a construct within a class non-static member function
and the variable is an accessible data member of the object for which the
non-static member function is invoked."
[OpenMP 5.2, 3.2.1 OpenMP Argument Lists [62:1-5])


Otherwise, if a clause should permit to use array sections or
structure-components as list item, that clause needs to explicit permit it. (In
5.2; in 5.1 it was the other way round.)

This wording still does not permit 'firstprivate(this->a)' - only
firstprivate(a).


I assume that for

 #pragma omp target map(this[:1]) firstprivate(a)
  { this->a = 5; a = 7;  printf("%d %d\n", a, this->a); }
  print ("%d\n", a);

the expected result would be: '7 5' and '5'.


* * *


Side remark: The following gives an ICE:

struct C {
   int a, b, c, d;
   void foo (void) {
     #pragma omp target defaultmap(firstprivate)
      { a = 1; b = 2; c = 3; d = 4; }
   }
};


Namely:

internal compiler error: in gimplify_adjust_omp_clauses, at gimplify.cc:13004
    4 |      #pragma omp target defaultmap(firstprivate)

* The ICE is new since GCC12

* GCC12+ accepts also 'map(this[:1])' - with the same ICE.
  while GCC11 rejected the latter with:
  'error: ‘this’ allowed in OpenMP only in ‘declare simd’ clauses'

* GCC 11 compiles the version above w/o map(this[:1]) and gimplifies it as
    defaultmap(firstprivate) firstprivate(this)
  and in the optimized dump, GCC 11 then has:
    .omp_data_arr.1.this = this_2(D); .omp_data_sizes.2[1] = {0};

[Bug c++/110347] [OpenMP] private/firstprivate of a C++ member variable mishandled

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=110347

--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> ---
I note that for the ICE example, the OG13 compiles it without an ICE and has as
result:
  firstprivate(this) map(tofrom:*this [len: 16])
  map(firstprivate:this [pointer assign, bias: 0])

For the main testcase of comment 0, OG13 gives the same result as maineline
(GCC 14).

[Bug c++/110347] [OpenMP] private/firstprivate of a C++ member variable mishandled

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=110347

--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
An explicit 'firstprivate(x)' will be turned in the compiler from a FIELD_DECL
to:

    int D.2935 [value-expr: ((struct t *) this)->x];
  #pragma omp target firstprivate(D.2934) firstprivate(D.2935)
    {
          (void) (D.2935 = 5)

in semantics.cc's omp_privatize_field, called by finish_omp_clauses for
handle_field_decl. The gimple dump then looks like:

  int x [value-expr: ((struct t *) this)->x];
  #pragma omp target ... firstprivate(x)
      map(alloc:MEM[(char *)this] [len: 0])
      map(firstprivate:this [pointer assign, bias: 0])
    {
      this->x = 5;

i.e. there is already a pointless 'this' mapping.

For 'private', we could do in omp_privatize_field a simple:
                  v = build_decl (input_location, VAR_DECL, DECL_NAME (t),
                                  TREE_TYPE (t));
but for 'firstprivate' that would miss the initialization - and adding a
pointless assignment is not really the best, especially not for larger objects
(like structs, arrays, reference types).

* * *

And for 'defaultmap(firstprivate)', the current code already adds 'this'
mapping in the original dump:

  #pragma omp target map(tofrom:*(struct t *) this [len: 44])
            map(firstprivate:(struct t *) this [pointer assign, bias: 0])
            defaultmap(firstprivate:all)
    {
      {
          (void) (((struct t *) this)->x = 5);

That's due to 'finish_omp_target_clauses_r' + data->this_expr_accessed = true;
it either needs to be suppressed here - or later in the ME removed again. The
former has the problem that 'defaultmap(firstprivate)' is not handled here, the
latter means a special case - ensuring that it is only removed if all member
accesses are for firstprivatized members.

While 'private' does not exist as defaultmap, compiler-internal handling or
(not checked) predefined/implicit mapping might.

[Bug c++/110347] [OpenMP] private/firstprivate of a C++ member variable mishandled

[cvs-commit at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=110347

--- Comment #3 from GCC Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Tobias Burnus <burnus@gcc.gnu.org>:

https://gcc.gnu.org/g:4f82d5a95a244d0aa4f8b2541b47a21bce8a191b

commit r14-9257-g4f82d5a95a244d0aa4f8b2541b47a21bce8a191b
Author: Jakub Jelinek <jakub@redhat.com>
Date:   Fri Mar 1 17:26:42 2024 +0100

    OpenMP/C++: Fix (first)private clause with member variables [PR110347]

    OpenMP permits '(first)private' for C++ member variables, which GCC handles
    by tagging those by DECL_OMP_PRIVATIZED_MEMBER, adding a temporary VAR_DECL
    and DECL_VALUE_EXPR pointing to the 'this->member_var' in the C++ front
end.

    The idea is that in omp-low.cc, the DECL_VALUE_EXPR is used before the
    region (for 'firstprivate'; ignored for 'private') while in the region,
    the DECL itself is used.

    In gimplify, the value expansion is suppressed and deferred if the
      lang_hooks.decls.omp_disregard_value_expr (decl, shared)
    returns true - which is never the case if 'shared' is true. In OpenMP 4.5,
    only 'map' and 'use_device_ptr' was permitted for the 'target' directive.
    And when OpenMP 5.0's 'private'/'firstprivate' clauses was added, the
    the update that now 'shared' argument could be false was missed. The
    respective check has now been added.

    2024-03-01  Jakub Jelinek  <jakub@redhat.com>
                Tobias Burnus  <tburnus@baylibre.com>

            PR c++/110347

    gcc/ChangeLog:

            * gimplify.cc (omp_notice_variable): Fix 'shared' arg to
            lang_hooks.decls.omp_disregard_value_expr for
            (first)private in target regions.

    libgomp/ChangeLog:

            * testsuite/libgomp.c++/target-lambda-3.C: Moved from
            gcc/testsuite/g++.dg/gomp/ and fixed is-mapped handling.
            * testsuite/libgomp.c++/target-lambda-1.C: Modify to also
            also work without offloading.
            * testsuite/libgomp.c++/firstprivate-1.C: New test.
            * testsuite/libgomp.c++/firstprivate-2.C: New test.
            * testsuite/libgomp.c++/private-1.C: New test.
            * testsuite/libgomp.c++/private-2.C: New test.
            * testsuite/libgomp.c++/target-lambda-4.C: New test.
            * testsuite/libgomp.c++/use_device_ptr-1.C: New test.

    gcc/testsuite/ChangeLog:

            * g++.dg/gomp/target-lambda-1.C: Moved to become a
            run-time test under testsuite/libgomp.c++.

    Co-authored-by: Tobias Burnus <tburnus@baylibre.com>

[Bug c++/110347] [OpenMP] private/firstprivate of a C++ member variable mishandled

[burnus at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=110347

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|UNCONFIRMED                 |RESOLVED
         Resolution|---                         |FIXED

--- Comment #4 from Tobias Burnus <burnus at gcc dot gnu.org> ---
FIXED on mainline (GCC 14).