This commit is contained in:
Ruipeng Li 2022-06-30 14:23:24 -07:00
parent 750d4877a4
commit aa153c9c89
6 changed files with 18 additions and 9 deletions

View File

@ -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_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 ); 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); HYPRE_BigInt *map, HYPRE_BigInt offset, char *SorA, HYPRE_Complex *y);
/* /*

View File

@ -12,10 +12,12 @@
#include "_hypre_utilities.hpp" #include "_hypre_utilities.hpp"
#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) #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_in,
HYPRE_Int *j_out); 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_Int *i_in,
HYPRE_Real *data_in, HYPRE_Real *Gx_data, HYPRE_Real *Gy_data, HYPRE_Real *Gz_data, HYPRE_Real *data_in, HYPRE_Real *Gx_data, HYPRE_Real *Gy_data, HYPRE_Real *Gz_data,
HYPRE_Real *data_out); HYPRE_Real *data_out);

View File

@ -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_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 ); 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_i,
HYPRE_Int *A_diag_j, HYPRE_Real *A_diag_data, HYPRE_Int *A_offd_i, HYPRE_Int *A_offd_j, 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, HYPRE_Real *A_offd_data, HYPRE_Int *Soc_diag_j, HYPRE_Int *Soc_offd_j, HYPRE_Int *CF_marker,

View File

@ -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 *A_offd_data, HYPRE_Int *CF_marker, HYPRE_Int *dof_func, HYPRE_Int *dof_func_offd,
HYPRE_Complex *row_sums ); 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 color,
HYPRE_Int *pass_order, HYPRE_Int *pass_marker, HYPRE_Int *pass_marker_offd, HYPRE_Int *S_diag_i, 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 *S_diag_j, HYPRE_Int *S_offd_i, HYPRE_Int *S_offd_j, HYPRE_Int *P_diag_i,
HYPRE_Int *P_offd_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 color,
HYPRE_Int *pass_order, HYPRE_Int *pass_marker, HYPRE_Int *pass_marker_offd, 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, 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_Int *Pi_offd_j, HYPRE_Real *Pi_offd_data, HYPRE_Int *P_offd_i, HYPRE_Int *P_offd_j,
HYPRE_Real *P_offd_data ); 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 color,
HYPRE_Int *pass_order, HYPRE_Int *pass_marker, HYPRE_Int *pass_marker_offd, 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, HYPRE_Int *fine_to_coarse, HYPRE_Int *fine_to_coarse_offd, HYPRE_Int *A_diag_i, HYPRE_Int *A_diag_j,

View File

@ -670,7 +670,8 @@ hypre_ParCSRMatrixPrintIJ( const hypre_ParCSRMatrix *matrix,
HYPRE_Int num_nonzeros_offd; HYPRE_Int num_nonzeros_offd;
HYPRE_BigInt ilower, iupper, jlower, jupper; 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) if (!matrix)
{ {

View File

@ -1314,7 +1314,8 @@ template HYPRE_Int hypreDevice_ScatterConstant(HYPRE_Complex *x, HYPRE_Int n, HY
HYPRE_Complex v); HYPRE_Complex v);
__global__ void __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_Complex *x, HYPRE_Complex beta, HYPRE_Complex *y)
{ {
HYPRE_Int i = hypre_gpu_get_grid_thread_id<1, 1>(item); HYPRE_Int i = hypre_gpu_get_grid_thread_id<1, 1>(item);