[clang] [llvm] [Clang][OpenMP] Capture mapped pointers on `target` by reference. (PR #145454)
Alexey Bataev via llvm-commits
llvm-commits at lists.llvm.org
Wed Jul 23 07:38:09 PDT 2025
alexey-bataev wrote:
> For the following:
>
> ```c
> int *p;
> #pragma omp target map(p[0]) // (A)
> (void)p;
>
> #pragma omp target map(p) // (B)
> (void)p;
>
> #pragma omp target map(p, p[0]) // (C)
> (void)p;
>
> #pragma omp target map(p[0], p) // (D)
> (void)p;
> ```
>
> For (A), the pointer `p` is predetermined `firstprivate`, so it should be (and is) captured by-copy. However, for (B), (C), and (D), since `p` is already listed in a `map` clause, it's not predetermined `firstprivate`, and hence, should be captured by-reference, like any other mapped variable.
>
> To ensure the correct handling of (C) and (D), the following changes were made:
>
> 1. In SemaOpenMP, we now ensure that `p` is marked to be captured by-reference in these cases.
> 2. We no longer ignore `map(p)` during codegen of `target` constructs, even if there's another map like `map(p[0])` that would have been mapped using a PTR_AND_OBJ map.
> 3. For cases like (D), we now handle `map(p)` before `map(p[0])`, so the former gets the TARGET_PARAM flag and sets the kernel argument.
I have one big question - why do we need this? If pointer is mapped explicitly, it becomes an attached pointer, which cannot be modified in target region. How we capture/process it, byref or byval, does not matter, we can do whatever is better. So, why do we need to process it byref?
https://github.com/llvm/llvm-project/pull/145454
More information about the llvm-commits
mailing list