Fix: cuda and rocm compiling error

Author: dyzheng

source/module_hsolver/kernels/cuda/math_kernel_op.cu

@@ -338,55 +338,6 @@ __global__ void vector_mul_vector_kernel(
     }
 }
 
-// vector operator: result[i] = vector[i] / constant
-template <>
-void vector_div_constant_op<double, base_device::DEVICE_GPU>::operator()(const int& dim,
-                                                                     double* result,
-                                                                     const double* vector,
-                                                                     const double constant)
-{
-    // In small cases, 1024 threads per block will only utilize 17 blocks, much less than 40
-    int thread = thread_per_block;
-    int block = (dim + thread - 1) / thread;
-    vector_div_constant_kernel<double><<<block, thread>>>(dim, result, vector, constant);
-
-    cudaCheckOnDebug();
-}
-
-template <typename FPTYPE>
-inline void vector_div_constant_wrapper(const int& dim,
-                                    std::complex<FPTYPE>* result,
-                                    const std::complex<FPTYPE>* vector,
-                                    const FPTYPE constant)
-{
-    thrust::complex<FPTYPE>* result_tmp = reinterpret_cast<thrust::complex<FPTYPE>*>(result);
-    const thrust::complex<FPTYPE>* vector_tmp = reinterpret_cast<const thrust::complex<FPTYPE>*>(vector);
-
-    int thread = thread_per_block;
-    int block = (dim + thread - 1) / thread;
-    vector_div_constant_kernel<thrust::complex<FPTYPE>><<<block, thread>>>(dim, result_tmp, vector_tmp, constant);
-
-    cudaCheckOnDebug();
-}
-
-template <>
-void vector_div_constant_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const int& dim,
-                                                                                  std::complex<float>* result,
-                                                                                  const std::complex<float>* vector,
-                                                                                  const float constant)
-{
-    vector_div_constant_wrapper(dim, result, vector, constant);
-}
-
-template <>
-void vector_div_constant_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const int& dim,
-                                                                                   std::complex<double>* result,
-                                                                                   const std::complex<double>* vector,
-                                                                                   const double constant)
-{
-    vector_div_constant_wrapper(dim, result, vector, constant);
-}
-
 template <typename T>
 __global__ void vector_div_vector_kernel(
     const int size,

source/module_hsolver/kernels/rocm/math_kernel_op.hip.cu

@@ -257,83 +257,6 @@ __global__ void vector_mul_vector_kernel(
     }
 }
 
-template <typename T>
-__launch_bounds__(1024) 
-__global__ void vector_div_constant_kernel(const int size,
-                                                                 T* result,
-                                                                 const T* vector,
-                                                                 const typename GetTypeReal<T>::type constant)
-{
-    int i = blockIdx.x * blockDim.x + threadIdx.x;
-    if (i < size)
-    {
-        result[i] = vector[i] / constant;
-    }
-}
-
-// vector operator: result[i] = vector[i] / constant
-template <>
-void vector_div_constant_op<double, base_device::DEVICE_GPU>::operator()(const int& dim,
-                                                                     double* result,
-                                                                     const double* vector,
-                                                                     const double constant)
-{
-    int thread = 1024;
-    int block = (dim + thread - 1) / thread;
-    hipLaunchKernelGGL(HIP_KERNEL_NAME(vector_div_constant_kernel<double>),
-                       dim3(block),
-                       dim3(thread),
-                       0,
-                       0,
-                       dim,
-                       result,
-                       vector,
-                       constant);
-
-    hipCheckOnDebug();
-}
-
-template <typename FPTYPE>
-inline void vector_div_constant_wrapper(const int& dim,
-                                    std::complex<FPTYPE>* result,
-                                    const std::complex<FPTYPE>* vector,
-                                    const FPTYPE constant)
-{
-    thrust::complex<FPTYPE>* result_tmp = reinterpret_cast<thrust::complex<FPTYPE>*>(result);
-    const thrust::complex<FPTYPE>* vector_tmp = reinterpret_cast<const thrust::complex<FPTYPE>*>(vector);
-    int thread = 1024;
-    int block = (dim + thread - 1) / thread;
-    hipLaunchKernelGGL(HIP_KERNEL_NAME(vector_div_constant_kernel<thrust::complex<FPTYPE>>),
-                       dim3(block),
-                       dim3(thread),
-                       0,
-                       0,
-                       dim,
-                       result_tmp,
-                       vector_tmp,
-                       constant);
-
-    hipCheckOnDebug();
-}
-
-template <>
-void vector_div_constant_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const int& dim,
-                                                                                  std::complex<float>* result,
-                                                                                  const std::complex<float>* vector,
-                                                                                  const float constant)
-{
-    vector_div_constant_wrapper(dim, result, vector, constant);
-}
-
-template <>
-void vector_div_constant_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const int& dim,
-                                                                                   std::complex<double>* result,
-                                                                                   const std::complex<double>* vector,
-                                                                                   const double constant)
-{
-    vector_div_constant_wrapper(dim, result, vector, constant);
-}
-
 template <typename T>
 __launch_bounds__(1024)
 __global__ void vector_div_vector_kernel(