[Openmp-dev] Avoid using is_device_ptr clause

[Itaru Kitayama via Openmp-dev]

Johannes,

The code like this?

#include <cstdlib>
#include <iostream>

extern "C" {
  void* llvm_omp_target_alloc_shared(size_t, int);
}

int main()
{
  std::intptr_t *p5 =
(std::intptr_t*)llvm_omp_target_alloc_shared(sizeof(int)*1024,0);
  p5[0]= 999;

#pragma omp target parallel for device(0)
  for (int i=0;i < 1024;i++) {
    std::printf("p5 %d %d\n", i, reinterpret_cast<int*>(p5)[0]);
  }
  return 0;
}

On Fri, Aug 27, 2021 at 7:40 AM Johannes Doerfert
<johannesdoerfert at gmail.com> wrote:
>
> The way I understood your problem is that you want to pass
> a device accessible pointer that is not mapped into a target
> region, right?
>
> When we implement 5.2 semantics you can pass pointers into
> a target region and they will not be nulled but instead retain
> their value if no mapping was found. That makes use of
> managed or device memory possible without declaring a mapping
> or adding the is_device_ptr clause.
>
> For now, you might get around is_device_ptr by type punning,
> e.g., cast the pointers into intptr_t before the target region
> and back inside.
>
> ~ Johannes
>
>
> On 8/26/21 3:55 PM, Itaru Kitayama via Openmp-dev wrote:
> > Joachim, all,
> > Is the example you pointed to me just an expansion of
> > llvm_omp_tarfet_alloc_shared, then?
> >
> > On Thu, Aug 26, 2021 at 20:59 Joachim Protze <protze.joachim at gmail.com>
> > wrote:
> >
> >> You can use the target pointer association API (see 5.12 in the new
> >> OpenMP Examples 5.1 document). This should help you avoid usage of
> >> is_device_ptr and actually get proper mapping behavior.
> >>
> >> - Joachim
> >>
> >> Am 26.08.21 um 01:41 schrieb Itaru Kitayama via Openmp-dev:
> >>> Is there a way to avoid adding many is_device_ptr clauses every time
> >>> program enters the target region? I like the llvm_target_alloc_shared API
> >>> can handle internally the device pointer and make it accessible from the
> >>> target region as well.
> >>> _______________________________________________
> >>> Openmp-dev mailing list
> >>> Openmp-dev at lists.llvm.org
> >>> https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev
> >>>
> >>
> >
> > _______________________________________________
> > Openmp-dev mailing list
> > Openmp-dev at lists.llvm.org
> > https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev