diff --git a/mindspore/lite/CMakeLists.txt b/mindspore/lite/CMakeLists.txt
index 48fd514df5d..0933f221f6e 100644
--- a/mindspore/lite/CMakeLists.txt
+++ b/mindspore/lite/CMakeLists.txt
@@ -234,7 +234,7 @@ else()
         set(MSLITE_GPU_BACKEND "off")
     endif()
     if((NOT MSLITE_GPU_BACKEND STREQUAL "tensorrt") AND (NOT MSLITE_GPU_BACKEND STREQUAL "off") AND
-    (NOT MSLITE_GPU_BACKEND STREQUAL "opencl"))
+    (NOT MSLITE_GPU_BACKEND STREQUAL "cuda") AND (NOT MSLITE_GPU_BACKEND STREQUAL "opencl"))
         message("invalid MSLITE_GPU_BACKEND value ${MSLITE_GPU_BACKEND} for x86_64, MSLITE_GPU_BACKEND is set to off.")
         set(MSLITE_GPU_BACKEND "off")
     endif()
@@ -344,6 +344,12 @@ endif()
 if(MSVC)
   set(MSLITE_ENABLE_CONVERTER off)
 endif()
+
+if(MSLITE_GPU_BACKEND STREQUAL cuda)
+    set(MSLITE_ENABLE_CONVERTER on)
+    set(MSLITE_ENABLE_RUNTIME_GLOG on)
+endif()
+
 if(MSLITE_ENABLE_CONVERTER)
     add_compile_definitions(ENABLE_CONVERTER)
 endif()
@@ -562,9 +568,16 @@ if(MSLITE_GPU_BACKEND STREQUAL opencl)
 endif()
 
 if(MSLITE_GPU_BACKEND STREQUAL cuda)
-    add_definitions(-DGPU_CUDA)
+    # add_definitions(-DGPU_CUDA)
     add_compile_definitions(SUPPORT_GPU)
+    set(SUPPORT_CUDA on)
+    if(DEFINED ENV{CUDA_HOME})
+        message("CUDA_HOME = $ENV{CUDA_HOME}")
+    else()
+        message(FATAL_ERROR "please set CUDA_HOME, example: export CUDA_HOME=/usr/local/cuda-10.1/")
+    endif()
 endif()
+
 if(MSLITE_GPU_BACKEND STREQUAL tensorrt)
     add_compile_definitions(SUPPORT_GPU)
     set(SUPPORT_TENSORRT on)
diff --git a/mindspore/lite/src/CMakeLists.txt b/mindspore/lite/src/CMakeLists.txt
index 6a2dfa448ef..e16b0e8449e 100644
--- a/mindspore/lite/src/CMakeLists.txt
+++ b/mindspore/lite/src/CMakeLists.txt
@@ -423,10 +423,6 @@ if(MSLITE_GPU_BACKEND STREQUAL opencl)
     add_subdirectory(runtime/kernel/opencl)
     target_link_libraries(mindspore-lite opencl_kernel_mid)
     target_link_libraries(mindspore-lite_static opencl_kernel_mid)
-elseif(MSLITE_GPU_BACKEND STREQUAL cuda)
-    add_subdirectory(runtime/kernel/cuda)
-    target_link_libraries(mindspore-lite cuda_kernel_mid)
-    target_link_libraries(mindspore-lite_static cuda_kernel_mid)
 endif()
 
 if(SUPPORT_NPU)
@@ -498,6 +494,21 @@ if(MSLITE_ENABLE_RUNTIME_CONVERT)
     endif()
 endif()
 
+if(SUPPORT_CUDA)
+    set(CUDA_PATH $ENV{CUDA_HOME})
+    include_directories(${CCSRC_DIR}/plugin/device/gpu/kernel)
+    set(ENABLE_GPU on)
+    add_definitions(-DENABLE_GPU)
+    set(CUDA_VERSION 11.1)
+    include_directories(${CUDA_PATH})
+    include_directories(${CUDA_PATH}/include)
+    find_package(CUDA)
+    add_subdirectory(runtime/kernel/cuda)
+    list(APPEND CUDA_NVCC_FLAGS -arch=sm_53 --expt-relaxed-constexpr)
+    add_subdirectory(${CCSRC_DIR}/plugin/device/gpu/kernel/cuda_impl/cuda_ops cuda_ops)
+    target_link_libraries(mindspore-lite cuda_lite_kernel_mid cuda_ops mindspore_core)
+endif()
+
 if(MSLITE_ENABLE_RUNTIME_GLOG)
     target_link_libraries(mindspore-lite mindspore::glog)
     target_link_libraries(mindspore-lite_static mindspore::glog)
diff --git a/mindspore/lite/src/delegate/tensorrt/CMakeLists.txt b/mindspore/lite/src/delegate/tensorrt/CMakeLists.txt
index 17f51d07e2e..7769c213503 100644
--- a/mindspore/lite/src/delegate/tensorrt/CMakeLists.txt
+++ b/mindspore/lite/src/delegate/tensorrt/CMakeLists.txt
@@ -82,6 +82,7 @@ find_package(CUDA)
 file(GLOB_RECURSE CUDA_KERNEL_SRC
     ${CMAKE_CURRENT_SOURCE_DIR}/cuda_impl/*.cu
 )
+
 set_source_files_properties(${CUDA_KERNEL_SRC} PROPERTIES CUDA_SOURCE_PROPERTY_FORMAT OBJ)
 SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGES} -std=c++14 -fPIC")
 SET(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-std=c++14;)
diff --git a/mindspore/lite/src/delegate/tensorrt/cuda_impl/cublas_utils.h b/mindspore/lite/src/delegate/tensorrt/cuda_impl/cublas_utils.h
index 27e74d38263..3047d1881e1 100644
--- a/mindspore/lite/src/delegate/tensorrt/cuda_impl/cublas_utils.h
+++ b/mindspore/lite/src/delegate/tensorrt/cuda_impl/cublas_utils.h
@@ -19,6 +19,7 @@
 
 #include <cublas_v2.h>
 #include "src/delegate/tensorrt/cuda_impl/cuda_helper.h"
+#include "src/common/log_util.h"
 
 // cublas API error checking
 #define CUBLAS_CHECK_VOID(err)                        \
diff --git a/mindspore/lite/src/delegate/tensorrt/cuda_impl/cuda_helper.cc b/mindspore/lite/src/delegate/tensorrt/cuda_impl/cuda_helper.cc
index e7e915baf75..ad9aeffd3eb 100644
--- a/mindspore/lite/src/delegate/tensorrt/cuda_impl/cuda_helper.cc
+++ b/mindspore/lite/src/delegate/tensorrt/cuda_impl/cuda_helper.cc
@@ -15,6 +15,7 @@
  */
 
 #include "src/delegate/tensorrt/cuda_impl/cuda_helper.h"
+#include "src/common/log_util.h"
 
 CudaHelper &CudaHelper::GetInstance() {
   static CudaHelper instance;
diff --git a/mindspore/lite/src/delegate/tensorrt/cuda_impl/cuda_helper.h b/mindspore/lite/src/delegate/tensorrt/cuda_impl/cuda_helper.h
index 79c6ee7c99c..92c902fd354 100644
--- a/mindspore/lite/src/delegate/tensorrt/cuda_impl/cuda_helper.h
+++ b/mindspore/lite/src/delegate/tensorrt/cuda_impl/cuda_helper.h
@@ -19,7 +19,6 @@
 
 #include <cuda_runtime.h>
 #include <algorithm>
-#include "src/common/log_util.h"
 
 class CudaHelper {
  public:
@@ -39,22 +38,20 @@ class CudaHelper {
 #define GET_BLOCKS(total_threads) CudaHelper::GetInstance().GetBlocksNum(total_threads)
 #define GET_THREADS CudaHelper::GetInstance().GetThreadNum()
 
-#define CUDA_CHECK(ret)                                                          \
-  do {                                                                           \
-    cudaError_t cuda_ret = (ret);                                                \
-    if ((cuda_ret) != cudaSuccess) {                                             \
-      MS_LOG(ERROR) << "cuda func call error: " << cudaGetErrorString(cuda_ret); \
-      return -1;                                                                 \
-    }                                                                            \
+#define CUDA_CHECK(ret)              \
+  do {                               \
+    cudaError_t cuda_ret = (ret);    \
+    if ((cuda_ret) != cudaSuccess) { \
+      return -1;                     \
+    }                                \
   } while (0)
 
-#define CUDA_CHECK_VOID(ret)                                                     \
-  do {                                                                           \
-    cudaError_t cuda_ret = (ret);                                                \
-    if ((cuda_ret) != cudaSuccess) {                                             \
-      MS_LOG(ERROR) << "cuda func call error: " << cudaGetErrorString(cuda_ret); \
-      return;                                                                    \
-    }                                                                            \
+#define CUDA_CHECK_VOID(ret)         \
+  do {                               \
+    cudaError_t cuda_ret = (ret);    \
+    if ((cuda_ret) != cudaSuccess) { \
+      return;                        \
+    }                                \
   } while (0)
 
 #endif  // MINDSPORE_LITE_SRC_DELEGATE_TENSORRT_CDUA_IMPL_CUDA_HELPER_H_
diff --git a/mindspore/lite/src/delegate/tensorrt/cuda_impl/cudnn_utils.h b/mindspore/lite/src/delegate/tensorrt/cuda_impl/cudnn_utils.h
index 831eae2824a..2e71363fe0d 100644
--- a/mindspore/lite/src/delegate/tensorrt/cuda_impl/cudnn_utils.h
+++ b/mindspore/lite/src/delegate/tensorrt/cuda_impl/cudnn_utils.h
@@ -20,6 +20,7 @@
 #include <cudnn.h>
 #include <NvInfer.h>
 #include "src/delegate/tensorrt/cuda_impl/cuda_helper.h"
+#include "src/common/log_util.h"
 
 #define CUDNN_CHECK_VOID(err)                                            \
   do {                                                                   \
diff --git a/mindspore/lite/src/runtime/kernel/cuda/CMakeLists.txt b/mindspore/lite/src/runtime/kernel/cuda/CMakeLists.txt
new file mode 100644
index 00000000000..714d8d23e71
--- /dev/null
+++ b/mindspore/lite/src/runtime/kernel/cuda/CMakeLists.txt
@@ -0,0 +1,6 @@
+file(GLOB CUDA_LITE_KERNEL_SRC LIST_DIRECTORIES false
+    ${CMAKE_CURRENT_SOURCE_DIR}/*.cc
+    ${CCSRC_DIR}/plugin/device/gpu/kernel/cuda_impl/cuda_class/*.cc
+)
+
+add_library(cuda_lite_kernel_mid OBJECT ${CUDA_LITE_KERNEL_SRC})
diff --git a/mindspore/lite/src/runtime/kernel/cuda/batchtospace.cc b/mindspore/lite/src/runtime/kernel/cuda/batchtospace.cc
new file mode 100644
index 00000000000..bd22f50ce0c
--- /dev/null
+++ b/mindspore/lite/src/runtime/kernel/cuda/batchtospace.cc
@@ -0,0 +1,47 @@
+/**
+ * Copyright 2022 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "src/runtime/kernel/cuda/batchtospace.h"
+#include <memory>
+#include "nnacl/batch_to_space.h"
+
+namespace mindspore::kernel {
+int BatchtoSpaceCudaKernel::Prepare() {
+  CudaKernel::Prepare();
+  if (batch_to_space_helper_ == nullptr) {
+    batch_to_space_helper_ = std::make_shared<cukernel::BatchToSpaceHelperGpuKernel<float>>(type_name_);
+    helper_ = batch_to_space_helper_;
+  }
+  cukernel::BatchToSpaceAttr attr;
+  auto param = reinterpret_cast<BatchToSpaceParameter *>(op_parameter_);
+  attr.block_size = param->block_shape_[0];
+  attr.crops.push_back({param->crops_[0], param->crops_[1]});
+  attr.crops.push_back({param->crops_[2], param->crops_[3]});
+  int ret = batch_to_space_helper_->CheckKernelParam(&attr);
+  CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
+  ret = ReSize();
+  CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
+  return RET_OK;
+}
+
+int BatchtoSpaceCudaKernel::Run() {
+  int ret = batch_to_space_helper_->Process(input_device_ptrs_, output_device_ptrs_, work_device_ptrs_, stream_);
+  CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
+  return RET_OK;
+}
+
+// REG_KERNEL(kGPU, kNumberTypeFloat32, PrimitiveType_BatchToSpace, CudaKernelCreator<BatchtoSpaceCudaKernel>)
+}  // namespace mindspore::kernel
diff --git a/mindspore/lite/src/runtime/kernel/cuda/batchtospace.h b/mindspore/lite/src/runtime/kernel/cuda/batchtospace.h
new file mode 100644
index 00000000000..8e30e92fcb5
--- /dev/null
+++ b/mindspore/lite/src/runtime/kernel/cuda/batchtospace.h
@@ -0,0 +1,39 @@
+/**
+ * Copyright 2022 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifndef MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
+#define MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
+
+#include <memory>
+#include <vector>
+#include "src/runtime/kernel/cuda/cuda_kernel.h"
+#include "cuda_impl/cuda_class/batchtospace_helper.h"
+
+namespace mindspore::kernel {
+class BatchtoSpaceCudaKernel : public CudaKernel {
+ public:
+  BatchtoSpaceCudaKernel(OpParameter *parameter, const std::vector<lite::Tensor *> &inputs,
+                         const std::vector<lite::Tensor *> &outputs, const lite::InnerContext *ctx)
+      : CudaKernel(parameter, inputs, outputs, ctx) {}
+  ~BatchtoSpaceCudaKernel() override = default;
+  int Prepare() override;
+  int Run() override;
+
+ private:
+  std::shared_ptr<cukernel::BatchToSpaceHelperGpuKernel<float>> batch_to_space_helper_{nullptr};
+};
+}  // namespace mindspore::kernel
+#endif
diff --git a/mindspore/lite/src/runtime/kernel/cuda/cuda_kernel.cc b/mindspore/lite/src/runtime/kernel/cuda/cuda_kernel.cc
new file mode 100644
index 00000000000..75178b9a082
--- /dev/null
+++ b/mindspore/lite/src/runtime/kernel/cuda/cuda_kernel.cc
@@ -0,0 +1,72 @@
+/**
+ * Copyright 2022 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "src/runtime/kernel/cuda/cuda_kernel.h"
+
+namespace mindspore::kernel {
+int CudaKernel::PreProcess() {
+  if (output_device_size_.size() == 0) {
+    for (size_t i = 0; i < out_tensors_.size(); i++) {
+      // allocator cudaMalloc mem_size: out_tensors_[i]->set_allocator(/*CudaAllocator*/)
+      output_device_size_.push_back(helper_->GetOutputSizeList()[i]);
+      output_device_ptrs_.push_back(out_tensors_[i]->MutableData());
+    }
+  } else {
+    for (size_t i = 0; i < out_tensors_.size(); i++) {
+      if (helper_->GetOutputSizeList()[i] > output_device_size_[i]) {
+        out_tensors_[i]->FreeData();
+        output_device_size_[i] = helper_->GetOutputSizeList()[i];
+        output_device_ptrs_[i] = out_tensors_[i]->MutableData();
+      }
+    }
+  }
+  for (size_t i = 0; i < in_tensors_.size(); i++) {
+    input_device_ptrs_[i] = in_tensors_[i]->data();
+  }
+  return RET_OK;
+}
+
+int CudaKernel::ReSize() {
+  // menory calculate
+  std::vector<std::vector<size_t>> input_shapes;
+  std::vector<std::vector<size_t>> output_shapes;
+  for (auto in : in_tensors_) {
+    std::vector<size_t> one_shape(in->shape().size());
+    for (size_t i = 0; i < in->shape().size(); i++) {
+      one_shape[i] = static_cast<size_t>(in->shape()[i]);
+    }
+    input_shapes.push_back(one_shape);
+  }
+  for (auto out : out_tensors_) {
+    std::vector<size_t> one_shape(out->shape().size());
+    for (size_t i = 0; i < out->shape().size(); i++) {
+      one_shape[i] = static_cast<size_t>(out->shape()[i]);
+    }
+    output_shapes.push_back(one_shape);
+  }
+  helper_->ResetResource();
+  auto ret = helper_->CalMemSize(input_shapes, output_shapes);
+  CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
+  return RET_OK;
+}
+int CudaKernel::PostProcess() {
+  for (size_t i = 0; i < in_tensors_.size(); i++) {
+    in_tensors_[i]->DecRefCount();
+  }
+  return RET_OK;
+}
+CudaKernel::~CudaKernel() { helper_ = nullptr; }
+}  // namespace mindspore::kernel
diff --git a/mindspore/lite/src/runtime/kernel/cuda/cuda_kernel.h b/mindspore/lite/src/runtime/kernel/cuda/cuda_kernel.h
new file mode 100644
index 00000000000..d87562560d5
--- /dev/null
+++ b/mindspore/lite/src/runtime/kernel/cuda/cuda_kernel.h
@@ -0,0 +1,62 @@
+/**
+ * Copyright 2022 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifndef MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_CUDA_KERNEL_H_
+#define MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_CUDA_KERNEL_H_
+
+#include <cuda_runtime_api.h>
+#include <vector>
+#include <string>
+#include <memory>
+#include "src/inner_kernel.h"
+#include "src/lite_kernel.h"
+#include "cuda_impl/cuda_class/helper_base.h"
+
+using mindspore::lite::RET_ERROR;
+using mindspore::lite::RET_OK;
+namespace mindspore::kernel {
+class CudaKernel : public InnerKernel {
+ public:
+  CudaKernel(OpParameter *parameter, const std::vector<lite::Tensor *> &inputs,
+             const std::vector<lite::Tensor *> &outputs, const lite::InnerContext *ctx)
+      : InnerKernel(parameter, inputs, outputs, ctx) {}
+  ~CudaKernel() override;
+  int Prepare() override {
+    type_name_ = std::string(EnumNamePrimitiveType(type()));
+    return RET_OK;
+  }
+  int PreProcess() override;
+  int PostProcess() override;
+  int ReSize() override;
+  int Run() override { return RET_ERROR; }
+
+ protected:
+  std::vector<size_t> output_device_size_;
+  std::vector<void *> input_device_ptrs_;
+  std::vector<void *> output_device_ptrs_;
+  std::vector<void *> work_device_ptrs_;
+  cudaStream_t stream_;
+  std::shared_ptr<cukernel::GpuKernelHelperBase> helper_{nullptr};
+  std::string type_name_;
+};
+template <class T>
+kernel::InnerKernel *CudaKernelCreator(const std::vector<lite::Tensor *> &inputs,
+                                       const std::vector<lite::Tensor *> &outputs, OpParameter *opParameter,
+                                       const lite::Context *ctx, const kernel::KernelKey &desc) {
+  return new (std::nothrow) T(opParameter, inputs, outputs, static_cast<const lite::InnerContext *>(ctx));
+}
+}  // namespace mindspore::kernel
+#endif
diff --git a/mindspore/lite/src/runtime/kernel/cuda/unary.cc b/mindspore/lite/src/runtime/kernel/cuda/unary.cc
new file mode 100644
index 00000000000..3ee741d1f7e
--- /dev/null
+++ b/mindspore/lite/src/runtime/kernel/cuda/unary.cc
@@ -0,0 +1,37 @@
+/**
+ * Copyright 2022 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "src/runtime/kernel/cuda/unary.h"
+#include <memory>
+
+namespace mindspore::kernel {
+int UnaryCudaKernel::Prepare() {
+  CudaKernel::Prepare();
+  if (unary_helper_ == nullptr) {
+    unary_helper_ = std::make_shared<cukernel::UnaryHelperGpuKernel<float>>(type_name_);
+    helper_ = unary_helper_;
+  }
+  int ret = ReSize();
+  CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
+  return RET_OK;
+}
+int UnaryCudaKernel::Run() {
+  int ret = unary_helper_->Process(input_device_ptrs_, output_device_ptrs_, work_device_ptrs_, stream_);
+  CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
+  return RET_OK;
+}
+// REG_KERNEL(kGPU, kNumberTypeFloat32, PrimitiveType_Log, CudaKernelCreator<UnaryCudaKernel>)
+}  // namespace mindspore::kernel
diff --git a/mindspore/lite/src/runtime/kernel/cuda/unary.h b/mindspore/lite/src/runtime/kernel/cuda/unary.h
new file mode 100644
index 00000000000..42080fd9ee6
--- /dev/null
+++ b/mindspore/lite/src/runtime/kernel/cuda/unary.h
@@ -0,0 +1,39 @@
+/**
+ * Copyright 2022 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifndef MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
+#define MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
+
+#include <memory>
+#include <vector>
+#include "src/runtime/kernel/cuda/cuda_kernel.h"
+#include "cuda_impl/cuda_class/unary_helper.h"
+
+namespace mindspore::kernel {
+class UnaryCudaKernel : public CudaKernel {
+ public:
+  UnaryCudaKernel(OpParameter *parameter, const std::vector<lite::Tensor *> &inputs,
+                  const std::vector<lite::Tensor *> &outputs, const lite::InnerContext *ctx)
+      : CudaKernel(parameter, inputs, outputs, ctx) {}
+  ~UnaryCudaKernel() override = default;  // cudaFree
+  int Prepare() override;
+  int Run() override;
+
+ private:
+  std::shared_ptr<cukernel::UnaryHelperGpuKernel<float>> unary_helper_{nullptr};
+};
+}  // namespace mindspore::kernel
+#endif
diff --git a/mindspore/lite/src/runtime/kernel/cuda/unique.cc b/mindspore/lite/src/runtime/kernel/cuda/unique.cc
new file mode 100644
index 00000000000..981643eaf66
--- /dev/null
+++ b/mindspore/lite/src/runtime/kernel/cuda/unique.cc
@@ -0,0 +1,47 @@
+/**
+ * Copyright 2022 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "src/runtime/kernel/cuda/unique.h"
+#include <memory>
+#include <vector>
+
+namespace mindspore::kernel {
+int UniqueCudaKernel::Prepare() {
+  CudaKernel::Prepare();
+  if (unique_helper_ == nullptr) {
+    unique_helper_ = std::make_shared<cukernel::UniqueHelperGpuKernel<int, int>>(type_name_);
+    helper_ = unique_helper_;
+  }
+  return RET_OK;
+}
+
+int UniqueCudaKernel::PostProcess() {
+  auto ret = CudaKernel::PostProcess();
+  CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
+  // set output tensor shape
+  std::vector<int> out_shape = out_tensors_[0]->shape();
+  out_shape[out_shape.size() - 1] = unique_helper_->GetOutSize();
+  out_tensors_[0]->set_shape(out_shape);
+  return RET_OK;
+}
+
+int UniqueCudaKernel::Run() {
+  int ret = unique_helper_->Process(input_device_ptrs_, output_device_ptrs_, work_device_ptrs_, stream_);
+  CHECK_NOT_EQUAL_RETURN(ret, RET_OK);
+  return RET_OK;
+}
+// REG_KERNEL(kGPU, kNumberTypeFloat32, PrimitiveType_Unique, CudaKernelCreator<UniqueCudaKernel>)
+}  // namespace mindspore::kernel
diff --git a/mindspore/lite/src/runtime/kernel/cuda/unique.h b/mindspore/lite/src/runtime/kernel/cuda/unique.h
new file mode 100644
index 00000000000..f34b5634948
--- /dev/null
+++ b/mindspore/lite/src/runtime/kernel/cuda/unique.h
@@ -0,0 +1,40 @@
+/**
+ * Copyright 2022 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#ifndef MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
+#define MINDSPORE_LITE_SRC_RUNTIME_KERNEL_CUDA_UNARY_H_
+
+#include <memory>
+#include <vector>
+#include "src/runtime/kernel/cuda/cuda_kernel.h"
+#include "cuda_impl/cuda_class/unique_helper.h"
+
+namespace mindspore::kernel {
+class UniqueCudaKernel : public CudaKernel {
+ public:
+  UniqueCudaKernel(OpParameter *parameter, const std::vector<lite::Tensor *> &inputs,
+                   const std::vector<lite::Tensor *> &outputs, const lite::InnerContext *ctx)
+      : CudaKernel(parameter, inputs, outputs, ctx) {}
+  ~UniqueCudaKernel() override = default;
+  int Prepare() override;
+  int PostProcess() override;
+  int Run() override;
+
+ private:
+  std::shared_ptr<cukernel::UniqueHelperGpuKernel<int, int>> unique_helper_{nullptr};
+};
+}  // namespace mindspore::kernel
+#endif
diff --git a/mindspore/lite/test/CMakeLists.txt b/mindspore/lite/test/CMakeLists.txt
index e48ee09a06f..2043b5bb9ea 100644
--- a/mindspore/lite/test/CMakeLists.txt
+++ b/mindspore/lite/test/CMakeLists.txt
@@ -72,6 +72,17 @@ if(MSLITE_GPU_BACKEND STREQUAL opencl)
     endif()
 endif()
 
+if(MSLITE_GPU_BACKEND STREQUAL cuda)
+    set(CUDA_PATH $ENV{CUDA_HOME})
+    include_directories(${CUDA_PATH}/include)
+    include_directories(${CCSRC_DIR}/plugin/device/gpu/kernel)
+    include_directories(${CUDA_PATH})
+    file(GLOB_RECURSE TEST_GPU_UT_SRC
+            ${TEST_DIR}/ut/src/runtime/kernel/cuda/*.cc
+            )
+    list(APPEND TEST_UT_SRC ${TEST_GPU_UT_SRC})
+endif()
+
 if(MSLITE_ENABLE_INT8)
     file(GLOB_RECURSE TEST_INT8_UT_SRC
             ${TEST_DIR}/ut/src/runtime/kernel/arm/int8/*.cc
diff --git a/mindspore/lite/test/ut/src/runtime/kernel/cuda/batchtospace_tests.cc b/mindspore/lite/test/ut/src/runtime/kernel/cuda/batchtospace_tests.cc
new file mode 100644
index 00000000000..18e7c5f325c
--- /dev/null
+++ b/mindspore/lite/test/ut/src/runtime/kernel/cuda/batchtospace_tests.cc
@@ -0,0 +1,84 @@
+/**
+ * Copyright 2022 Huawei Technologies Co., Ltd
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "common/common_test.h"
+#include "schema/ops_generated.h"
+#include "src/runtime/kernel/cuda/batchtospace.h"
+#include "ut/src/runtime/kernel/cuda/common.h"
+#include "nnacl/batch_to_space.h"
+
+namespace mindspore {
+class CudaTest_BatchToSpace : public CommonTest {
+ public:
+  CudaTest_BatchToSpace() {}
+};
+namespace {
+// input: [batch*block_size*block_size, height_pad/block_size, width_pad/block_size, depth]
+// crops = [[crop_top, crop_bottom], [crop_left, crop_right]]
+// height = height_pad - crop_top - crop_bottom
+// width = width_pad - crop_left - crop_right
+OpParameter *CreateParameter() {
+  auto *param = mindspore::lite::cuda::test::CreateParameter<BatchToSpaceParameter>(schema::PrimitiveType_BatchToSpace);
+  param->block_shape_[0] = 2;
+  param->block_shape_[1] = 2;
+  param->crops_[0] = 0;
+  param->crops_[1] = 0;
+  param->crops_[THIRD_INPUT] = 0;
+  param->crops_[FOURTH_INPUT] = 0;
+  param->no_crop_ = true;
+  return reinterpret_cast<OpParameter *>(param);
+}
+}  // namespace
+TEST_F(CudaTest_BatchToSpace, basic) {
+  std::vector<int> input_shape{4, 1, 1, 1};
+  lite::Tensor *in_tensor = new (std::nothrow) lite::Tensor(TypeId::kNumberTypeFloat32, input_shape);
+  std::vector<lite::Tensor *> inputs{in_tensor};
+  std::vector<int> output_shape{1, 2, 2, 1};
+  lite::Tensor *out_tensor = new (std::nothrow) lite::Tensor(TypeId::kNumberTypeFloat32, output_shape);
+  std::vector<lite::Tensor *> outputs{out_tensor};
+  lite::InnerContext *ctx = new (std::nothrow) lite::InnerContext();
+  kernel::BatchtoSpaceCudaKernel *kernel =
+    new (std::nothrow) kernel::BatchtoSpaceCudaKernel(CreateParameter(), inputs, outputs, ctx);
+
+  void *input_device_ptr = nullptr;
+  cudaMalloc(&input_device_ptr, 4 * sizeof(float));
+  float input_host_ptr[4]{1, 2, 3, 4};
+  cudaMemcpy(input_device_ptr, input_host_ptr, 4 * sizeof(float), cudaMemcpyHostToDevice);
+  in_tensor->set_data(input_device_ptr);
+
+  void *output_device_ptr = nullptr;
+  cudaMalloc(&output_device_ptr, 4 * sizeof(float));
+  float output_host_ptr[4];
+  out_tensor->set_data(output_device_ptr);
+
+  kernel->Prepare();
+
+  kernel->Run();
+
+  cudaMemcpy(output_host_ptr, output_device_ptr, 4 * sizeof(float), cudaMemcpyDeviceToHost);
+  for (int i = 0; i < 4; i++) {
+    MS_LOG(ERROR) << "BatchtoSpaceCudaKernel out: " << output_host_ptr[i];
+  }
+
+  cudaFree(output_device_ptr);
+  cudaFree(input_device_ptr);
+  in_tensor->set_data(nullptr);
+  out_tensor->set_data(nullptr);
+  delete in_tensor;
+  delete out_tensor;
+  delete kernel;
+}
+}  // namespace mindspore