fix undefined hspi

Author: Qianruipku

source/module_hamilt_pw/hamilt_pwdft/kernels/cuda/ekinetic_op.cu

@@ -27,6 +27,10 @@ __global__ void ekinetic_pw(
       {
           hpsi[block_idx * max_npw + ii] = gk2[ii] * tpiba2 * psi[block_idx * max_npw + ii];
       }
+      for (int ii = npw + thread_idx; ii < max_npw; ii += blockDim.x)
+      {
+          hpsi[block_idx * max_npw + ii] = 0.0;
+      }
   }
   else
   {

source/module_hamilt_pw/hamilt_pwdft/kernels/ekinetic_op.cpp

@@ -26,6 +26,13 @@ struct ekinetic_pw_op<FPTYPE, base_device::DEVICE_CPU>
                 {
                     tmhpsi[ig] = gk2_ik[ig] * tpiba2 * tmpsi_in[ig];
                 }
+#ifdef _OPENMP
+#pragma omp parallel for
+#endif
+                for (int ig = npw; ig < max_npw; ++ig)
+                {
+                    tmhpsi[ig] = 0.0;
+                }
                 tmpsi_in += max_npw;
                 tmhpsi += max_npw;
             }

source/module_hamilt_pw/hamilt_pwdft/kernels/rocm/ekinetic_op.hip.cu

@@ -27,6 +27,10 @@ __global__ void ekinetic_pw(
       {
           hpsi[block_idx * max_npw + ii] = gk2[ii] * tpiba2 * psi[block_idx * max_npw + ii];
       }
+      for (int ii = npw + thread_idx; ii < max_npw; ii += blockDim.x)
+      {
+          hpsi[block_idx * max_npw + ii] = 0.0;
+      }
   }
   else
   {