Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Don't call the predicate twice in remove_if #702

Open
mcarilli opened this issue Sep 15, 2016 · 7 comments
Open

Don't call the predicate twice in remove_if #702

mcarilli opened this issue Sep 15, 2016 · 7 comments
Assignees
Labels
thrust For all items related to Thrust.

Comments

@mcarilli
Copy link

mcarilli commented Sep 15, 2016

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, 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 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

Thanks in advance for your insight,
Michael

@3gx
Copy link

3gx commented Sep 16, 2016

This is Thrust implementation detail of how reduce_if is implemented: the first time predicate is called for scan operation, and the second time to check whether the element needs to be copied (remove_if is a wrapper call to copy_if with a negated predicate). This implementation detail may change in future and one should not rely that reduce_if is called twice.

Hope this helps.

@mcarilli
Copy link
Author

I don't think you need to call the functor the second time. I think you could do something like the following: In the hit array record a 1 for each element you want to keep, as I described above. Then run the scan on hit. On the second pass over the data elements, instead of rerunning the functor on element i, you can just compare hits_scanned[i] with hits_scanned[i+1]. If hits_scanned[i+1] = hits_scanned[i]+1, it means the exclusive scan recorded a hit at location i, and the element should be copied. There's nothing stopping you from doing the scan in-place so hits_scanned could be the same memory as hit but I renamed it to hopefully make my point clearer.

Thank you for the information in any case. I was mostly seeking reassurance that I hadn't somehow screwed up my usage of the thrust functions.

@3gx
Copy link

3gx commented Sep 16, 2016

Agree, there is no need having two predicate invocations, and as I said, this is something that may change in future.

@brycelelbach brycelelbach self-assigned this Sep 15, 2017
@brycelelbach
Copy link
Collaborator

brycelelbach commented Oct 13, 2017

Yah, this is and remains a problem in many places. Note that, by design, the C++17 parallel algorithms relax the complexity guarantees of the sequential algorithms, to allow flexibility and redundant computations. However, there are many places in Thrust where we go overboard with that, and assume (either by design or by mistake) that all user-defined types are TriviallyCopyable, all Callables are pure and that they can be called on invalid data.

@brycelelbach brycelelbach added the nvbug Has an associated internal NVIDIA NVBug. label Oct 13, 2017
@brycelelbach brycelelbach changed the title Why does remove_if need to call the predicate filter twice?? Don't call the predicate twice in remove_if Feb 13, 2018
@brycelelbach
Copy link
Collaborator

Tracked by NVBug 2062288.

@mcarilli
Copy link
Author

Ha, i’m actually at Nvidia now. I joined Frameworks two and a half months ago. Come to think of it, i sit pretty close to @egaburov. I’ll say hi tomorrow.

@brycelelbach
Copy link
Collaborator

brycelelbach commented Feb 14, 2018 via email

@jrhemstad jrhemstad added the thrust For all items related to Thrust. label Feb 22, 2023
@jrhemstad jrhemstad removed the nvbug Has an associated internal NVIDIA NVBug. label Mar 7, 2023
@jarmak-nv jarmak-nv transferred this issue from NVIDIA/thrust Nov 8, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
thrust For all items related to Thrust.
Projects
Status: Todo
Development

No branches or pull requests

4 participants