NVIDIA / NVIDIA/cccl

Don't call the predicate twice in remove_if

Open
#702 7 comments 0 reactions 1 assignee Claimed by @brycelelbach View on GitHub
thrust
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

Hello Thrust Team,

While playing around with learning Thrust, I happened to put a printf statement in the predicate-functor passed to `thrust::remove_if()`, and noticed some curious behavior. When run on host_vectors (ie on the CPU) the printf statement is output once per element, but when run on device vectors, the printf statement is output twice per element.

According to [this](https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-optimized-filtering-warp-aggregated-atomics/), Thrust uses a parallel scan to determine output positions for the non-removed elements in a stable way, so it must do something like run the predicate on each element, then create an array of "hit" tags corresponding to each element with `hit[i] = 1` if `predicate( array[i] ) == false` (don't remove) and `hit[i] = 0` if `predicate( array[i] ) == true` (do remove). An exclusive scan on `hit` then gives the output positions for filtered elements of `array`...I'm just saying all this to highlight that I can't for the life of me imagine why the predicate functor would need to be called twice per element. What am I missing??

I have attached a minimal working sample, identical to the [remove_points2d.cu](https://github.com/thrust/thrust/blob/master/examples/remove_points2d.cu) example except that it has a printf statement in the predicate functor and includes device_vector.h in addition to host_vector.h. If x and y are host_vectors, "In here" gets printed 20 times, but if x and y are device_vectors, it gets printed 40 times.

[remove_points_devicevshost.cu.txt](https://github.com/thrust/thrust/files/475825/remove_points_devicevshost.cu.txt)

Thanks in advance for your insight,
Michael

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.