message: omp-target for Polybench_3MM
diff: @@ -388,7 +388,58 @@ void POLYBENCH_3MM::runKernel(VariantID vid)
  break;
  }
  -#endif
  +
  +
  +#if defined(RAJA_ENABLE_TARGET_OPENMP)
  +#define NUMTEAMS 128
  + case RAJA_OpenMPTarget: {
  + POLYBENCH_3MM_DATA;
  +
  + #pragma omp target enter data map(to: A[0:m_ni * m_nk],B[0:m_nk * m_nj], C[0:m_nj * m_nm], D[0:m_nm * m_nl], E[0: m_ni * m_nj], F[0: m_nj * m_nl], G[0: m_ni * m_nl])
  + startTimer();
  +
  + #pragma omp target data use_device_ptr(A,B,C,D,E,F,G)
  + for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
  +
  + RAJA::forall<RAJA::policy::omp::omp_target_parallel_for_exec<NUMTEAMS>>(0,ni * nj, [=](Index_type ii) {
  + Index_type i,j,k;
  + *(E + ii) = 0.0;
  + i = ii/nj; j = ii % nj;
  + for(k=0;k<nk;k++) {
  + POLYBENCH_3MM_BODY2;
  + }
  + });
  +
  + // #pragma omp target data use_device_ptr(C,D,F)
  +
  + RAJA::forall<RAJA::policy::omp::omp_target_parallel_for_exec<NUMTEAMS>>(0,nj * nl, [=](Index_type ii) {
  + Index_type j,l,m;
  + *(F + ii) = 0.0;
  + j = ii/nl; l = ii % nl;
  + for(m=0;m<nm;m++) {
  + POLYBENCH_3MM_BODY4;
  + }
  + });
  +
  + //#pragma omp target data use_device_ptr(E,F,G)
  +
  + RAJA::forall<RAJA::policy::omp::omp_target_parallel_for_exec<NUMTEAMS>>(0,ni * nl, [=](Index_type ii) {
  + Index_type i,l,j;
  + *(G + ii) = 0.0;
  + i = ii/nl; l = ii % nl;
  + for(j=0;j<nj;j++) {
  + POLYBENCH_3MM_BODY6;
  + }
  + });
  + } // for run_reps
  +
  + stopTimer();
  + #pragma omp target exit data map(from:G[0:m_ni * m_nl]) map(delete: A[0:m_ni * m_nk],B[0:m_nk * m_nj], C[0:m_nj * m_nm], D[0:m_nm * m_nl], E[0: m_ni * m_nj], F[0: m_nj * m_nl])
  +
  + break;
  + }
  +#endif //RAJA_ENABLE_TARGET_OPENMP
  +#endif //RAJA_ENABLE_OMP
  #if defined(RAJA_ENABLE_CUDA)
  case Base_CUDA : {
  