diff --git a/onnxruntime/core/providers/cuda/generator/range.cc b/onnxruntime/core/providers/cuda/generator/range.cc
index 347d54ee85..cc57d9d2a5 100644
--- a/onnxruntime/core/providers/cuda/generator/range.cc
+++ b/onnxruntime/core/providers/cuda/generator/range.cc
@@ -18,11 +18,15 @@ ONNX_OPERATOR_KERNEL_EX(
     kOnnxDomain,
     11,
     kCudaExecutionProvider,
-    KernelDefBuilder().TypeConstraint("T", {DataTypeImpl::GetTensorType<float>(),
-                                            DataTypeImpl::GetTensorType<double>(),
-                                            DataTypeImpl::GetTensorType<int16_t>(),
-                                            DataTypeImpl::GetTensorType<int32_t>(),
-                                            DataTypeImpl::GetTensorType<int64_t>()}),
+    KernelDefBuilder()
+        .InputMemoryType<OrtMemTypeCPUInput>(0)  // start
+        .InputMemoryType<OrtMemTypeCPUInput>(1)  // limit
+        .InputMemoryType<OrtMemTypeCPUInput>(2)  // delta
+        .TypeConstraint("T", {DataTypeImpl::GetTensorType<float>(),
+                              DataTypeImpl::GetTensorType<double>(),
+                              DataTypeImpl::GetTensorType<int16_t>(),
+                              DataTypeImpl::GetTensorType<int32_t>(),
+                              DataTypeImpl::GetTensorType<int64_t>()}),
     Range);
 
 template <typename T>
@@ -47,17 +51,13 @@ static Status ComputeRange(OpKernelContext* ctx) {
                            delta_tensor_ptr->Shape());
   }
 
-  // Start, Limit and Delta are stored in GPU. So we need copy it to CPU to read.
-  // It is better to store these tensors in pinned memory or CPU for better performance.
-  T start;
-  CUDA_RETURN_IF_ERROR(cudaMemcpy(&start, start_tensor.template Data<T>(), sizeof(T), cudaMemcpyDeviceToHost));
-
-  T limit;
-  CUDA_RETURN_IF_ERROR(cudaMemcpy(&limit, limit_tensor.template Data<T>(), sizeof(T), cudaMemcpyDeviceToHost));
+  // Start, Limit and Delta are stored in CPU. 
+  T start = *(start_tensor.template Data<T>());
+  T limit = *(limit_tensor.template Data<T>());
 
   T delta = T(1);
   if (delta_tensor_ptr != nullptr) {
-    CUDA_RETURN_IF_ERROR(cudaMemcpy(&delta, delta_tensor_ptr->template Data<T>(), sizeof(T), cudaMemcpyDeviceToHost));
+    delta = *(delta_tensor_ptr->template Data<T>());
   }
 
   if (delta == T(0)) {