                                               const TArray<int64_t> input_pitches,
                                               const TArray<fast_divmod> output_div_pitches,
                                               const TArray<fast_divmod> scales_div,
-                                              const T* input_data,
-                                              T* output_data,
+                                              const T* __restrict__ input_data,
+                                              T* __restrict__ output_data,
                                               const size_t N) {
   CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N);
   CUDA_LONG input_index = 0;

                                       const TArray<int64_t> input_pitches,
                                       const TArray<fast_divmod> output_div_pitches,
                                       const TArray<fast_divmod> scales_div,
                                       const T* input_data,
                                       T* output_data,
                                       const size_t N) {
  CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N);
  CUDA_LONG input_index = 0;

                 const size_t N) {
  int blocksPerGrid = (int)(ceil(static_cast<float>(N) / GridDim::maxThreadsPerBlock));
  if (onnxruntime::UpsampleMode::NN == upsample_mode) {
    if (rank == 4) {
      _UpampleNearestKernel<T,4><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(
          input_pitches, output_div_pitches, scales_div,
          input_data, output_data, N);
    } else if (rank == 3) {
      _UpampleNearestKernel<T,3><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(
          input_pitches, output_div_pitches, scales_div,
          input_data, output_data, N);
    } else if (rank == 2) {
      _UpampleNearestKernel<T,2><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(
          input_pitches, output_div_pitches, scales_div,
          input_data, output_data, N);
    } else if (rank == 1) {
      _UpampleNearestKernel<T,1><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(
          input_pitches, output_div_pitches, scales_div,
          input_data, output_data, N);
    } else {
      ORT_THROW("Unsupported rank by the Upsample CUDA kernel");
    }
  } else if (onnxruntime::UpsampleMode::LINEAR) {
    if (rank == 4) {
      _UpampleBilinear4DInputKernel<T><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(
          input_dim2, input_pitches, output_div_pitches, scales_div,
          input_data, output_data, N);
    } else if (rank == 2) {
      _UpampleBilinear2DInputKernel<T><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(
          input_dim2, input_pitches, output_div_pitches, scales_div,
          input_data, output_data, N);
    } else {
      ORT_THROW("Unsupported rank by the Upsample CUDA kernel");
    }
  }
}

