message: Move OMP target variant to use kernel.
diff: @@ -77,41 +77,23 @@ void NESTED_INIT::runOpenMPTargetVariant(VariantID vid)
  } else if ( vid == RAJA_OpenMPTarget ) {
  -#if 1 // temporary implementation until RAJA::kernel works with OpenMP target
  -
  - NESTED_INIT_DATA_SETUP_OMP_TARGET;
  -
  - startTimer();
  - for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
  -
  - RAJA::forall<RAJA::omp_target_parallel_for_exec<NUMTEAMS>>(
  - RAJA::RangeSegment(0, nk), [=](Index_type k) {
  - for (Index_type j = 0; j < nj; ++j ) {
  - for (Index_type i = 0; i < ni; ++i ) {
  - NESTED_INIT_BODY;
  - }
  - }
  - });
  -
  - }
  - stopTimer();
  -
  - NESTED_INIT_DATA_TEARDOWN_OMP_TARGET;
  -
  -#else
  -
  NESTED_INIT_DATA_SETUP_OMP_TARGET;
  - using EXEC_POL = ...
  + using EXEC_POL =
  + RAJA::KernelPolicy<
  + RAJA::statement::Collapse<RAJA::omp_target_parallel_collapse_exec,
  + RAJA::ArgList<2, 1, 0>, // k, j, i
  + RAJA::statement::Lambda<0>
  + >
  + >;
  startTimer();
  for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
  - RAJA::kernel...
  - RAJA::make_tuple(RAJA::RangeSegment(0, ni),
  + RAJA::kernel<EXEC_POL>( RAJA::make_tuple(RAJA::RangeSegment(0, ni),
  RAJA::RangeSegment(0, nj),
  RAJA::RangeSegment(0, nk)),
  - [=] __device__ (Index_type i, Index_type j, Index_type k) {
  + [=](Index_type i, Index_type j, Index_type k) {
  NESTED_INIT_BODY;
  });
  @@ -120,8 +102,6 @@ void NESTED_INIT::runOpenMPTargetVariant(VariantID vid)
  NESTED_INIT_DATA_TEARDOWN_OMP_TARGET;
  -#endif
  -
  } else {
  std::cout << "\n NESTED_INIT : Unknown variant id = " << vid << std::endl;
  }
  