diff --git a/src/parcsr_ls/par_mgr.c b/src/parcsr_ls/par_mgr.c index b5b20fb7a..718c022fb 100644 --- a/src/parcsr_ls/par_mgr.c +++ b/src/parcsr_ls/par_mgr.c @@ -1281,7 +1281,7 @@ hypre_MGRBuildPHost( hypre_ParCSRMatrix *A, if (method > 0) { - hypre_ParCSRMatrixGenerateFFFC(A, CF_marker, num_cpts_global, NULL, &A_FC, &A_FF); + hypre_ParCSRMatrixGenerateFFFCHost(A, CF_marker, num_cpts_global, NULL, &A_FC, &A_FF); diag = hypre_CTAlloc(HYPRE_Complex, nfpoints, memory_location_P); if (method == 1) { diff --git a/src/parcsr_ls/par_mod_lr_interp.c b/src/parcsr_ls/par_mod_lr_interp.c index fc222bc43..5b371b74e 100644 --- a/src/parcsr_ls/par_mod_lr_interp.c +++ b/src/parcsr_ls/par_mod_lr_interp.c @@ -122,7 +122,7 @@ hypre_BoomerAMGBuildModExtInterpHost(hypre_ParCSRMatrix *A, hypre_MPI_Bcast(&total_global_cpts, 1, HYPRE_MPI_BIG_INT, num_procs - 1, comm); n_Cpts = num_cpts_global[1] - num_cpts_global[0]; - hypre_ParCSRMatrixGenerateFFFC(A, CF_marker, num_cpts_global, S, &As_FC, &As_FF); + hypre_ParCSRMatrixGenerateFFFCHost(A, CF_marker, num_cpts_global, S, &As_FC, &As_FF); As_FC_diag = hypre_ParCSRMatrixDiag(As_FC); As_FC_diag_i = hypre_CSRMatrixI(As_FC_diag); @@ -706,7 +706,7 @@ hypre_BoomerAMGBuildModExtPIInterpHost(hypre_ParCSRMatrix *A, hypre_MPI_Bcast(&total_global_cpts, 1, HYPRE_MPI_BIG_INT, num_procs - 1, comm); n_Cpts = num_cpts_global[1] - num_cpts_global[0]; - hypre_ParCSRMatrixGenerateFFFC(A, CF_marker, num_cpts_global, S, &As_FC, &As_FF); + hypre_ParCSRMatrixGenerateFFFCHost(A, CF_marker, num_cpts_global, S, &As_FC, &As_FF); if (num_procs > 1) { @@ -1371,7 +1371,7 @@ hypre_BoomerAMGBuildModExtPEInterpHost(hypre_ParCSRMatrix *A, hypre_MPI_Bcast(&total_global_cpts, 1, HYPRE_MPI_BIG_INT, num_procs - 1, comm); n_Cpts = num_cpts_global[1] - num_cpts_global[0]; - hypre_ParCSRMatrixGenerateFFFC(A, CF_marker, num_cpts_global, S, &As_FC, &As_FF); + hypre_ParCSRMatrixGenerateFFFCHost(A, CF_marker, num_cpts_global, S, &As_FC, &As_FF); As_FC_diag = hypre_ParCSRMatrixDiag(As_FC); As_FC_diag_i = hypre_CSRMatrixI(As_FC_diag); diff --git a/src/parcsr_mv/_hypre_parcsr_mv.h b/src/parcsr_mv/_hypre_parcsr_mv.h index 0cf680a61..be6fde904 100644 --- a/src/parcsr_mv/_hypre_parcsr_mv.h +++ b/src/parcsr_mv/_hypre_parcsr_mv.h @@ -742,10 +742,15 @@ HYPRE_Int HYPRE_VectorToParVector ( MPI_Comm comm, HYPRE_Vector b, HYPRE_BigInt HYPRE_Int HYPRE_ParVectorGetValues ( HYPRE_ParVector vector, HYPRE_Int num_values, HYPRE_BigInt *indices, HYPRE_Complex *values); -/*gen_fffc.c */ -HYPRE_Int hypre_ParCSRMatrixGenerateFFFC(hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, - HYPRE_BigInt *cpts_starts, hypre_ParCSRMatrix *S, hypre_ParCSRMatrix **A_FC_ptr, - hypre_ParCSRMatrix **A_FF_ptr ) ; +/* gen_fffc.c */ +HYPRE_Int hypre_ParCSRMatrixGenerateFFFCHost( hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, + HYPRE_BigInt *cpts_starts, hypre_ParCSRMatrix *S, + hypre_ParCSRMatrix **A_FC_ptr, + hypre_ParCSRMatrix **A_FF_ptr ) ; +HYPRE_Int hypre_ParCSRMatrixGenerateFFFC( hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, + HYPRE_BigInt *cpts_starts, hypre_ParCSRMatrix *S, + hypre_ParCSRMatrix **A_FC_ptr, + hypre_ParCSRMatrix **A_FF_ptr ) ; HYPRE_Int hypre_ParCSRMatrixGenerateFFFC3(hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, HYPRE_BigInt *cpts_starts, hypre_ParCSRMatrix *S, hypre_ParCSRMatrix **A_FC_ptr, hypre_ParCSRMatrix **A_FF_ptr ) ; @@ -1216,7 +1221,6 @@ HYPRE_Int hypre_ParVectorGetValuesDevice(hypre_ParVector *vector, HYPRE_Int num_ HYPRE_BigInt *indices, HYPRE_BigInt base, HYPRE_Complex *values); - #ifdef __cplusplus } #endif diff --git a/src/parcsr_mv/gen_fffc.c b/src/parcsr_mv/gen_fffc.c index 28c793575..a39e8009f 100644 --- a/src/parcsr_mv/gen_fffc.c +++ b/src/parcsr_mv/gen_fffc.c @@ -10,17 +10,17 @@ #include "_hypre_lapack.h" #include "_hypre_blas.h" -/* ----------------------------------------------------------------------------- - * generate AFF or AFC - * ----------------------------------------------------------------------------- */ +/*-------------------------------------------------------------------------- + * hypre_ParCSRMatrixGenerateFFFCHost + *--------------------------------------------------------------------------*/ HYPRE_Int -hypre_ParCSRMatrixGenerateFFFC( hypre_ParCSRMatrix *A, - HYPRE_Int *CF_marker, - HYPRE_BigInt *cpts_starts, - hypre_ParCSRMatrix *S, - hypre_ParCSRMatrix **A_FC_ptr, - hypre_ParCSRMatrix **A_FF_ptr) +hypre_ParCSRMatrixGenerateFFFCHost( hypre_ParCSRMatrix *A, + HYPRE_Int *CF_marker, + HYPRE_BigInt *cpts_starts, + hypre_ParCSRMatrix *S, + hypre_ParCSRMatrix **A_FC_ptr, + hypre_ParCSRMatrix **A_FF_ptr) { MPI_Comm comm = hypre_ParCSRMatrixComm(A); HYPRE_MemoryLocation memory_location_P = hypre_ParCSRMatrixMemoryLocation(A); @@ -487,10 +487,41 @@ hypre_ParCSRMatrixGenerateFFFC( hypre_ParCSRMatrix *A, return hypre_error_flag; } +/*-------------------------------------------------------------------------- + * hypre_ParCSRMatrixGenerateFFFC + * + * Generate AFF or AFC + *--------------------------------------------------------------------------*/ -/* ----------------------------------------------------------------------------- +HYPRE_Int +hypre_ParCSRMatrixGenerateFFFC( hypre_ParCSRMatrix *A, + HYPRE_Int *CF_marker, + HYPRE_BigInt *cpts_starts, + hypre_ParCSRMatrix *S, + hypre_ParCSRMatrix **A_FC_ptr, + hypre_ParCSRMatrix **A_FF_ptr) +{ +#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) || defined(HYPRE_USING_SYCL) + HYPRE_ExecutionPolicy exec = hypre_GetExecPolicy1( hypre_ParCSRMatrixMemoryLocation(A) ); + + if (exec == HYPRE_EXEC_DEVICE) + { + hypre_ParCSRMatrixGenerateFFFCDevice(A, CF_marker, cpts_starts, S, A_FC_ptr, A_FF_ptr); + } + else +#endif + { + hypre_ParCSRMatrixGenerateFFFCHost(A, CF_marker, cpts_starts, S, A_FC_ptr, A_FF_ptr); + } + + return hypre_error_flag; +} + +/*-------------------------------------------------------------------------- + * hypre_ParCSRMatrixGenerateFFFC3 + * * generate AFF, AFC, for 2 stage extended interpolation - * ----------------------------------------------------------------------------- */ + *--------------------------------------------------------------------------*/ HYPRE_Int hypre_ParCSRMatrixGenerateFFFC3( hypre_ParCSRMatrix *A, @@ -510,6 +541,7 @@ hypre_ParCSRMatrixGenerateFFFC3( hypre_ParCSRMatrix *A, HYPRE_Complex *A_diag_data = hypre_CSRMatrixData(A_diag); HYPRE_Int *A_diag_i = hypre_CSRMatrixI(A_diag); HYPRE_Int *A_diag_j = hypre_CSRMatrixJ(A_diag); + /* off-diag part of A */ hypre_CSRMatrix *A_offd = hypre_ParCSRMatrixOffd(A); HYPRE_Complex *A_offd_data = hypre_CSRMatrixData(A_offd); @@ -523,6 +555,7 @@ hypre_ParCSRMatrixGenerateFFFC3( hypre_ParCSRMatrix *A, hypre_CSRMatrix *S_diag = hypre_ParCSRMatrixDiag(S); HYPRE_Int *S_diag_i = hypre_CSRMatrixI(S_diag); HYPRE_Int *S_diag_j = hypre_CSRMatrixJ(S_diag); + /* off-diag part of S */ hypre_CSRMatrix *S_offd = hypre_ParCSRMatrixOffd(S); HYPRE_Int *S_offd_i = hypre_CSRMatrixI(S_offd); @@ -698,9 +731,11 @@ hypre_ParCSRMatrixGenerateFFFC3( hypre_ParCSRMatrix *A, } index = 0; num_sends = hypre_ParCSRCommPkgNumSends(comm_pkg); - int_buf_data = hypre_CTAlloc(HYPRE_Int, hypre_ParCSRCommPkgSendMapStart(comm_pkg, num_sends), + int_buf_data = hypre_CTAlloc(HYPRE_Int, + hypre_ParCSRCommPkgSendMapStart(comm_pkg, num_sends), HYPRE_MEMORY_HOST); - big_buf_data = hypre_CTAlloc(HYPRE_BigInt, hypre_ParCSRCommPkgSendMapStart(comm_pkg, num_sends), + big_buf_data = hypre_CTAlloc(HYPRE_BigInt, + hypre_ParCSRCommPkgSendMapStart(comm_pkg, num_sends), HYPRE_MEMORY_HOST); for (i = 0; i < num_sends; i++) { @@ -716,7 +751,7 @@ hypre_ParCSRMatrixGenerateFFFC3( hypre_ParCSRMatrix *A, hypre_ParCSRCommHandleDestroy(comm_handle); - comm_handle = hypre_ParCSRCommHandleCreate( 21, comm_pkg, big_buf_data, big_convert_offd); + comm_handle = hypre_ParCSRCommHandleCreate(21, comm_pkg, big_buf_data, big_convert_offd); hypre_ParCSRCommHandleDestroy(comm_handle); @@ -1021,9 +1056,12 @@ hypre_ParCSRMatrixGenerateFFFC3( hypre_ParCSRMatrix *A, return hypre_error_flag; } -/* ----------------------------------------------------------------------------- - * generate AFF, AFC, AFFC for 2 stage extended+i(e)interpolation - * ----------------------------------------------------------------------------- */ + +/*-------------------------------------------------------------------------- + * hypre_ParCSRMatrixGenerateFFFCD3 + * + * Generate AFF, AFC, AFFC for 2 stage extended+i(e)interpolation + *--------------------------------------------------------------------------*/ HYPRE_Int hypre_ParCSRMatrixGenerateFFFCD3( hypre_ParCSRMatrix *A, @@ -1190,7 +1228,8 @@ hypre_ParCSRMatrixGenerateFFFCD3( hypre_ParCSRMatrix *A, big_new_Fpts = n_new_Fpts; hypre_MPI_Scan(&big_Fpts, fpts_starts + 1, 1, HYPRE_MPI_BIG_INT, hypre_MPI_SUM, comm); - hypre_MPI_Scan(&big_new_Fpts, new_fpts_starts + 1, 1, HYPRE_MPI_BIG_INT, hypre_MPI_SUM, comm); + hypre_MPI_Scan(&big_new_Fpts, new_fpts_starts + 1, 1, HYPRE_MPI_BIG_INT, + hypre_MPI_SUM, comm); fpts_starts[0] = fpts_starts[1] - big_Fpts; new_fpts_starts[0] = new_fpts_starts[1] - big_new_Fpts; if (my_id == num_procs - 1) @@ -1233,9 +1272,11 @@ hypre_ParCSRMatrixGenerateFFFCD3( hypre_ParCSRMatrix *A, } index = 0; num_sends = hypre_ParCSRCommPkgNumSends(comm_pkg); - int_buf_data = hypre_CTAlloc(HYPRE_Int, hypre_ParCSRCommPkgSendMapStart(comm_pkg, num_sends), + int_buf_data = hypre_CTAlloc(HYPRE_Int, + hypre_ParCSRCommPkgSendMapStart(comm_pkg, num_sends), HYPRE_MEMORY_HOST); - big_buf_data = hypre_CTAlloc(HYPRE_BigInt, hypre_ParCSRCommPkgSendMapStart(comm_pkg, num_sends), + big_buf_data = hypre_CTAlloc(HYPRE_BigInt, + hypre_ParCSRCommPkgSendMapStart(comm_pkg, num_sends), HYPRE_MEMORY_HOST); for (i = 0; i < num_sends; i++) { @@ -1247,11 +1288,11 @@ hypre_ParCSRMatrixGenerateFFFCD3( hypre_ParCSRMatrix *A, } } - comm_handle = hypre_ParCSRCommHandleCreate( 11, comm_pkg, int_buf_data, CF_marker_offd); + comm_handle = hypre_ParCSRCommHandleCreate(11, comm_pkg, int_buf_data, CF_marker_offd); hypre_ParCSRCommHandleDestroy(comm_handle); - comm_handle = hypre_ParCSRCommHandleCreate( 21, comm_pkg, big_buf_data, big_convert_offd); + comm_handle = hypre_ParCSRCommHandleCreate(21, comm_pkg, big_buf_data, big_convert_offd); hypre_ParCSRCommHandleDestroy(comm_handle); diff --git a/src/parcsr_mv/par_csr_fffc_device.c b/src/parcsr_mv/par_csr_fffc_device.c index 87acceac8..ab2be376f 100644 --- a/src/parcsr_mv/par_csr_fffc_device.c +++ b/src/parcsr_mv/par_csr_fffc_device.c @@ -1346,6 +1346,10 @@ hypre_ParCSRMatrixGenerateFFFCDevice_core( hypre_ParCSRMatrix *A, return hypre_error_flag; } +/*-------------------------------------------------------------------------- + * hypre_ParCSRMatrixGenerateFFFCDevice + *--------------------------------------------------------------------------*/ + HYPRE_Int hypre_ParCSRMatrixGenerateFFFCDevice( hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, @@ -1354,10 +1358,15 @@ hypre_ParCSRMatrixGenerateFFFCDevice( hypre_ParCSRMatrix *A, hypre_ParCSRMatrix **AFC_ptr, hypre_ParCSRMatrix **AFF_ptr ) { - return hypre_ParCSRMatrixGenerateFFFCDevice_core(A, CF_marker, cpts_starts, S, AFC_ptr, AFF_ptr, + return hypre_ParCSRMatrixGenerateFFFCDevice_core(A, CF_marker, cpts_starts, S, + AFC_ptr, AFF_ptr, NULL, NULL, 1); } +/*-------------------------------------------------------------------------- + * hypre_ParCSRMatrixGenerateFFFC3Device + *--------------------------------------------------------------------------*/ + HYPRE_Int hypre_ParCSRMatrixGenerateFFFC3Device( hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, @@ -1366,10 +1375,15 @@ hypre_ParCSRMatrixGenerateFFFC3Device( hypre_ParCSRMatrix *A, hypre_ParCSRMatrix **AFC_ptr, hypre_ParCSRMatrix **AFF_ptr) { - return hypre_ParCSRMatrixGenerateFFFCDevice_core(A, CF_marker, cpts_starts, S, AFC_ptr, AFF_ptr, + return hypre_ParCSRMatrixGenerateFFFCDevice_core(A, CF_marker, cpts_starts, S, + AFC_ptr, AFF_ptr, NULL, NULL, 2); } +/*-------------------------------------------------------------------------- + * hypre_ParCSRMatrixGenerateFFCFDevice + *--------------------------------------------------------------------------*/ + HYPRE_Int hypre_ParCSRMatrixGenerateFFCFDevice( hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, @@ -1378,10 +1392,14 @@ hypre_ParCSRMatrixGenerateFFCFDevice( hypre_ParCSRMatrix *A, hypre_ParCSRMatrix **ACF_ptr, hypre_ParCSRMatrix **AFF_ptr ) { - return hypre_ParCSRMatrixGenerateFFFCDevice_core(A, CF_marker, cpts_starts, S, NULL, AFF_ptr, + return hypre_ParCSRMatrixGenerateFFFCDevice_core(A, CF_marker, cpts_starts, S, + NULL, AFF_ptr, ACF_ptr, NULL, 1); } +/*-------------------------------------------------------------------------- + * hypre_ParCSRMatrixGenerateCFDevice + *--------------------------------------------------------------------------*/ HYPRE_Int hypre_ParCSRMatrixGenerateCFDevice( hypre_ParCSRMatrix *A, @@ -1390,10 +1408,15 @@ hypre_ParCSRMatrixGenerateCFDevice( hypre_ParCSRMatrix *A, hypre_ParCSRMatrix *S, hypre_ParCSRMatrix **ACF_ptr) { - return hypre_ParCSRMatrixGenerateFFFCDevice_core(A, CF_marker, cpts_starts, S, NULL, NULL, ACF_ptr, - NULL, 1); + return hypre_ParCSRMatrixGenerateFFFCDevice_core(A, CF_marker, cpts_starts, S, + NULL, NULL, + ACF_ptr, NULL, 1); } +/*-------------------------------------------------------------------------- + * hypre_ParCSRMatrixGenerateCCDevice + *--------------------------------------------------------------------------*/ + HYPRE_Int hypre_ParCSRMatrixGenerateCCDevice( hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, @@ -1401,10 +1424,15 @@ hypre_ParCSRMatrixGenerateCCDevice( hypre_ParCSRMatrix *A, hypre_ParCSRMatrix *S, hypre_ParCSRMatrix **ACC_ptr) { - return hypre_ParCSRMatrixGenerateFFFCDevice_core(A, CF_marker, cpts_starts, S, NULL, NULL, NULL, - ACC_ptr, 1); + return hypre_ParCSRMatrixGenerateFFFCDevice_core(A, CF_marker, cpts_starts, S, + NULL, NULL, + NULL, ACC_ptr, 1); } +/*-------------------------------------------------------------------------- + * hypre_ParCSRMatrixGenerate1DCFDevice + *--------------------------------------------------------------------------*/ + HYPRE_Int hypre_ParCSRMatrixGenerate1DCFDevice( hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, diff --git a/src/parcsr_mv/protos.h b/src/parcsr_mv/protos.h index 72d1b6e11..b4eadafe0 100644 --- a/src/parcsr_mv/protos.h +++ b/src/parcsr_mv/protos.h @@ -92,10 +92,15 @@ HYPRE_Int HYPRE_VectorToParVector ( MPI_Comm comm, HYPRE_Vector b, HYPRE_BigInt HYPRE_Int HYPRE_ParVectorGetValues ( HYPRE_ParVector vector, HYPRE_Int num_values, HYPRE_BigInt *indices, HYPRE_Complex *values); -/*gen_fffc.c */ -HYPRE_Int hypre_ParCSRMatrixGenerateFFFC(hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, - HYPRE_BigInt *cpts_starts, hypre_ParCSRMatrix *S, hypre_ParCSRMatrix **A_FC_ptr, - hypre_ParCSRMatrix **A_FF_ptr ) ; +/* gen_fffc.c */ +HYPRE_Int hypre_ParCSRMatrixGenerateFFFCHost( hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, + HYPRE_BigInt *cpts_starts, hypre_ParCSRMatrix *S, + hypre_ParCSRMatrix **A_FC_ptr, + hypre_ParCSRMatrix **A_FF_ptr ) ; +HYPRE_Int hypre_ParCSRMatrixGenerateFFFC( hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, + HYPRE_BigInt *cpts_starts, hypre_ParCSRMatrix *S, + hypre_ParCSRMatrix **A_FC_ptr, + hypre_ParCSRMatrix **A_FF_ptr ) ; HYPRE_Int hypre_ParCSRMatrixGenerateFFFC3(hypre_ParCSRMatrix *A, HYPRE_Int *CF_marker, HYPRE_BigInt *cpts_starts, hypre_ParCSRMatrix *S, hypre_ParCSRMatrix **A_FC_ptr, hypre_ParCSRMatrix **A_FF_ptr ) ; @@ -565,4 +570,3 @@ HYPRE_Int hypre_ParVectorElmdivpyMarked( hypre_ParVector *x, hypre_ParVector *b, HYPRE_Int hypre_ParVectorGetValuesDevice(hypre_ParVector *vector, HYPRE_Int num_values, HYPRE_BigInt *indices, HYPRE_BigInt base, HYPRE_Complex *values); - diff --git a/src/test/ij_device.c b/src/test/ij_device.c index 9469149e9..0da8a96bf 100644 --- a/src/test/ij_device.c +++ b/src/test/ij_device.c @@ -2300,8 +2300,8 @@ testFFFC(HYPRE_ParCSRMatrix parcsr_A) parcsr_A_h = hypre_ParCSRMatrixClone_v2(parcsr_A, 1, HYPRE_MEMORY_HOST); parcsr_S_h = hypre_ParCSRMatrixClone_v2(parcsr_S_device, 0, HYPRE_MEMORY_HOST); hypre_MatvecCommPkgCreate(parcsr_A_h); - hypre_ParCSRMatrixGenerateFFFC(parcsr_A_h, hypre_IntArrayData(h_CF_marker), coarse_pnts_global, - parcsr_S_h, &AFC_h, &AFF_h); + hypre_ParCSRMatrixGenerateFFFCHost(parcsr_A_h, hypre_IntArrayData(h_CF_marker), + coarse_pnts_global, parcsr_S_h, &AFC_h, &AFF_h); /* AFF * AFC */ W_h = hypre_ParCSRMatMatHost(AFF_h, AFC_h);