Skip to content

Commit

Permalink
nvector: test atomic ternary
Browse files Browse the repository at this point in the history
  • Loading branch information
jsdomine committed Jul 7, 2023
1 parent ed1739b commit 70119e4
Showing 1 changed file with 30 additions and 22 deletions.
52 changes: 30 additions & 22 deletions src/nvector/parhyp/nvector_parhyp.c
Original file line number Diff line number Diff line change
Expand Up @@ -1144,28 +1144,36 @@ realtype N_VWSqrSumLocal_ParHyp(N_Vector x, N_Vector w)
const size_t buffer_size = atomic ? 1 : grid;
NV_CATCH_ERR_PH(InitializeReductionBuffer(x, sum, buffer_size))

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

// Get result from the GPU
Expand Down

0 comments on commit 70119e4

Please sign in to comment.