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

MGR bug fix and no longer throw error for zero len kernel launches for sycl

Browse Source
This commit is contained in:
Wayne Mitchell 2022-03-25 19:41:45 +00:00
parent 6b5a13cedf
commit b5de7f7dfd
5 changed files with 31 additions and 54 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
src/parcsr_ls/par_amg_setup.c
View File

@ -1555,7 +1555,6 @@ hypre_BoomerAMGSetup( void *amg_vdata,
HYPRE_ANNOTATE_REGION_END("%s", "Coarsening"); HYPRE_ANNOTATE_REGION_END("%s", "Coarsening");
HYPRE_ANNOTATE_REGION_BEGIN("%s", "Interpolation"); HYPRE_ANNOTATE_REGION_BEGIN("%s", "Interpolation");
hypre_printf("WM: debug - agg_num_levels = %d, nodal = %d\n", agg_num_levels, nodal);
if (level < agg_num_levels) if (level < agg_num_levels)
{ {
if (nodal == 0) if (nodal == 0)

54
src/parcsr_ls/par_mgr.c
View File

@ -3206,19 +3206,19 @@ hypre_MGRBuildInterp(hypre_ParCSRMatrix *A,
/* Interpolation for each level */ /* Interpolation for each level */
if (interp_type < 3) if (interp_type < 3)
{ {
if (exec == HYPRE_EXEC_HOST) #if defined(HYPRE_USING_CUDA)
if (exec == HYPRE_EXEC_DEVICE)
{
hypre_MGRBuildPDevice(A, CF_marker, num_cpts_global, interp_type, &P_ptr);
//hypre_ParCSRMatrixPrintIJ(P_ptr, 0, 0, "P_device");
}
else
#endif
{ {
// hypre_MGRBuildP(A, CF_marker, num_cpts_global, interp_type, debug_flag, &P_ptr); // hypre_MGRBuildP(A, CF_marker, num_cpts_global, interp_type, debug_flag, &P_ptr);
hypre_MGRBuildPHost(A, CF_marker, num_cpts_global, interp_type, &P_ptr); hypre_MGRBuildPHost(A, CF_marker, num_cpts_global, interp_type, &P_ptr);
//hypre_ParCSRMatrixPrintIJ(P_ptr, 0, 0, "P_host"); //hypre_ParCSRMatrixPrintIJ(P_ptr, 0, 0, "P_host");
} }
#if defined(HYPRE_USING_CUDA)
else
{
hypre_MGRBuildPDevice(A, CF_marker, num_cpts_global, interp_type, &P_ptr);
//hypre_ParCSRMatrixPrintIJ(P_ptr, 0, 0, "P_device");
}
#endif
/* Could do a few sweeps of Jacobi to further improve Jacobi interpolation P */ /* Could do a few sweeps of Jacobi to further improve Jacobi interpolation P */
/* /*
if(interp_type == 2) if(interp_type == 2)
@ -3233,17 +3233,17 @@ hypre_MGRBuildInterp(hypre_ParCSRMatrix *A,
} }
else if (interp_type == 4) else if (interp_type == 4)
{ {
if (exec == HYPRE_EXEC_HOST) #if defined(HYPRE_USING_CUDA)
if (exec == HYPRE_EXEC_DEVICE)
{
hypre_NoGPUSupport("interpolation");
}
else
#endif
{ {
hypre_MGRBuildInterpApproximateInverse(A, CF_marker, num_cpts_global, debug_flag, &P_ptr); hypre_MGRBuildInterpApproximateInverse(A, CF_marker, num_cpts_global, debug_flag, &P_ptr);
hypre_BoomerAMGInterpTruncation(P_ptr, trunc_factor, max_elmts); hypre_BoomerAMGInterpTruncation(P_ptr, trunc_factor, max_elmts);
} }
#if defined(HYPRE_USING_CUDA)
else
{
hypre_NoGPUSupport("interpolation");
}
#endif
} }
/* /*
else if (interp_type == 99) else if (interp_type == 99)
@ -3322,33 +3322,33 @@ hypre_MGRBuildRestrict(hypre_ParCSRMatrix *A,
/* Restriction for each level */ /* Restriction for each level */
if (restrict_type == 0) if (restrict_type == 0)
{ {
if (exec == HYPRE_EXEC_HOST)
{
hypre_MGRBuildP(A, CF_marker, num_cpts_global, restrict_type, debug_flag, &R_ptr);
//hypre_ParCSRMatrixPrintIJ(R_ptr, 0, 0, "R_host");
}
#if defined(HYPRE_USING_CUDA) #if defined(HYPRE_USING_CUDA)
else if (exec == HYPRE_EXEC_DEVICE)
{ {
hypre_MGRBuildPDevice(A, CF_marker, num_cpts_global, restrict_type, &R_ptr); hypre_MGRBuildPDevice(A, CF_marker, num_cpts_global, restrict_type, &R_ptr);
//hypre_ParCSRMatrixPrintIJ(R_ptr, 0, 0, "R_device"); //hypre_ParCSRMatrixPrintIJ(R_ptr, 0, 0, "R_device");
} }
else
#endif #endif
{
hypre_MGRBuildP(A, CF_marker, num_cpts_global, restrict_type, debug_flag, &R_ptr);
//hypre_ParCSRMatrixPrintIJ(R_ptr, 0, 0, "R_host");
}
} }
else if (restrict_type == 1 || restrict_type == 2) else if (restrict_type == 1 || restrict_type == 2)
{ {
if (exec == HYPRE_EXEC_HOST)
{
hypre_MGRBuildP(AT, CF_marker, num_cpts_global, restrict_type, debug_flag, &R_ptr);
//hypre_ParCSRMatrixPrintIJ(R_ptr, 0, 0, "R_host");
}
#if defined(HYPRE_USING_CUDA) #if defined(HYPRE_USING_CUDA)
else if (exec == HYPRE_EXEC_DEVICE)
{ {
hypre_MGRBuildPDevice(AT, CF_marker, num_cpts_global, restrict_type, &R_ptr); hypre_MGRBuildPDevice(AT, CF_marker, num_cpts_global, restrict_type, &R_ptr);
//hypre_ParCSRMatrixPrintIJ(R_ptr, 0, 0, "R_device"); //hypre_ParCSRMatrixPrintIJ(R_ptr, 0, 0, "R_device");
} }
else
#endif #endif
{
hypre_MGRBuildP(AT, CF_marker, num_cpts_global, restrict_type, debug_flag, &R_ptr);
//hypre_ParCSRMatrixPrintIJ(R_ptr, 0, 0, "R_host");
}
} }
else if (restrict_type == 3) else if (restrict_type == 3)
{ {

20
src/parcsr_mv/par_csr_triplemat_device.c
View File

@ -553,13 +553,6 @@ hypre_ParCSRMatrix*
hypre_ParCSRMatMatDevice( hypre_ParCSRMatrix *A, hypre_ParCSRMatMatDevice( hypre_ParCSRMatrix *A,
hypre_ParCSRMatrix *B ) hypre_ParCSRMatrix *B )
{ {
/* WM: debug */
hypre_printf("WM: debug - print P's\n");
/* hypre_ParCSRMatrixPrint(A, "P1"); */
/* hypre_ParCSRMatrixPrint(B, "P2"); */
HYPRE_Int my_id;
hypre_MPI_Comm_rank(hypre_MPI_COMM_WORLD, &my_id);
hypre_printf("WM: debug - rank %d, inside hypre_ParCSRMatMatDevice()\n", my_id);
hypre_ParCSRMatrix *C; hypre_ParCSRMatrix *C;
hypre_CSRMatrix *C_diag; hypre_CSRMatrix *C_diag;
hypre_CSRMatrix *C_offd; hypre_CSRMatrix *C_offd;
@ -745,8 +738,6 @@ hypre_ParCSRMatMatDevice( hypre_ParCSRMatrix *A,
HYPRE_MEMORY_HOST, HYPRE_MEMORY_DEVICE); HYPRE_MEMORY_HOST, HYPRE_MEMORY_DEVICE);
} }
hypre_printf("WM: debug - rank %d, finished hypre_ParCSRMatMatDevice()\n", my_id);
/* hypre_ParCSRMatrixPrint(C, "P"); */
return C; return C;
} }
@ -769,7 +760,6 @@ hypre_ParCSRTMatMatKTDevice( hypre_ParCSRMatrix *A,
hypre_MPI_Comm_size(hypre_ParCSRMatrixComm(A), &num_procs); hypre_MPI_Comm_size(hypre_ParCSRMatrixComm(A), &num_procs);
HYPRE_Int my_id; HYPRE_Int my_id;
hypre_MPI_Comm_rank(hypre_ParCSRMatrixComm(A), &my_id); hypre_MPI_Comm_rank(hypre_ParCSRMatrixComm(A), &my_id);
hypre_printf("WM: debug - rank %d, inside hypre_ParCSRTMatMatKTDevice()\n", my_id);
if (hypre_ParCSRMatrixGlobalNumRows(A) != hypre_ParCSRMatrixGlobalNumRows(B) || if (hypre_ParCSRMatrixGlobalNumRows(A) != hypre_ParCSRMatrixGlobalNumRows(B) ||
hypre_ParCSRMatrixNumRows(A) != hypre_ParCSRMatrixNumRows(B)) hypre_ParCSRMatrixNumRows(A) != hypre_ParCSRMatrixNumRows(B))
@ -1141,7 +1131,6 @@ hypre_ParCSRTMatMatKTDevice( hypre_ParCSRMatrix *A,
hypre_SyncComputeStream(hypre_handle()); hypre_SyncComputeStream(hypre_handle());
hypre_printf("WM: debug - rank %d, finished hypre_ParCSRTMatMatKTDevice()\n", my_id);
return C; return C;
} }
@ -1167,7 +1156,6 @@ hypre_ParCSRMatrixRAPKTDevice( hypre_ParCSRMatrix *R,
hypre_MPI_Comm_size(hypre_ParCSRMatrixComm(A), &num_procs); hypre_MPI_Comm_size(hypre_ParCSRMatrixComm(A), &num_procs);
HYPRE_Int my_id; HYPRE_Int my_id;
hypre_MPI_Comm_rank(hypre_ParCSRMatrixComm(A), &my_id); hypre_MPI_Comm_rank(hypre_ParCSRMatrixComm(A), &my_id);
hypre_printf("WM: debug - rank %d, inside hypre_ParCSRMatrixRAPKTDevice()\n", my_id);
if ( hypre_ParCSRMatrixGlobalNumRows(R) != hypre_ParCSRMatrixGlobalNumRows(A) || if ( hypre_ParCSRMatrixGlobalNumRows(R) != hypre_ParCSRMatrixGlobalNumRows(A) ||
hypre_ParCSRMatrixGlobalNumCols(A) != hypre_ParCSRMatrixGlobalNumRows(P) ) hypre_ParCSRMatrixGlobalNumCols(A) != hypre_ParCSRMatrixGlobalNumRows(P) )
@ -1555,14 +1543,6 @@ hypre_ParCSRMatrixRAPKTDevice( hypre_ParCSRMatrix *R,
hypre_SyncComputeStream(hypre_handle()); hypre_SyncComputeStream(hypre_handle());
/* hypre_printf("WM: debug - rank %d, finished hypre_ParCSRMatrixRAPKTDevice()\n", my_id); */
/* hypre_MatvecCommPkgCreate(C); */
/* hypre_ParCSRMatrixPrint(C, "C"); */
/* char my_filename[256]; */
/* hypre_sprintf(my_filename, "commPkg%d", my_id); */
/* hypre_ParCSRCommPkgPrint(hypre_ParCSRMatrixCommPkg(C), my_filename); */
/* hypre_MPI_Finalize(); */
/* exit(0); */
return C; return C;
} }

5
src/utilities/_hypre_utilities.hpp
View File

@ -211,10 +211,9 @@ using dim3 = sycl::range<1>;
{ \ { \
if ( gridsize[0] == 0 || blocksize[0] == 0 ) \ if ( gridsize[0] == 0 || blocksize[0] == 0 ) \
{ \ { \
hypre_printf("Error %s %d: Invalid SYCL 1D launch parameters grid/block (%d) (%d)\n", \ /* hypre_printf("Warning %s %d: Zero SYCL 1D launch parameters grid/block (%d) (%d)\n", \
__FILE__, __LINE__, \ __FILE__, __LINE__, \
gridsize[0], blocksize[0]); \ gridsize[0], blocksize[0]); */ \
assert(0); exit(1); \
} \ } \
else \ else \
{ \ { \

5
src/utilities/device_utils.h
View File

@ -154,10 +154,9 @@ using dim3 = sycl::range<1>;
{ \ { \
if ( gridsize[0] == 0 || blocksize[0] == 0 ) \ if ( gridsize[0] == 0 || blocksize[0] == 0 ) \
{ \ { \
hypre_printf("Error %s %d: Invalid SYCL 1D launch parameters grid/block (%d) (%d)\n", \ /* hypre_printf("Warning %s %d: Zero SYCL 1D launch parameters grid/block (%d) (%d)\n", \
__FILE__, __LINE__, \ __FILE__, __LINE__, \
gridsize[0], blocksize[0]); \ gridsize[0], blocksize[0]); */ \
assert(0); exit(1); \
} \ } \
else \ else \
{ \ { \