diff --git a/src/IJ_mv/IJVector_parcsr_device.c b/src/IJ_mv/IJVector_parcsr_device.c index 595e6e718..8f5636c49 100644 --- a/src/IJ_mv/IJVector_parcsr_device.c +++ b/src/IJ_mv/IJVector_parcsr_device.c @@ -35,7 +35,8 @@ HYPRE_Int hypre_IJVectorAssembleSortAndReduce3(HYPRE_Int N0, HYPRE_BigInt *I0, c HYPRE_Int hypre_IJVectorAssembleSortAndReduce1(HYPRE_Int N0, HYPRE_BigInt *I0, char *X0, HYPRE_Complex *A0, HYPRE_Int *N1, HYPRE_BigInt **I1, char **X1, HYPRE_Complex **A1 ); -__global__ void hypreCUDAKernel_IJVectorAssemblePar(hypre_DeviceItem &item, HYPRE_Int n, HYPRE_Complex *x, +__global__ void hypreCUDAKernel_IJVectorAssemblePar(hypre_DeviceItem &item, HYPRE_Int n, + HYPRE_Complex *x, HYPRE_BigInt *map, HYPRE_BigInt offset, char *SorA, HYPRE_Complex *y); /* diff --git a/src/parcsr_ls/ads.c b/src/parcsr_ls/ads.c index fb87b12b4..01992a847 100644 --- a/src/parcsr_ls/ads.c +++ b/src/parcsr_ls/ads.c @@ -12,10 +12,12 @@ #include "_hypre_utilities.hpp" #if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) -__global__ void hypreCUDAKernel_AMSComputePi_copy1(hypre_DeviceItem &item, HYPRE_Int nnz, HYPRE_Int dim, +__global__ void hypreCUDAKernel_AMSComputePi_copy1(hypre_DeviceItem &item, HYPRE_Int nnz, + HYPRE_Int dim, HYPRE_Int *j_in, HYPRE_Int *j_out); -__global__ void hypreCUDAKernel_AMSComputePi_copy2(hypre_DeviceItem &item, HYPRE_Int nrows, HYPRE_Int dim, +__global__ void hypreCUDAKernel_AMSComputePi_copy2(hypre_DeviceItem &item, HYPRE_Int nrows, + HYPRE_Int dim, HYPRE_Int *i_in, HYPRE_Real *data_in, HYPRE_Real *Gx_data, HYPRE_Real *Gy_data, HYPRE_Real *Gz_data, HYPRE_Real *data_out); diff --git a/src/parcsr_ls/par_interp_device.c b/src/parcsr_ls/par_interp_device.c index 0d29f8e14..88b544166 100644 --- a/src/parcsr_ls/par_interp_device.c +++ b/src/parcsr_ls/par_interp_device.c @@ -24,7 +24,8 @@ __global__ void hypre_BoomerAMGBuildDirInterp_getcoef( hypre_DeviceItem &item, H HYPRE_Int *P_diag_i, HYPRE_Int *P_diag_j, HYPRE_Real *P_diag_data, HYPRE_Int *P_offd_i, HYPRE_Int *P_offd_j, HYPRE_Real *P_offd_data, HYPRE_Int *fine_to_coarse ); -__global__ void hypre_BoomerAMGBuildDirInterp_getcoef_v2( hypre_DeviceItem &item, HYPRE_Int nr_of_rows, +__global__ void hypre_BoomerAMGBuildDirInterp_getcoef_v2( hypre_DeviceItem &item, + HYPRE_Int nr_of_rows, HYPRE_Int *A_diag_i, HYPRE_Int *A_diag_j, HYPRE_Real *A_diag_data, HYPRE_Int *A_offd_i, HYPRE_Int *A_offd_j, HYPRE_Real *A_offd_data, HYPRE_Int *Soc_diag_j, HYPRE_Int *Soc_offd_j, HYPRE_Int *CF_marker, diff --git a/src/parcsr_ls/par_mod_multi_interp_device.c b/src/parcsr_ls/par_mod_multi_interp_device.c index e0729d28a..298b790b5 100644 --- a/src/parcsr_ls/par_mod_multi_interp_device.c +++ b/src/parcsr_ls/par_mod_multi_interp_device.c @@ -75,13 +75,15 @@ __global__ void hypreCUDAKernel_cfmarker_masked_rowsum( hypre_DeviceItem &item, HYPRE_Complex *A_offd_data, HYPRE_Int *CF_marker, HYPRE_Int *dof_func, HYPRE_Int *dof_func_offd, HYPRE_Complex *row_sums ); -__global__ void hypreCUDAKernel_generate_Pdiag_i_Poffd_i( hypre_DeviceItem &item, HYPRE_Int num_points, +__global__ void hypreCUDAKernel_generate_Pdiag_i_Poffd_i( hypre_DeviceItem &item, + HYPRE_Int num_points, HYPRE_Int color, HYPRE_Int *pass_order, HYPRE_Int *pass_marker, HYPRE_Int *pass_marker_offd, HYPRE_Int *S_diag_i, HYPRE_Int *S_diag_j, HYPRE_Int *S_offd_i, HYPRE_Int *S_offd_j, HYPRE_Int *P_diag_i, HYPRE_Int *P_offd_i ); -__global__ void hypreCUDAKernel_generate_Pdiag_j_Poffd_j( hypre_DeviceItem &item, HYPRE_Int num_points, +__global__ void hypreCUDAKernel_generate_Pdiag_j_Poffd_j( hypre_DeviceItem &item, + HYPRE_Int num_points, HYPRE_Int color, HYPRE_Int *pass_order, HYPRE_Int *pass_marker, HYPRE_Int *pass_marker_offd, HYPRE_Int *fine_to_coarse, HYPRE_Int *fine_to_coarse_offd, HYPRE_Int *A_diag_i, HYPRE_Int *A_diag_j, @@ -97,7 +99,8 @@ __global__ void hypreCUDAKernel_insert_remaining_weights( hypre_DeviceItem &item HYPRE_Int *Pi_offd_j, HYPRE_Real *Pi_offd_data, HYPRE_Int *P_offd_i, HYPRE_Int *P_offd_j, HYPRE_Real *P_offd_data ); -__global__ void hypreCUDAKernel_generate_Qdiag_j_Qoffd_j( hypre_DeviceItem &item, HYPRE_Int num_points, +__global__ void hypreCUDAKernel_generate_Qdiag_j_Qoffd_j( hypre_DeviceItem &item, + HYPRE_Int num_points, HYPRE_Int color, HYPRE_Int *pass_order, HYPRE_Int *pass_marker, HYPRE_Int *pass_marker_offd, HYPRE_Int *fine_to_coarse, HYPRE_Int *fine_to_coarse_offd, HYPRE_Int *A_diag_i, HYPRE_Int *A_diag_j, diff --git a/src/parcsr_mv/par_csr_matrix.c b/src/parcsr_mv/par_csr_matrix.c index 2bd785484..874fe92a4 100644 --- a/src/parcsr_mv/par_csr_matrix.c +++ b/src/parcsr_mv/par_csr_matrix.c @@ -670,7 +670,8 @@ hypre_ParCSRMatrixPrintIJ( const hypre_ParCSRMatrix *matrix, HYPRE_Int num_nonzeros_offd; HYPRE_BigInt ilower, iupper, jlower, jupper; - HYPRE_MemoryLocation memory_location = hypre_ParCSRMatrixMemoryLocation((hypre_ParCSRMatrix*) matrix); + HYPRE_MemoryLocation memory_location = + hypre_ParCSRMatrixMemoryLocation((hypre_ParCSRMatrix*) matrix); if (!matrix) { diff --git a/src/utilities/device_utils.c b/src/utilities/device_utils.c index 4e03a2e18..df02ee458 100644 --- a/src/utilities/device_utils.c +++ b/src/utilities/device_utils.c @@ -1314,7 +1314,8 @@ template HYPRE_Int hypreDevice_ScatterConstant(HYPRE_Complex *x, HYPRE_Int n, HY HYPRE_Complex v); __global__ void -hypreGPUKernel_DiagScaleVector(hypre_DeviceItem &item, HYPRE_Int n, HYPRE_Int *A_i, HYPRE_Complex *A_data, +hypreGPUKernel_DiagScaleVector(hypre_DeviceItem &item, HYPRE_Int n, HYPRE_Int *A_i, + HYPRE_Complex *A_data, HYPRE_Complex *x, HYPRE_Complex beta, HYPRE_Complex *y) { HYPRE_Int i = hypre_gpu_get_grid_thread_id<1, 1>(item);