Skip to content

Commit

Permalink
bugfix: wrong assert() makes GPU kernel crashed
Browse files Browse the repository at this point in the history
issue reported at heterodb#702
  • Loading branch information
kaigai committed Dec 25, 2023
1 parent 346cd25 commit b73ff0b
Show file tree
Hide file tree
Showing 2 changed files with 0 additions and 2 deletions.
1 change: 0 additions & 1 deletion src/cuda_gpupreagg.cu
Original file line number Diff line number Diff line change
Expand Up @@ -576,7 +576,6 @@ __update_nogroups__psum_int(kern_context *kcxt,
kagg_state__psum_int_packed *r =
(kagg_state__psum_int_packed *)buffer;

assert(ival >= 0);
pgstrom_stair_sum_int64(ival, &sum);
if (get_local_id() == 0)
{
Expand Down
1 change: 0 additions & 1 deletion src/cuda_gpuscan.cu
Original file line number Diff line number Diff line change
Expand Up @@ -95,7 +95,6 @@ pgstrom_stair_sum_binary(bool predicate, uint32_t *p_total_count)
\
assert(__activemask() == ~0U); \
warp_sum = __stair_sum_warp_common(value); \
assert(warp_sum >= value); \
if (LaneId() == warpSize - 1) \
__stair_sum_buffer.FIELD[warp_id] = warp_sum; \
__syncthreads(); \
Expand Down

0 comments on commit b73ff0b

Please sign in to comment.