                                               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;

namespace onnxruntime {
namespace cuda {

template <typename T, int RANK>
__global__ void _UpampleNearestKernel(const TArray<int64_t> input_pitches,
                                      const TArray<fast_divmod> output_div_pitches,
                                      const TArray<fast_divmod> scales_div,
                                      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;
  CUDA_LONG output_index = id;

  int div, mod;
  for (int dim = 0; dim < RANK; ++dim) {
    output_div_pitches[dim].divmod(output_index, div, mod);
    output_index = mod;
    if (scales_div[dim].d_ != 1 && div > 0) {
