[BUG][OpenMP][runtime]:The OpenMP runtime memory mapping table cannot be properly deleted.

[bmq7911]

Hello, I’ve encountered an issue with OpenMP in LLVM 17. While testing OpenMP memory mapping, I found that when using multi-level chained pointer access, the runtime mapping table cannot be correctly deleted due to reference counting.

For example, consider the following OpenMP directives:

#pragma omp target enter data map(to: s.ps->ps->ps->ps->i)
#pragma omp target exit data map(from: s.ps->ps->ps->ps->i)

After the enter data directive, the reference counts of the mapped objects are:

s.ps(1) -> ps(2) -> ps(2) -> ps(2) -> i(1)

Then, during exit data, the directive reduces the reference counts as follows:

s.ps(0) -> ps(1) -> ps(1) -> ps(1) -> i(0)

As a result, only the mappings for s.ps and i are deleted, while the intermediate mappings remain, which prevents proper cleanup.

[ppanchad-amd]

Hi @bmq7911. Internal ticket has been created to investigate this issue. Thanks!

[adityas-amd]

Hi @bmq7911, I reproduced a minimal script "test.txt" and "log.txt" confirms this as the intermediate values are still alive.
ROCm: 6.4
Clang/LLVM: 19.x
Device: (gfx1100)
log.txt
test.txt

this needs further investigation; we will update you.

[adityas-amd]

@bmq7911 llvm#153683 PR is in progress

[adityas-amd]

tested with the new large OpenMP patch llvm#153683, with the test file
test (1).txt
the compiler and runtime skip all intermediate mappings and works like " one map + one attach check + clean detach". I have the test results attached before and after LIBOMPTARGET_DEBUG=3 logs. After the patch, the mapping table is empty