message: Converted CUDA variant to use more kernel features.
diff: @@ -135,7 +135,10 @@ void POLYBENCH_2MM::runCudaVariant(VariantID vid)
  RAJA::statement::CudaKernel<
  RAJA::statement::For<0, RAJA::cuda_block_exec,
  RAJA::statement::For<1, RAJA::cuda_thread_exec,
  - RAJA::statement::Lambda<0>
  + RAJA::statement::Lambda<0>,
  + RAJA::statement::For<2, RAJA::seq_exec,
  + RAJA::statement::Lambda<1>
  + >
  >
  >
  >
  @@ -145,30 +148,29 @@ void POLYBENCH_2MM::runCudaVariant(VariantID vid)
  for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
  RAJA::kernel<EXEC_POL>( RAJA::make_tuple(RAJA::RangeSegment(0, ni),
  - RAJA::RangeSegment(0, nj)),
  - [=] __device__ (Index_type i, Index_type j) {
  -
  + RAJA::RangeSegment(0, nj),
  + RAJA::RangeSegment{0, nk}),
  + [=] __device__ (Index_type i, Index_type j, Index_type k) {
  POLYBENCH_2MM_BODY1;
  - for (Index_type k=0;k<nk;k++) {
  + },
  + [=] __device__ (Index_type i, Index_type j, Index_type k) {
  POLYBENCH_2MM_BODY2;
  }
  -
  - });
  + );
  memcpy(m_D,m_DD,m_ni * m_nl * sizeof(Real_type));
  initCudaDeviceData(D,m_D,m_ni * m_nl );
  RAJA::kernel<EXEC_POL>( RAJA::make_tuple(RAJA::RangeSegment(0, ni),
  - RAJA::RangeSegment(0, nl)),
  - [=] __device__ (Index_type i, Index_type l) {
  -
  + RAJA::RangeSegment(0, nl),
  + RAJA::RangeSegment{0, nj}),
  + [=] __device__ (Index_type i, Index_type l, Index_type j) {
  POLYBENCH_2MM_BODY3;
  - for (Index_type j=0;j<nj;j++) {
  + },
  + [=] __device__ (Index_type i, Index_type l, Index_type j) {
  POLYBENCH_2MM_BODY4;
  }
  -
  - });
  -
  + );
  }
  stopTimer();
  