Skip to content

GPU mixed krylov -- WIP #1423

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
oseikuffuor1 wants to merge 36 commits into
base: master
Choose a base branch
from
Open

GPU mixed krylov -- WIP #1423

Show file tree
Hide file tree
Changes from 13 commits
Commits
Show all changes
36 commits
Select commit Hold shift + click to select a range
31928ae
Merge branch 'master' into gpu-mixed-krylov
oseikuffuor1 Oct 21, 2025
a8835bb
Update autoconf build to prevent use of long double with GPUs.
oseikuffuor1 Oct 23, 2025
57d6ebb
Add utility functions for mixed-precision function development.
oseikuffuor1 Nov 14, 2025
542180e
Add new mixed-precision matrix and vector functionality.
oseikuffuor1 Nov 14, 2025
cd69eba
Updates to support use of hypre_double for dnum_nonzeros in parcsr_ma…
oseikuffuor1 Nov 14, 2025
7781cad
Minor updates to test drivers.
oseikuffuor1 Nov 14, 2025
5ce7060
Fixes for GPU build errors.
oseikuffuor1 Nov 14, 2025
91a169f
Cleanup typos and unnecessary print statements.
oseikuffuor1 Nov 14, 2025
cb030d2
Change utility function name GetSizeOfPrecision to GetSizeOfReal.
oseikuffuor1 Nov 14, 2025
8bfdbb9
Generate configure script with correct autoconf version.
oseikuffuor1 Nov 15, 2025
887c4c3
Cleanup unused variable warnings.
oseikuffuor1 Nov 15, 2025
dfea312
Some fixes and cleanup and minor refactoring for GPU build.
oseikuffuor1 Nov 18, 2025
18e458a
Add gpu function lists for utilities folder.
oseikuffuor1 Nov 18, 2025
7c44083
Added GPU-specific function lists for remaining directories.
oseikuffuor1 Nov 19, 2025
7163b69
Revert changes to MPI function location and remove unnecessary guards.
oseikuffuor1 Nov 19, 2025
1c74ee1
Updating gen_code.sh script to build code for gpu functions
rfalgout Nov 20, 2025
40f8a3a
Merge branch 'gpu-mixed-krylov' of github.com:hypre-space/hypre into …
rfalgout Nov 20, 2025
ac3e7f3
Fixed astyle changes in previous commits
oseikuffuor1 Nov 20, 2025
fd1e9e0
Merge branch 'gpu-mixed-krylov' of github.com:hypre-space/hypre into …
rfalgout Nov 20, 2025
366e620
Minor changes to config/gen_code.sh script
rfalgout Nov 20, 2025
05baa66
Removed empty gpu function lists
oseikuffuor1 Nov 21, 2025
489e1fa
Minor edit to suppress warnings for directories with no GPU function …
oseikuffuor1 Nov 21, 2025
3264a49
Minor changes to gen_code.sh to make it easier to follow
rfalgout Dec 3, 2025
1fecea4
Cosmetic changes to gen_code.sh
rfalgout Dec 3, 2025
f052718
Changed mup.*.gpu to mup.*_gpu in utilities and modified the gen_code…
rfalgout Dec 3, 2025
34762e3
Renaming mup.*.gpu to mup.*_gpu
rfalgout Dec 4, 2025
5bc2ddb
Modified gen_code.sh script to write 'mup_*_gpu.c' files
rfalgout Dec 4, 2025
3996555
First draft of gen_code.sh to put the gpu functions in separate files
rfalgout Dec 4, 2025
e20afec
Updating gen_code.sh to fix a couple of issues
rfalgout Dec 5, 2025
3fd5776
Added gpu if test to gen_code.sh script before writing '.hpp' file
rfalgout Dec 5, 2025
69c34e8
Minor edit to include internal header in wrapper function file.
oseikuffuor1 Dec 9, 2025
d4fe2a5
Changes to unify some vendor-specific return types.
oseikuffuor1 Dec 9, 2025
a3c320e
Generated new MuP wrapper code and headers for GPU support and some r…
oseikuffuor1 Dec 9, 2025
49f8e70
Fixed linker issues and updated gen_code.sh script.
oseikuffuor1 Dec 15, 2025
905f1b7
Some minor re-org. of gen_code.sh script and gix for test driver.
oseikuffuor1 Dec 16, 2025
feb75f2
Merge branch 'master' into gpu-mixed-krylov
oseikuffuor1 Dec 16, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
21 changes: 17 additions & 4 deletions src/config/configure.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -347,10 +347,6 @@ AS_HELP_STRING([--enable-longdouble],
esac],
[hypre_using_longdouble=no]
)
if test "$hypre_using_longdouble" = "yes"
then
AC_DEFINE(HYPRE_LONG_DOUBLE, 1, [Define to 1 if using quad precision values for HYPRE_Real])
fi

AC_ARG_ENABLE(complex,
AS_HELP_STRING([--enable-complex],
Expand Down Expand Up @@ -2424,6 +2420,23 @@ AS_IF([ test x"$hypre_using_hip" == x"yes" ],
[AC_MSG_ERROR([unable to find ${HYPRE_ROCM_PREFIX}/include/hip/hip_common.h ... Ensure ROCm is installed and set ROCM_PATH environment variable to ROCm installation path.])] )],
[])

dnl *********************************************************************
dnl * Check for longdouble support
dnl *********************************************************************

if test "$hypre_using_longdouble" = "yes"
then
if [test "x$hypre_using_cuda" = "xyes" || test "x$hypre_using_device_openmp" = "xyes" || test "x$hypre_using_hip" = "xyes" || test "x$hypre_using_sycl" = "xyes"]
then
AC_MSG_ERROR([******************** Incompatible precision on device **********************
Long double data format is not supported on device.
For GPU builds, please use the default double precision or --enable-single.
****************************************************************************])
else
AC_DEFINE(HYPRE_LONG_DOUBLE, 1, [Define to 1 if using quad precision values for HYPRE_Real])
fi
fi

dnl *********************************************************************
dnl * Set GPU warp size
dnl *********************************************************************
Expand Down
22 changes: 16 additions & 6 deletions src/configure
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3608,12 +3608,6 @@ else $as_nop

fi

if test "$hypre_using_longdouble" = "yes"
then

printf "%s\n" "#define HYPRE_LONG_DOUBLE 1" >>confdefs.h

fi

# Check whether --enable-complex was given.
if test ${enable_complex+y}
Expand Down Expand Up @@ -10651,6 +10645,22 @@ fi
fi


if test "$hypre_using_longdouble" = "yes"
then
if test "x$hypre_using_cuda" = "xyes" || test "x$hypre_using_device_openmp" = "xyes" || test "x$hypre_using_hip" = "xyes" || test "x$hypre_using_sycl" = "xyes"
then
as_fn_error $? "******************** Incompatible precision on device **********************
Long double data format is not supported on device.
For GPU builds, please use the default double precision or --enable-single.
****************************************************************************" "$LINENO" 5
else

printf "%s\n" "#define HYPRE_LONG_DOUBLE 1" >>confdefs.h

fi
fi


if test "x$hypre_warp_size" = "xauto"
then
if test "x$hypre_using_cuda" = "xyes" || test "x$hypre_using_sycl" = "xyes"
Expand Down
2 changes: 1 addition & 1 deletion src/parcsr_block_mv/_hypre_parcsr_block_mv.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -243,7 +243,7 @@ typedef struct
HYPRE_Int owns_data;

HYPRE_BigInt num_nonzeros;
HYPRE_Real d_num_nonzeros;
hypre_double d_num_nonzeros;

/* Buffers used by GetRow to hold row currently being accessed. AJC, 4/99 */
HYPRE_Int *rowindices;
Expand Down
2 changes: 1 addition & 1 deletion src/parcsr_block_mv/par_csr_block_matrix.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -200,7 +200,7 @@ hypre_ParCSRBlockMatrixSetDNumNonzeros( hypre_ParCSRBlockMatrix *matrix)
local_num_nonzeros = (HYPRE_Real) diag_i[local_num_rows] + (HYPRE_Real) offd_i[local_num_rows];
hypre_MPI_Allreduce(&local_num_nonzeros, &total_num_nonzeros, 1,
HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRBlockMatrixDNumNonzeros(matrix) = total_num_nonzeros;
hypre_ParCSRBlockMatrixDNumNonzeros(matrix) = (hypre_double)total_num_nonzeros;

return ierr;
}
Expand Down
2 changes: 1 addition & 1 deletion src/parcsr_block_mv/par_csr_block_matrix.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,7 +59,7 @@ typedef struct
HYPRE_Int owns_data;

HYPRE_BigInt num_nonzeros;
HYPRE_Real d_num_nonzeros;
hypre_double d_num_nonzeros;

/* Buffers used by GetRow to hold row currently being accessed. AJC, 4/99 */
HYPRE_Int *rowindices;
Expand Down
4 changes: 2 additions & 2 deletions src/parcsr_ls/par_2s_interp_device.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -313,7 +313,7 @@ hypre_BoomerAMGBuildModPartialExtInterpDevice( hypre_ParCSRMatrix *A,

hypre_ParCSRMatrixNumNonzeros(P) = hypre_ParCSRMatrixNumNonzeros(W) +
hypre_ParCSRMatrixGlobalNumCols(W);
hypre_ParCSRMatrixDNumNonzeros(P) = (HYPRE_Real) hypre_ParCSRMatrixNumNonzeros(P);
hypre_ParCSRMatrixDNumNonzeros(P) = (hypre_double) hypre_ParCSRMatrixNumNonzeros(P);

hypre_ParCSRMatrixDestroy(W);

Expand Down Expand Up @@ -666,7 +666,7 @@ hypre_BoomerAMGBuildModPartialExtPEInterpDevice( hypre_ParCSRMatrix *A,

hypre_ParCSRMatrixNumNonzeros(P) = hypre_ParCSRMatrixNumNonzeros(W) +
hypre_ParCSRMatrixGlobalNumCols(W);
hypre_ParCSRMatrixDNumNonzeros(P) = (HYPRE_Real) hypre_ParCSRMatrixNumNonzeros(P);
hypre_ParCSRMatrixDNumNonzeros(P) = (hypre_double) hypre_ParCSRMatrixNumNonzeros(P);

hypre_ParCSRMatrixDestroy(W);

Expand Down
6 changes: 3 additions & 3 deletions src/parcsr_ls/par_amg_setup.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3806,12 +3806,12 @@ hypre_BoomerAMGSetup( void *amg_vdata,

if (cum_nnz_AP > 0.0)
{
cum_nnz_AP = hypre_ParCSRMatrixDNumNonzeros(A_array[0]);
cum_nnz_AP = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(A_array[0]);
for (j = 0; j < num_levels - 1; j++)
{
hypre_ParCSRMatrixSetDNumNonzeros(P_array[j]);
cum_nnz_AP += hypre_ParCSRMatrixDNumNonzeros(P_array[j]);
cum_nnz_AP += hypre_ParCSRMatrixDNumNonzeros(A_array[j + 1]);
cum_nnz_AP += (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(P_array[j]);
cum_nnz_AP += (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(A_array[j + 1]);
}
hypre_ParAMGDataCumNnzAP(amg_data) = cum_nnz_AP;
}
Expand Down
2 changes: 1 addition & 1 deletion src/parcsr_ls/par_amg_solve.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -359,7 +359,7 @@ hypre_BoomerAMGSolve( void *amg_vdata,
{
num_coeffs = hypre_CTAlloc(HYPRE_Real, num_levels, HYPRE_MEMORY_HOST);
num_variables = hypre_CTAlloc(HYPRE_Real, num_levels, HYPRE_MEMORY_HOST);
num_coeffs[0] = hypre_ParCSRMatrixDNumNonzeros(A);
num_coeffs[0] = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(A);
num_variables[0] = (HYPRE_Real) hypre_ParCSRMatrixGlobalNumRows(A);

if (block_mode)
Expand Down
8 changes: 4 additions & 4 deletions src/parcsr_ls/par_amg_solveT.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -92,7 +92,7 @@ hypre_BoomerAMGSolveT( void *amg_vdata,

num_coeffs = hypre_CTAlloc(HYPRE_Real, num_levels, HYPRE_MEMORY_HOST);
num_variables = hypre_CTAlloc(HYPRE_BigInt, num_levels, HYPRE_MEMORY_HOST);
num_coeffs[0] = hypre_ParCSRMatrixDNumNonzeros(A_array[0]);
num_coeffs[0] = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(A_array[0]);
num_variables[0] = hypre_ParCSRMatrixGlobalNumRows(A_array[0]);

A_array[0] = A;
Expand All @@ -110,7 +110,7 @@ hypre_BoomerAMGSolveT( void *amg_vdata,
Vtemp = hypre_ParAMGDataVtemp(amg_data);
for (j = 1; j < num_levels; j++)
{
num_coeffs[j] = hypre_ParCSRMatrixDNumNonzeros(A_array[j]);
num_coeffs[j] = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(A_array[j]);
num_variables[j] = hypre_ParCSRMatrixGlobalNumRows(A_array[j]);
}

Expand Down Expand Up @@ -396,11 +396,11 @@ hypre_BoomerAMGCycleT( void *amg_vdata,
if (grid_relax_points) { old_version = 1; }

num_coeffs = hypre_CTAlloc(HYPRE_Real, num_levels, HYPRE_MEMORY_HOST);
num_coeffs[0] = hypre_ParCSRMatrixDNumNonzeros(A_array[0]);
num_coeffs[0] = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(A_array[0]);

for (j = 1; j < num_levels; j++)
{
num_coeffs[j] = hypre_ParCSRMatrixDNumNonzeros(A_array[j]);
num_coeffs[j] = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(A_array[j]);
}

/*---------------------------------------------------------------------
Expand Down
4 changes: 2 additions & 2 deletions src/parcsr_ls/par_cycle.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -169,7 +169,7 @@ hypre_BoomerAMGCycle( void *amg_vdata,
}

num_coeffs = hypre_CTAlloc(HYPRE_Real, num_levels, HYPRE_MEMORY_HOST);
num_coeffs[0] = hypre_ParCSRMatrixDNumNonzeros(A_array[0]);
num_coeffs[0] = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(A_array[0]);
comm = hypre_ParCSRMatrixComm(A_array[0]);
hypre_MPI_Comm_rank(comm, &my_id);

Expand All @@ -184,7 +184,7 @@ hypre_BoomerAMGCycle( void *amg_vdata,
{
for (j = 1; j < num_levels; j++)
{
num_coeffs[j] = hypre_ParCSRMatrixDNumNonzeros(A_array[j]);
num_coeffs[j] = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(A_array[j]);
}
}

Expand Down
4 changes: 2 additions & 2 deletions src/parcsr_ls/par_fsai_setup.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1080,8 +1080,8 @@ hypre_FSAIPrintStats( void *fsai_vdata,
/* Compute density */
hypre_ParCSRMatrixSetDNumNonzeros(G);
hypre_ParCSRMatrixSetDNumNonzeros(A);
density = hypre_ParCSRMatrixDNumNonzeros(G) /
hypre_ParCSRMatrixDNumNonzeros(A);
density = (HYPRE_Real)(hypre_ParCSRMatrixDNumNonzeros(G) /
hypre_ParCSRMatrixDNumNonzeros(A));
hypre_ParFSAIDataDensity(fsai_data) = density;

if (!my_id)
Expand Down
44 changes: 22 additions & 22 deletions src/parcsr_ls/par_ilu_setup.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1104,11 +1104,11 @@ hypre_ILUSetup( void *ilu_vdata,
if (matS)
{
hypre_ParCSRMatrixSetDNumNonzeros(matS);
nnzS = hypre_ParCSRMatrixDNumNonzeros(matS);
nnzS = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(matS);
/* if we have Schur system need to reduce it from size_C */
}
hypre_ParILUDataOperatorComplexity(ilu_data) = ((HYPRE_Real)nnzG + nnzS) /
hypre_ParCSRMatrixDNumNonzeros(matA);
(HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(matA);
}
else if (ilu_type == 50)
{
Expand All @@ -1132,19 +1132,19 @@ hypre_ILUSetup( void *ilu_vdata,
if (matS)
{
hypre_ParCSRMatrixSetDNumNonzeros(matS);
nnzS = hypre_ParCSRMatrixDNumNonzeros(matS);
nnzS = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(matS);
/* if we have Schur system need to reduce it from size_C */
}
hypre_ParILUDataOperatorComplexity(ilu_data) = ((HYPRE_Real)nnzG + nnzS) /
hypre_ParCSRMatrixDNumNonzeros(matA);
(HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(matA);
}
else
#endif
{
if (matS)
{
hypre_ParCSRMatrixSetDNumNonzeros(matS);
nnzS = hypre_ParCSRMatrixDNumNonzeros(matS);
nnzS = (HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(matS);

/* If we have Schur system need to reduce it from size_C */
size_C -= hypre_ParCSRMatrixGlobalNumRows(matS);
Expand Down Expand Up @@ -1172,9 +1172,9 @@ hypre_ILUSetup( void *ilu_vdata,
}

hypre_ParILUDataOperatorComplexity(ilu_data) = ((HYPRE_Real)size_C + nnzS +
hypre_ParCSRMatrixDNumNonzeros(matL) +
hypre_ParCSRMatrixDNumNonzeros(matU)) /
hypre_ParCSRMatrixDNumNonzeros(matA);
(HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(matL) +
(HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(matU)) /
(HYPRE_Real)hypre_ParCSRMatrixDNumNonzeros(matA);
}

/* TODO (VPM): Move ILU statistics printout to its own function */
Expand Down Expand Up @@ -2800,7 +2800,7 @@ hypre_ILUSetupMILU0(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) ctrL;
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matL) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matL) = (hypre_double)total_nnz;

matU = hypre_ParCSRMatrixCreate( comm,
hypre_ParCSRMatrixGlobalNumRows(A),
Expand All @@ -2827,7 +2827,7 @@ hypre_ILUSetupMILU0(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) ctrU;
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matU) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matU) = (hypre_double)total_nnz;
/* free memory */
hypre_TFree(wL, HYPRE_MEMORY_HOST);
hypre_TFree(iw, HYPRE_MEMORY_HOST);
Expand Down Expand Up @@ -3807,7 +3807,7 @@ hypre_ILUSetupILUK(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) (L_diag_i[n]);
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matL) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matL) = (hypre_double)total_nnz;

matU = hypre_ParCSRMatrixCreate( comm,
hypre_ParCSRMatrixGlobalNumRows(A),
Expand All @@ -3833,7 +3833,7 @@ hypre_ILUSetupILUK(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) (U_diag_i[n]);
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matU) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matU) = (hypre_double)total_nnz;

/* free */
hypre_TFree(iw, HYPRE_MEMORY_HOST);
Expand Down Expand Up @@ -4686,7 +4686,7 @@ hypre_ILUSetupILUT(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) (L_diag_i[n]);
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matL) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matL) = (hypre_double)total_nnz;

matU = hypre_ParCSRMatrixCreate( comm,
hypre_ParCSRMatrixGlobalNumRows(A),
Expand Down Expand Up @@ -4715,7 +4715,7 @@ hypre_ILUSetupILUT(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) (U_diag_i[n]);
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matU) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matU) = (hypre_double)total_nnz;

/* free working array */
hypre_TFree(iw, HYPRE_MEMORY_HOST);
Expand Down Expand Up @@ -4834,8 +4834,8 @@ hypre_NSHSetup( void *nsh_vdata,
hypre_ParCSRMatrixSetDNumNonzeros(matM);

/* Compute complexity */
hypre_ParNSHDataOperatorComplexity(nsh_data) = hypre_ParCSRMatrixDNumNonzeros(matM) /
hypre_ParCSRMatrixDNumNonzeros(matA);
hypre_ParNSHDataOperatorComplexity(nsh_data) = (HYPRE_Real)(hypre_ParCSRMatrixDNumNonzeros(matM) /
hypre_ParCSRMatrixDNumNonzeros(matA));
if (my_id == 0 && print_level > 0)
{
hypre_printf("NSH SETUP: operator complexity = %f \n",
Expand Down Expand Up @@ -5475,7 +5475,7 @@ hypre_ILUSetupILU0RAS(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) ctrL;
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matL) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matL) = (hypre_double)total_nnz;

matU = hypre_ParCSRMatrixCreate( comm,
global_num_rows,
Expand All @@ -5502,7 +5502,7 @@ hypre_ILUSetupILU0RAS(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) ctrU;
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matU) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matU) = (hypre_double)total_nnz;
/* free memory */
hypre_TFree(wL, HYPRE_MEMORY_HOST);
hypre_TFree(iw, HYPRE_MEMORY_HOST);
Expand Down Expand Up @@ -6594,7 +6594,7 @@ hypre_ILUSetupILUKRAS(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) (L_diag_i[total_rows]);
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matL) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matL) = (hypre_double)total_nnz;

matU = hypre_ParCSRMatrixCreate( comm,
global_num_rows,
Expand All @@ -6620,7 +6620,7 @@ hypre_ILUSetupILUKRAS(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) (U_diag_i[total_rows]);
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matU) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matU) = (hypre_double)total_nnz;

/* free */
hypre_TFree(iw, HYPRE_MEMORY_HOST);
Expand Down Expand Up @@ -7539,7 +7539,7 @@ hypre_ILUSetupILUTRAS(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) (L_diag_i[total_rows]);
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matL) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matL) = (hypre_double)total_nnz;

matU = hypre_ParCSRMatrixCreate( comm,
global_num_rows,
Expand All @@ -7566,7 +7566,7 @@ hypre_ILUSetupILUTRAS(hypre_ParCSRMatrix *A,
/* store (global) total number of nonzeros */
local_nnz = (HYPRE_Real) (U_diag_i[total_rows]);
hypre_MPI_Allreduce(&local_nnz, &total_nnz, 1, HYPRE_MPI_REAL, hypre_MPI_SUM, comm);
hypre_ParCSRMatrixDNumNonzeros(matU) = total_nnz;
hypre_ParCSRMatrixDNumNonzeros(matU) = (hypre_double)total_nnz;

/* free working array */
hypre_TFree(iw, HYPRE_MEMORY_HOST);
Expand Down
6 changes: 3 additions & 3 deletions src/parcsr_ls/par_lr_interp_device.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -970,7 +970,7 @@ hypre_BoomerAMGBuildExtInterpDevice(hypre_ParCSRMatrix *A,

hypre_ParCSRMatrixNumNonzeros(P) = hypre_ParCSRMatrixNumNonzeros(W) +
hypre_ParCSRMatrixGlobalNumCols(W);
hypre_ParCSRMatrixDNumNonzeros(P) = (HYPRE_Real) hypre_ParCSRMatrixNumNonzeros(P);
hypre_ParCSRMatrixDNumNonzeros(P) = (hypre_double) hypre_ParCSRMatrixNumNonzeros(P);

hypre_GpuProfilingPushRange("Truncation");
if (trunc_factor != 0.0 || max_elmts > 0)
Expand Down Expand Up @@ -1248,7 +1248,7 @@ hypre_BoomerAMGBuildExtPIInterpDevice( hypre_ParCSRMatrix *A,

hypre_ParCSRMatrixNumNonzeros(P) = hypre_ParCSRMatrixNumNonzeros(W) +
hypre_ParCSRMatrixGlobalNumCols(W);
hypre_ParCSRMatrixDNumNonzeros(P) = (HYPRE_Real) hypre_ParCSRMatrixNumNonzeros(P);
hypre_ParCSRMatrixDNumNonzeros(P) = (hypre_double) hypre_ParCSRMatrixNumNonzeros(P);

hypre_GpuProfilingPushRange("Truncation");
if (trunc_factor != 0.0 || max_elmts > 0)
Expand Down Expand Up @@ -1525,7 +1525,7 @@ hypre_BoomerAMGBuildExtPEInterpDevice(hypre_ParCSRMatrix *A,

hypre_ParCSRMatrixNumNonzeros(P) = hypre_ParCSRMatrixNumNonzeros(W) +
hypre_ParCSRMatrixGlobalNumCols(W);
hypre_ParCSRMatrixDNumNonzeros(P) = (HYPRE_Real) hypre_ParCSRMatrixNumNonzeros(P);
hypre_ParCSRMatrixDNumNonzeros(P) = (hypre_double) hypre_ParCSRMatrixNumNonzeros(P);

hypre_GpuProfilingPushRange("Truncation");
if (trunc_factor != 0.0 || max_elmts > 0)
Expand Down
Loading