OpenMP target misses LoopData privatization

[khaki3]

Hi. I'm working on the NVHPC compiler. We are testing RAJAPerf since it is a part of OpenMP 4.5 Test Suite. I'm aware that RAJA does not actively support OpenMP target.

It seems to me that policy/openmp_target/forall.hpp#L130 is missing the privatization of LoopData. If I modify pattern/kernel/internal/LoopData.hpp#L206 to declare a value type (data_t data), the code pasted below works. But I know the implementation of LoopData is shared among policies, and forall.hpp needs another workaround.

#include "RAJA/RAJA.hpp"
#include "RAJA/policy/cuda/raja_cudaerrchk.hpp"

int main()
{
  int *a;
  int N = 10;
  cudaMallocManaged(&a, sizeof(int) * N, cudaMemAttachGlobal);

  using Index_type = std::ptrdiff_t;

  using EXEC_POL =
    RAJA::KernelPolicy<
      RAJA::statement::For<0, RAJA::omp_target_parallel_for_exec<256>,
        RAJA::statement::Lambda<0>>>;

  RAJA::kernel<EXEC_POL>(
    RAJA::make_tuple(RAJA::RangeSegment{1,5}),
    [=] (Index_type i) {
      a[i]=i;
      printf("%d\n", (int)i);
    });

  if (a[4] == 4)
    puts("PASS");
  else
    puts("FAIL");

  return 0;
}