CFD/hypre
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Generalize hypre_ParCSRMatrixGenerateFFFC (#772)

Browse Source 
This PR modifies hypre_ParCSRMatrixGenerateFFFC to act as a wrapper between the host and device implementations. Consequently, hypre_ParCSRMatrixGenerateFFFCHost has been added.
This commit is contained in:
Victor A. Paludetto Magri 2022-11-30 09:13:39 -08:00 committed by GitHub
parent b93beb9465
commit dcbd6d724b
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
7 changed files with 122 additions and 45 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/parcsr_ls/par_mgr.c
View File

@ -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)
{

6
src/parcsr_ls/par_mod_lr_interp.c
View File

@ -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);

14
src/parcsr_mv/_hypre_parcsr_mv.h
View File

@ -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

85
src/parcsr_mv/gen_fffc.c
View File

@ -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);

42
src/parcsr_mv/par_csr_fffc_device.c
View File

@ -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,

14
src/parcsr_mv/protos.h
View File

@ -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);

4
src/test/ij_device.c
View File

@ -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);