This is an archive of the discontinued LLVM Phabricator instance.

[OPENMP]Fix PR49698: OpenMP declare mapper causes segmentation fault.
AcceptedPublic

Authored by ABataev on Apr 16 2021, 11:33 AM.

Details

Summary

The implicitly generated mappings for allocation/deallocation in mappers
runtime should be mapped as implicit, also no need to clear member_of
flag to avoid ref counter increment. Also, the ref counter should not be
incremented for the very first element that comes from the mapper
function.

Diff Detail

Event Timeline

ABataev created this revision.Apr 16 2021, 11:33 AM
ABataev requested review of this revision.Apr 16 2021, 11:33 AM
Herald added projects: Restricted Project, Restricted Project. · View Herald TranscriptApr 16 2021, 11:33 AM

Can we use the reproducer from the bug report, I want an outermost array section with objects that have a declare mapper.

abhinavgaba added inline comments.Apr 16 2021, 12:50 PM
openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp
7

Thanks for the fixes, Alexey. Are you planning on handling this case separately?

Can we use the reproducer from the bug report, I want an outermost array section with objects that have a declare mapper.

Sure, will add.

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp
7

Yes

ABataev updated this revision to Diff 338210.Apr 16 2021, 1:00 PM

Added repro from PR49698

jdoerfert accepted this revision.Apr 16 2021, 2:08 PM

LG, thanks for the quick fix.

This revision is now accepted and ready to land.Apr 16 2021, 2:08 PM
protze.joachim requested changes to this revision.Apr 19 2021, 4:29 AM

I tested the patch and still get wrong results (I modified the declare_mapper_nested_default_mappers_complex_structure.cpp test to use my verbose printf from bugzilla):

Modified Data Hierarchy:
Object C (0x623e50, 2) Contents:
        Object B (0x623e80, 2) Contents:
                Object A (0x6243f0) Contents:
                        data1 = 6439680  data2 = 0
                Object A (0x6243f8) Contents:
                        data1 = -784902216  data2 = 11062
        Object B (0x623e90, 2) Contents:
                Object A (0x6249e0) Contents:
                        data1 = 6441856  data2 = 0
                Object A (0x6249e8) Contents:
                        data1 = 6439904  data2 = 0
Object C (0x623e60, 2) Contents:
        Object B (0x623ef0, 2) Contents:
                Object A (0x624b90) Contents:
                        data1 = 6438736  data2 = 0
                Object A (0x624b98) Contents:
                        data1 = 6441344  data2 = 0
        Object B (0x623f00, 2) Contents:
                Object A (0x624c60) Contents:
                        data1 = 6444032  data2 = 0
                Object A (0x624c68) Contents:
                        data1 = 6441856  data2 = 0

The values inside the target region look good, but things break when mapping back.

When I remove the teams distribute for and only leave #pragma omp target, I get a different result and the test mistakenly passes:

Modified Data Hierarchy:
Object C (0x622e40, 2) Contents:
        Object B (0x622e70, 2) Contents:
                Object A (0x6233e0) Contents:
                        data1 = 6433120  data2 = 0
                Object A (0x6233e8) Contents:
                        data1 = 11  data2 = 22
        Object B (0x622e80, 2) Contents:
                Object A (0x6239d0) Contents:
                        data1 = 6435792  data2 = 0
                Object A (0x6239d8) Contents:
                        data1 = 11  data2 = 22
Object C (0x622e50, 2) Contents:
        Object B (0x622ee0, 2) Contents:
                Object A (0x623b80) Contents:
                        data1 = 6437312  data2 = 0
                Object A (0x623b88) Contents:
                        data1 = 11  data2 = 22
        Object B (0x622ef0, 2) Contents:
                Object A (0x623c50) Contents:
                        data1 = 6437744  data2 = 0
                Object A (0x623c58) Contents:
                        data1 = 11  data2 = 22
Testing for correctness...
outer[1].arr[1].arr[1].data2 = 22.

I'm currently only testing x86_64 offloading.

This revision now requires changes to proceed.Apr 19 2021, 4:29 AM

I tested the patch and still get wrong results (I modified the declare_mapper_nested_default_mappers_complex_structure.cpp test to use my verbose printf from bugzilla):

Modified Data Hierarchy:
Object C (0x623e50, 2) Contents:
        Object B (0x623e80, 2) Contents:
                Object A (0x6243f0) Contents:
                        data1 = 6439680  data2 = 0
                Object A (0x6243f8) Contents:
                        data1 = -784902216  data2 = 11062
        Object B (0x623e90, 2) Contents:
                Object A (0x6249e0) Contents:
                        data1 = 6441856  data2 = 0
                Object A (0x6249e8) Contents:
                        data1 = 6439904  data2 = 0
Object C (0x623e60, 2) Contents:
        Object B (0x623ef0, 2) Contents:
                Object A (0x624b90) Contents:
                        data1 = 6438736  data2 = 0
                Object A (0x624b98) Contents:
                        data1 = 6441344  data2 = 0
        Object B (0x623f00, 2) Contents:
                Object A (0x624c60) Contents:
                        data1 = 6444032  data2 = 0
                Object A (0x624c68) Contents:
                        data1 = 6441856  data2 = 0

The values inside the target region look good, but things break when mapping back.

When I remove the teams distribute for and only leave #pragma omp target, I get a different result and the test mistakenly passes:

Modified Data Hierarchy:
Object C (0x622e40, 2) Contents:
        Object B (0x622e70, 2) Contents:
                Object A (0x6233e0) Contents:
                        data1 = 6433120  data2 = 0
                Object A (0x6233e8) Contents:
                        data1 = 11  data2 = 22
        Object B (0x622e80, 2) Contents:
                Object A (0x6239d0) Contents:
                        data1 = 6435792  data2 = 0
                Object A (0x6239d8) Contents:
                        data1 = 11  data2 = 22
Object C (0x622e50, 2) Contents:
        Object B (0x622ee0, 2) Contents:
                Object A (0x623b80) Contents:
                        data1 = 6437312  data2 = 0
                Object A (0x623b88) Contents:
                        data1 = 11  data2 = 22
        Object B (0x622ef0, 2) Contents:
                Object A (0x623c50) Contents:
                        data1 = 6437744  data2 = 0
                Object A (0x623c58) Contents:
                        data1 = 11  data2 = 22
Testing for correctness...
outer[1].arr[1].arr[1].data2 = 22.

I'm currently only testing x86_64 offloading.

The problem is in transferring from the device. Looks like we need not only to invert the list of data in one data transfer block but also invert the list of blocks. Trying to implement this.

ABataev updated this revision to Diff 338624.Apr 19 2021, 1:28 PM

Fixes and rebase

protze.joachim accepted this revision.Apr 21 2021, 7:15 AM

The latest changes fix the issue. I tested with x86_64 and nvidia offloading.

The behavior and tests looks good to me, just one comment.

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp
123–126 ↗(On Diff #338624)

Do these asserts make sense with the assert in line 104?

This revision is now accepted and ready to land.Apr 21 2021, 7:15 AM
ABataev added inline comments.Apr 21 2021, 8:11 AM
openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp
123–126 ↗(On Diff #338624)

Ah, just forgot to remove these lines, thanks!

This revision was landed with ongoing or failed builds.Apr 21 2021, 10:39 AM
This revision was automatically updated to reflect the committed changes.
jdoerfert reopened this revision.Apr 23 2021, 6:25 PM
This revision is now accepted and ready to land.Apr 23 2021, 6:25 PM

I think this broke GridMini (https://github.com/meifeng/GridMini/tree/openmp) :(

Will investigate this on Monday.

I think this broke GridMini (https://github.com/meifeng/GridMini/tree/openmp) :(

Hmm, this app does not use mappers. Are you sure that this commit is the actual cause of the problem? Also, could you provide detailed debug log for the investigation (with LIBOMPTARGET_DEBUG set)?

jdoerfert added a comment.EditedApr 26 2021, 8:43 AM

I think this broke GridMini (https://github.com/meifeng/GridMini/tree/openmp) :(

Hmm, this app does not use mappers. Are you sure that this commit is the actual cause of the problem? Also, could you provide detailed debug log for the investigation (with LIBOMPTARGET_DEBUG set)?

I did debug something else and the clang-omp branch failed until I reverted this patch, at least that is what I thought. I will look again as soon as I find the time. For sure it run fine until very recently and we should make sure it continues to do so ;)

I think this broke GridMini (https://github.com/meifeng/GridMini/tree/openmp) :(

Hmm, this app does not use mappers. Are you sure that this commit is the actual cause of the problem? Also, could you provide detailed debug log for the investigation (with LIBOMPTARGET_DEBUG set)?

I did debug something else and the clang-omp branch failed until I reverted this patch, at least that is what I thought. I will look again as soon as I find the time. For sure it run fine until very recently and we should make sure it continues to do so ;)

Definitely! That's why I'm asking for the debug log for the investigation. It requires CUDA but I don't have an access to the machine with cuda.

cchen added a subscriber: cchen.Oct 1 2021, 3:28 PM
cchen added inline comments.
openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp
55

Do we really expect sa[1].f.b be the same as &x[0]? Shouldn't sa[1].f.b be the same as &y[0]?

69

Same here.

ABataev added inline comments.Oct 1 2021, 4:00 PM
openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp
55

Yes, looks so, missed it.