This is an archive of the discontinued LLVM Phabricator instance.

[OpenMP] Add map info for dereference pointer.
ClosedPublic

Authored by jyu2 on Mar 1 2023, 11:20 AM.

Details

Summary

This is to fix run time problem when use:

int **a;
map((*a)[:3]) or (*a)[1].

current we skip generate map info for dereference pointer:
&(*a), &(*a)[0], 3*sizeof(int), TARGET_PARAM | TO | FROM

One way to fix runtime problem is to generate map info for dereference
pointer.

&(*a), &(*a), sizeof(pointer), ALLOC | TARGET_PARAM
&(*a), &(*a)[0], 3*sizeof(int), PTR_AND_OBJ | TO | FROM

The change in CGOpenMPRuntime.cpp add that.

The change in SemaOpenMP is to fix variable of dereference pointer to array
captured by reference. That is wrong. That cause run time to fail.

The rule is:
If variable is identified in a map clause it is always captured by
reference except if it is a pointer that is dereferenced somehow.

Diff Detail

Event Timeline

jyu2 created this revision.Mar 1 2023, 11:20 AM
Herald added a project: Restricted Project. · View Herald TranscriptMar 1 2023, 11:20 AM
jyu2 requested review of this revision.Mar 1 2023, 11:20 AM
ABataev added inline comments.Mar 1 2023, 12:52 PM
clang/test/OpenMP/target_map_deref_array_codegen.cpp
17

Is this correct at all? This is not a pointer to the array, it is an array of pointers, IIRC.

jyu2 added inline comments.Mar 1 2023, 4:24 PM
clang/test/OpenMP/target_map_deref_array_codegen.cpp
17

I am confuse about those two(pointer to array or array of pointer). I thinks

int **tid ===  int (*t1d)[3],

it is pointer to array.

ABataev added inline comments.Mar 1 2023, 4:30 PM
clang/test/OpenMP/target_map_deref_array_codegen.cpp
17

Nope. The first one is an array of pointers, the second - the pointer to the array of ints. They are different.

https://stackoverflow.com/questions/20120054/pointer-to-an-array-and-array-of-pointers

jyu2 edited the summary of this revision. (Show Details)Mar 1 2023, 6:05 PM

What result produces map(a[0][:3]?

jyu2 added a comment.Mar 2 2023, 8:37 AM

What result produces map(a[0][:3]?

Yes, that would be another way to fix the runtime problem. However the difficulty is when process array section, section base is different.
with a[0][:3]
the section base is a[0]
with (*a)[:3], the section base is (*a); It is hard to set a[0] as section base during processing the array section.

That is why I am adding dereference pointer not a[0].

What result produces map(a[0][:3]?

Yes, that would be another way to fix the runtime problem. However the difficulty is when process array section, section base is different.
with a[0][:3]
the section base is a[0]
with (*a)[:3], the section base is (*a); It is hard to set a[0] as section base during processing the array section.

That is why I am adding dereference pointer not a[0].

I mean we shall emit the same mapping for (*a)[:3] and for a[0][:3]

jyu2 added a comment.Mar 2 2023, 8:57 AM

I mean we shall emit the same mapping for (*a)[:3] and for a[0][:3]

Yes I mean emit (*a)[:3] as a[0][:3]

The difficulty is during the process array section of
OMPArraySectionExpr 0x12aa37e0 '<OpenMP array section type>' lvalue

-ImplicitCastExpr 0x12a8a918 'int *' <ArrayToPointerDecay>
`-ArraySubscriptExpr 0x12a8a8d8 'int[3]':'int[3]' lvalue
-ImplicitCastExpr 0x12a8a8c0 'int (*)[3]' <LValueToRValue>
`-DeclRefExpr 0x12a8a880 'int (*)[3]' lvalue ParmVar 0x12a8a000 'a' 'int (*)[3]'
`-IntegerLiteral 0x12a8a8a0 'int' 0
-<<<NULL>>>
-IntegerLiteral 0x12a8a8f8 'int' 3

`-<<<NULL>>>

It is not easy to set section base as a[0].

I mean we shall emit the same mapping for (*a)[:3] and for a[0][:3]

Yes I mean emit (*a)[:3] as a[0][:3]

The difficulty is during the process array section of
OMPArraySectionExpr 0x12aa37e0 '<OpenMP array section type>' lvalue

-ImplicitCastExpr 0x12a8a918 'int *' <ArrayToPointerDecay>
`-ArraySubscriptExpr 0x12a8a8d8 'int[3]':'int[3]' lvalue
-ImplicitCastExpr 0x12a8a8c0 'int (*)[3]' <LValueToRValue>
`-DeclRefExpr 0x12a8a880 'int (*)[3]' lvalue ParmVar 0x12a8a000 'a' 'int (*)[3]'
`-IntegerLiteral 0x12a8a8a0 'int' 0
-<<<NULL>>>
-IntegerLiteral 0x12a8a8f8 'int' 3

`-<<<NULL>>>

It is not easy to set section base as a[0].

Yes, I understand. And did not ask for it. I just mean that a[0][:3] emits different mapping data - TARGET_PAPARM|TO|FROM, PTR_AND_OBJ|TO|FROM

jyu2 updated this revision to Diff 501921.Mar 2 2023, 10:47 AM

Thanks Alexey for the review.

jyu2 added a comment.Mar 2 2023, 10:47 AM

Yes, I understand. And did not ask for it. I just mean that a[0][:3] emits different mapping data - TARGET_PAPARM|TO|FROM, PTR_AND_OBJ|TO|FROM

Oh I see. I was think alloc for adding pointer. But I now think I should just orignal maptype. Changed thanks.

ABataev added inline comments.Mar 2 2023, 11:00 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7488–7498

Why need a loop here? Can you somehow merge analysis for (*p) expression with the pointer subscript analysis?

clang/lib/Sema/SemaOpenMP.cpp
2206–2216

Why do you need a loop here?

jyu2 added inline comments.Mar 2 2023, 11:33 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7488–7498

What about if you have three dereference pointers like (***a)[0:3] or four pointers...

clang/lib/Sema/SemaOpenMP.cpp
2206–2216

What about if you have three dereference pointers like (***a)[0:3] or four pointers...

ABataev added inline comments.Mar 2 2023, 11:56 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7488–7498

Why you can't iterate through the required components just like for the array subscrit expression?

jyu2 added inline comments.Mar 2 2023, 12:30 PM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7488–7498

Because currently, for unarray operator the map info is skipped following the rule:

bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous;

So I need to check if it is dereference to array some how to not skip it.

ABataev added inline comments.Mar 2 2023, 12:31 PM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7488–7498
  1. Same problem in Sema too.
  2. Can you try to fix it to avoid those loops?
jyu2 updated this revision to Diff 502144.Mar 3 2023, 9:10 AM

Thanks Alexey for the code review. I removed loop by using last expression in the component for checking.

jyu2 marked an inline comment as done.Mar 3 2023, 9:11 AM
jyu2 added inline comments.
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7488–7498

Thanks. I removed loop.

ABataev added inline comments.Mar 3 2023, 9:17 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7490–7494

What if we have something like:

typedef int (T)[3];

T** p;
map(**p)

? Shall it work? I.e. mapping of the whole array by dereferening the pointer to the array.

jyu2 marked an inline comment as done.Mar 3 2023, 2:57 PM
jyu2 added inline comments.
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7490–7494

No, it is not work currently. What about *(*(p+a)+b)...

ABataev added inline comments.Mar 3 2023, 3:14 PM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7490–7494

My question - shall it work? Mapping (a)[:3] should result in the same in the same code for a if a is

typedef int(T)[3];
T** a;
jyu2 updated this revision to Diff 502842.Mar 6 2023, 3:52 PM

Add additional test as Alexey asked.

clang/lib/CodeGen/CGOpenMPRuntime.cpp
7490–7494

Oh, that should work. I add additional test for that.

typedef int(T)[3];
void bar()
{

T** a;
int b[2][3];
int (*p)[3] = b;
a =  &p;
printf("%p %p p %p\n", &a, b, *p);
for (int i = 0; i< 3; i++) {
  (**a)[1] = i;
}
#pragma omp target map((**a)[:3])
{
 (**a)[1]=5;
  // CHECK: 5
  printf("a = %d\n", (**a)[1]);
}
}
ABataev added inline comments.Mar 7 2023, 5:51 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7490–7494

No, I meant different thing:

T **a;
.. map (**a)
jyu2 added inline comments.Mar 7 2023, 7:51 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7490–7494

Yes, that is not working currently. That is something need to be fixed include *(*(a+i)+j) ....

ABataev added inline comments.Mar 7 2023, 7:53 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
7490–7494

Not sure need to implement such complex expression, but would be good to support (**a) (for arrays, array sections, just a regular scalars, etc.) here

jdoerfert retitled this revision from Add map info for dereference pointer. to [OpenMP] Add map info for dereference pointer..Mar 7 2023, 8:39 AM
jyu2 updated this revision to Diff 503413.Mar 8 2023, 10:13 AM

Okay, I add code to get **a to work. Thanks Alexey for the review.

This revision is now accepted and ready to land.Mar 8 2023, 10:17 AM
This revision was landed with ongoing or failed builds.Mar 8 2023, 5:59 PM
This revision was automatically updated to reflect the committed changes.

had not heard back , so reverted it for now
commit 8cf85a0cadb033fed3d96aa5283deb4bfbbaf2c8 (HEAD -> main, origin/main, origin/HEAD)
Author: Ron Lieberman <ron.lieberman@amd.com>
Date: Wed Mar 8 22:01:22 2023 -0600

Revert "Add map info for dereference pointer."

breaks amdgpu buildbot
jyu2 added a comment.Mar 9 2023, 11:33 AM

I turn off the test run to amdgpu. And rechecked in with
c5b060900e98

had not heard back , so reverted it for now
commit 8cf85a0cadb033fed3d96aa5283deb4bfbbaf2c8 (HEAD -> main, origin/main, origin/HEAD)
Author: Ron Lieberman <ron.lieberman@amd.com>
Date: Wed Mar 8 22:01:22 2023 -0600

Revert "Add map info for dereference pointer."

breaks amdgpu buildbot