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

Silence uninitialized var HIP warnings

Browse Source 
This should be a benign change. What happens is that the first one or two
workitems/threads in each workgroup/block read a value and then broadcast
it (with __shfl_sync or similar) and then code branching happens
based on this value. But the compiler can't see it all the way
through it so we get some uninitialized var warnings.
This commit is contained in:
Paul T. Bauman 2022-03-09 13:26:36 -06:00
parent 251cd3d269
commit d55a409bdb
12 changed files with 35 additions and 36 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
src/parcsr_ls/ams.c
View File

@ -388,7 +388,7 @@ hypreCUDAKernel_ParCSRMatrixFixZeroRows( HYPRE_Int nrows,
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Real eps = 0.0; /* DBL_EPSILON * 1e+4; */
HYPRE_Real l1_norm = 0.0;
HYPRE_Int p1, q1, p2 = 0, q2 = 0;
HYPRE_Int p1=0, q1, p2 = 0, q2 = 0;
if (lane < 2)
{
@ -1495,7 +1495,7 @@ hypreCUDAKernel_AMSComputePi_copy2(HYPRE_Int nrows,
}
const HYPRE_Int lane_id = hypre_cuda_get_lane_id<1>();
HYPRE_Int j, istart, iend;
HYPRE_Int j=0, istart, iend;
HYPRE_Real t, G[3], *Gdata[3];
Gdata[0] = Gx_data;
@ -1781,7 +1781,7 @@ hypreCUDAKernel_AMSComputePixyz_copy(HYPRE_Int nrows,
}
const HYPRE_Int lane_id = hypre_cuda_get_lane_id<1>();
HYPRE_Int j, istart, iend;
HYPRE_Int j=0, istart, iend;
HYPRE_Real t, G[3], *Gdata[3], *Odata[3];
Gdata[0] = Gx_data;
@ -2363,7 +2363,7 @@ hypreCUDAKernel_AMSComputeGPi_copy2(HYPRE_Int nrows,
}
const HYPRE_Int lane_id = hypre_cuda_get_lane_id<1>();
HYPRE_Int j, istart, iend;
HYPRE_Int j=0, istart, iend;
HYPRE_Real t, G[3], *Gdata[3];
Gdata[0] = Gx_data;
@ -2670,7 +2670,7 @@ hypreCUDAKernel_FixInterNodes( HYPRE_Int nrows,
return;
}
HYPRE_Int p1, q1, p2 = 0, q2 = 0;
HYPRE_Int p1=0, q1, p2 = 0, q2 = 0;
bool nonempty_offd = G0t_offd_data != NULL;
if (lane < 2)
@ -2720,7 +2720,7 @@ hypreCUDAKernel_AMSSetupScaleGGt( HYPRE_Int Gt_num_rows,
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Real h2 = 0.0;
HYPRE_Int ne, p1, q1, p2 = 0, q2 = 0;
HYPRE_Int ne, p1=0, q1, p2 = 0, q2 = 0;
if (lane < 2)
{

12
src/parcsr_ls/par_2s_interp_device.c
View File

@ -575,8 +575,8 @@ void hypreCUDAKernel_MMInterpScaleAFF( HYPRE_Int AFF_nrows,
}
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int ib_diag, ie_diag;
HYPRE_Int rowF;
HYPRE_Int ib_diag=0, ie_diag;
HYPRE_Int rowF=0;
if (lane == 0)
{
@ -624,7 +624,7 @@ void hypreCUDAKernel_MMInterpScaleAFF( HYPRE_Int AFF_nrows,
}
}
HYPRE_Int ib_offd, ie_offd;
HYPRE_Int ib_offd=0, ie_offd;
if (lane < 2)
{
@ -707,8 +707,8 @@ void hypreCUDAKernel_MMPEInterpScaleAFF( HYPRE_Int AFF_nrows,
}
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int ib_diag, ie_diag;
HYPRE_Int rowF;
HYPRE_Int ib_diag=0, ie_diag;
HYPRE_Int rowF=0;
if (lane == 0)
{
@ -757,7 +757,7 @@ void hypreCUDAKernel_MMPEInterpScaleAFF( HYPRE_Int AFF_nrows,
}
}
HYPRE_Int ib_offd, ie_offd;
HYPRE_Int ib_offd=0, ie_offd;
if (lane < 2)
{

3
src/parcsr_ls/par_coarsen_device.c
View File

@ -381,7 +381,7 @@ hypreCUDAKernel_PMISCoarseningUpdateCF(HYPRE_Int graph_diag_size,
}
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int row, i, marker_row, row_start, row_end;
HYPRE_Int row, i=0, marker_row, row_start, row_end;
if (lane < 2)
{
@ -543,4 +543,3 @@ hypre_PMISCoarseningUpdateCFDevice( hypre_ParCSRMatrix *S, /* in
}
#endif // #if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)

4
src/parcsr_ls/par_indepset_device.c
View File

@ -35,8 +35,8 @@ hypreCUDAKernel_IndepSetMain(HYPRE_Int graph_diag_size,
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int row, row_start, row_end;
HYPRE_Int i, j;
HYPRE_Real t, measure_row;
HYPRE_Int i=0, j;
HYPRE_Real t=0.0, measure_row;
HYPRE_Int marker_row = 1;
if (lane < 2)

16
src/parcsr_ls/par_interp_device.c
View File

@ -426,7 +426,7 @@ hypre_BoomerAMGBuildDirInterp_getnnz( HYPRE_Int nr_of_rows,
return;
}
HYPRE_Int p, q, dof_func_i;
HYPRE_Int p=0, q, dof_func_i=0;
HYPRE_Int jPd = 0, jPo = 0;
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
@ -570,7 +570,7 @@ hypre_BoomerAMGBuildDirInterp_getcoef( HYPRE_Int nr_of_rows,
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int k, dof_func_i;
HYPRE_Int k=0, dof_func_i=0;
if (lane == 0)
{
@ -608,7 +608,7 @@ hypre_BoomerAMGBuildDirInterp_getcoef( HYPRE_Int nr_of_rows,
HYPRE_Real diagonal = 0.0, sum_N_pos = 0.0, sum_N_neg = 0.0, sum_P_pos = 0.0, sum_P_neg = 0.0;
/* diag part */
HYPRE_Int p_diag_A, q_diag_A, p_diag_P, q_diag_P;
HYPRE_Int p_diag_A=0, q_diag_A, p_diag_P=0, q_diag_P;
if (lane < 2)
{
p_diag_A = read_only_load(A_diag_i + i + lane);
@ -677,7 +677,7 @@ hypre_BoomerAMGBuildDirInterp_getcoef( HYPRE_Int nr_of_rows,
hypre_device_assert(k == q_diag_P);
/* offd part */
HYPRE_Int p_offd_A, q_offd_A, p_offd_P, q_offd_P;
HYPRE_Int p_offd_A=0, q_offd_A, p_offd_P=0, q_offd_P;
if (lane < 2)
{
p_offd_A = read_only_load(A_offd_i + i + lane);
@ -842,7 +842,7 @@ hypre_BoomerAMGBuildDirInterp_getcoef_v2( HYPRE_Int nr_of_rows,
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int k, dof_func_i;
HYPRE_Int k=0, dof_func_i=0;
if (lane == 0)
{
@ -880,7 +880,7 @@ hypre_BoomerAMGBuildDirInterp_getcoef_v2( HYPRE_Int nr_of_rows,
HYPRE_Real diagonal = 0.0, sum_F = 0.0;
/* diag part */
HYPRE_Int p_diag_A, q_diag_A, p_diag_P, q_diag_P;
HYPRE_Int p_diag_A=0, q_diag_A, p_diag_P=0, q_diag_P;
if (lane < 2)
{
p_diag_A = read_only_load(A_diag_i + i + lane);
@ -941,7 +941,7 @@ hypre_BoomerAMGBuildDirInterp_getcoef_v2( HYPRE_Int nr_of_rows,
hypre_device_assert(k == q_diag_P);
/* offd part */
HYPRE_Int p_offd_A, q_offd_A, p_offd_P, q_offd_P;
HYPRE_Int p_offd_A=0, q_offd_A, p_offd_P=0, q_offd_P;
if (lane < 2)
{
p_offd_A = read_only_load(A_offd_i + i + lane);
@ -1383,7 +1383,7 @@ hypre_BoomerAMGBuildInterpOnePnt_getnnz( HYPRE_Int nr_of_rows,
return;
}
HYPRE_Int p, q;
HYPRE_Int p=0, q;
HYPRE_Int max_j_diag = -1, max_j_offd = -1;
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Real max_diag = -1.0, max_offd = -1.0;

2
src/parcsr_ls/par_interp_trunc_device.c
View File

@ -26,7 +26,7 @@ hypreCUDAKernel_InterpTruncation( HYPRE_Int nrows,
return;
}
HYPRE_Int lane = hypre_cuda_get_lane_id<1>(), p, q;
HYPRE_Int lane = hypre_cuda_get_lane_id<1>(), p=0, q;
/* 1. compute row max, rowsum */
if (lane < 2)

2
src/parcsr_ls/par_lr_restr_device.c
View File

@ -339,7 +339,7 @@ hypre_BoomerAMGBuildRestrNeumannAIR_assembleRdiag( HYPRE_Int nr_of_rows,
return;
}
HYPRE_Int p, q, pZ;
HYPRE_Int p=0, q, pZ=0;
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
/* diag part */

8
src/parcsr_ls/par_mod_multi_interp_device.c
View File

@ -1667,7 +1667,7 @@ void hypreCUDAKernel_generate_Pdiag_j_Poffd_j( HYPRE_Int num_points,
HYPRE_Int i1 = read_only_load(&pass_order[row_i]);
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int p_diag_A, q_diag_A, p_diag_P, q_diag_P;
HYPRE_Int p_diag_A=0, q_diag_A, p_diag_P=0, q_diag_P;
HYPRE_Int k;
HYPRE_Complex row_sum_C = 0.0, diagonal = 0.0;
@ -1717,7 +1717,7 @@ void hypreCUDAKernel_generate_Pdiag_j_Poffd_j( HYPRE_Int num_points,
hypre_device_assert(k == q_diag_P);
// S_offd
HYPRE_Int p_offd_A, q_offd_A, p_offd_P, q_offd_P;
HYPRE_Int p_offd_A=0, q_offd_A, p_offd_P=0, q_offd_P;
if (lane < 2)
{
@ -1903,7 +1903,7 @@ void hypreCUDAKernel_generate_Qdiag_j_Qoffd_j( HYPRE_Int num_points,
HYPRE_Int i1 = read_only_load(&pass_order[row_i]);
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int p_diag_A, q_diag_A, p_diag_P;
HYPRE_Int p_diag_A=0, q_diag_A, p_diag_P=0;
#ifdef HYPRE_DEBUG
HYPRE_Int q_diag_P;
#endif
@ -1990,7 +1990,7 @@ void hypreCUDAKernel_generate_Qdiag_j_Qoffd_j( HYPRE_Int num_points,
#endif
// S_offd
HYPRE_Int p_offd_A, q_offd_A, p_offd_P;
HYPRE_Int p_offd_A=0, q_offd_A, p_offd_P=0;
#ifdef HYPRE_DEBUG
HYPRE_Int q_offd_P;
#endif

2
src/parcsr_ls/par_relax_more_device.c
View File

@ -41,7 +41,7 @@ hypreCUDAKernel_CSRMaxEigEstimate(HYPRE_Int nrows,
}
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int p, q;
HYPRE_Int p=0, q;
HYPRE_Complex diag_value = 0.0;
HYPRE_Complex row_sum_i = 0.0;

4
src/parcsr_ls/par_strength_device.c
View File

@ -268,7 +268,7 @@ __global__ void hypre_BoomerAMGCreateS_rowcount( HYPRE_Int nr_of_rows,
}
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int p_diag, q_diag, p_offd, q_offd;
HYPRE_Int p_diag=0, q_diag, p_offd=0, q_offd;
/* diag part */
if (lane < 2)
@ -467,7 +467,7 @@ __global__ void hypre_BoomerAMGCreateSabs_rowcount( HYPRE_Int nr_of_rows,
}
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int p_diag, q_diag, p_offd, q_offd;
HYPRE_Int p_diag=0, q_diag, p_offd=0, q_offd;
/* diag part */
if (lane < 2)

4
src/parcsr_mv/par_csr_matop_device.c
View File

@ -565,7 +565,7 @@ hypreCUDAKernel_ConcatDiagAndOffd(HYPRE_Int nrows, HYPRE_Int diag_ncol,
/* lane id inside the warp */
const HYPRE_Int lane_id = hypre_cuda_get_lane_id<1>();
HYPRE_Int i, j, k, p, istart, iend, bstart;
HYPRE_Int i, j=0, k=0, p, istart, iend, bstart;
/* diag part */
if (lane_id < 2)
@ -1078,7 +1078,7 @@ hypre_ParCSRMatrixDropSmallEntriesDevice_getElmtTols( HYPRE_Int nrows,
}
HYPRE_Int lane = hypre_cuda_get_lane_id<1>();
HYPRE_Int p_diag, p_offd, q_diag, q_offd;
HYPRE_Int p_diag=0, p_offd=0, q_diag=0, q_offd=0;
/* sum row norm over diag part */
if (lane < 2)

2
src/utilities/device_utils.c
View File

@ -376,7 +376,7 @@ hypreCUDAKernel_CopyParCSRRows(HYPRE_Int nrows,
/* lane id inside the warp */
const HYPRE_Int lane_id = hypre_cuda_get_lane_id<1>();
HYPRE_Int i, j, k, p, row, istart, iend, bstart;
HYPRE_Int i, j=0, k=0, p, row, istart, iend, bstart;
/* diag part */
if (lane_id < 2)