Skip to content

Commit

Permalink
nvector: ternary atomic test 2
Browse files Browse the repository at this point in the history
  • Loading branch information
jsdomine committed Jul 7, 2023
1 parent 204d292 commit 26c41f8
Showing 1 changed file with 33 additions and 24 deletions.
57 changes: 33 additions & 24 deletions src/nvector/parhyp/nvector_parhyp.c
Original file line number Diff line number Diff line change
Expand Up @@ -1212,30 +1212,39 @@ realtype N_VWSqrSumMaskLocal_ParHyp(N_Vector x, N_Vector w, N_Vector id)
const size_t buffer_size = atomic ? 1 : grid;
NV_CATCH_ERR_PH(InitializeReductionBuffer(x, sum, buffer_size))

if (atomic)
{
wL2NormSquareMaskKernel<realtype, sunindextype, GridReducerAtomic><<<grid, block, shMemSize, stream>>>
(
xd,
wd,
idd,
NV_DBUFFERp_PH(x),
N,
nullptr
);
}
else
{
wL2NormSquareMaskKernel<realtype, sunindextype, GridReducerLDS><<<grid, block, shMemSize, stream>>>
(
xd,
wd,
idd,
NV_DBUFFERp_PH(x),
N,
NV_DCOUNTERp_PH(x)
);
}
wL2NormSquareMaskKernel<realtype, sunindextype, (atomic?GridReducerAtomic:GridReducerLDS)><<<grid, block, shMemSize, stream>>>
(
xd,
wd,
idd,
NV_DBUFFERp_PH(x),
N,
(atomic?nullptr:NV_DCOUNTERp_PH(x))
);
// if (atomic)
// {
// wL2NormSquareMaskKernel<realtype, sunindextype, GridReducerAtomic><<<grid, block, shMemSize, stream>>>
// (
// xd,
// wd,
// idd,
// NV_DBUFFERp_PH(x),
// N,
// nullptr
// );
// }
// else
// {
// wL2NormSquareMaskKernel<realtype, sunindextype, GridReducerLDS><<<grid, block, shMemSize, stream>>>
// (
// xd,
// wd,
// idd,
// NV_DBUFFERp_PH(x),
// N,
// NV_DCOUNTERp_PH(x)
// );
// }
PostKernelLaunch();

// Get result from the GPU
Expand Down

0 comments on commit 26c41f8

Please sign in to comment.