                 index_of_dim1 * input_pitches[1] +
                 index_of_input_dim2 * input_pitches[2] +
                 index_of_input_dim3;
-  
+
   T x00 = input_data[input_index];
   T x10, x01, x11;
 

  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,

    // It's the end in dimension 3
    x10 = x00;
    x11 = x01;
  } else {
    x10 = input_data[input_index + 1];
    x11 = end_of_dim2 ? x10 : input_data[input_index + input_pitches[2] + 1];
  }
