Skip to content

Commit

Permalink
Explicitly build the dim3 objects for the kernel dimensions to keep H…
Browse files Browse the repository at this point in the history
…IP happy.
  • Loading branch information
xavierandrade committed Oct 11, 2024
1 parent dd0fe52 commit 6b12600
Showing 1 changed file with 8 additions and 2 deletions.
10 changes: 8 additions & 2 deletions external_libs/gpurun/include/gpu/reduce.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -171,7 +171,10 @@ auto run(reduce const & redx, reduce const & redy, kernel_type kernel) -> declty

gpu::array<type, 2> result({nblockx, nblocky});

reduce_kernel_rr<<<{nblockx, nblocky}, {bsizex, bsizey}, bsizex*bsizey*sizeof(type)>>>(sizex, sizey, kernel, begin(result));
struct dim3 dg{nblockx, nblocky};
struct dim3 db{bsizex, bsizey};

reduce_kernel_rr<<<dg, db, bsizex*bsizey*sizeof(type)>>>(sizex, sizey, kernel, begin(result));
check_error(last_error());

if(nblockx*nblocky == 1) {
Expand Down Expand Up @@ -256,7 +259,10 @@ auto run(reduce const & redx, reduce const & redy, reduce const & redz, kernel_t

gpu::array<type, 3> result({nblockx, nblocky, nblockz});

reduce_kernel_rrr<<<{nblockx, nblocky, nblockz}, {bsizex, bsizey, bsizez}, bsizex*bsizey*bsizez*sizeof(type)>>>(sizex, sizey, sizez, kernel, begin(result));
struct dim3 dg{nblockx, nblocky, nblockz};
struct dim3 db{bsizex, bsizey, bsizez};

reduce_kernel_rrr<<<dg, db, bsizex*bsizey*bsizez*sizeof(type)>>>(sizex, sizey, sizez, kernel, begin(result));
check_error(last_error());

if(nblockx*nblocky*nblockz == 1) {
Expand Down

0 comments on commit 6b12600

Please sign in to comment.