Add device non-shmem pair_retrieve functions
#122
Conversation
PointKernel
added
type: feature request
|
@vyasr @jrhemstad Here is preliminary work mainly to get your feedback. Will add corresponding tests shortly. |
| // TODO: `=` operator cannot work here | ||
| // *(probe_output_begin + output_idx) = pair; | ||
| // *(contained_output_begin + output_idx) = arr[0]; | ||
| thrust::get<0>(*(probe_output_begin + output_idx)) = pair.first; | ||
| thrust::get<1>(*(probe_output_begin + output_idx)) = pair.second; | ||
| thrust::get<0>(*(contained_output_begin + output_idx)) = arr[0].first; | ||
| thrust::get<1>(*(contained_output_begin + output_idx)) = arr[0].second; |
There was a problem hiding this comment.
no idea by what miracle the = operator works just fine right now.
There was a problem hiding this comment.
By writing unit tests for non-shmem pair_retrieve, I see the error again:
/home/yunsongw/Work/cuCollections/include/cuco/detail/static_multimap/device_view_impl.inl(1120): error: no operator "=" matches these operands
operand types are: thrust::detail::tuple_of_iterator_references<thrust::detail::any_assign &, thrust::detail::any_assign &> = const cuco::pair<int32_t, int32_t>
and in this case, both probe_output_begin and contained_output_begin are defined as
auto begin = thrust::make_zip_iterator(
thrust::make_tuple(thrust::make_discard_iterator(), thrust::make_discard_iterator()));
There was a problem hiding this comment.
In light of our discussion offline, let's add a comment as to why we must assign the iterator elements individually rather than as a pair.
|
FYI I plan to review this, I'm just trying to get my code far enough along to be able to actually use this feature so that I can actually test it in action in addition to reviewing the code. |
|
@vyasr are you still planning to review? |
|
@jrhemstad yes, sorry for the long delay. I'm very close to having a reviewable prototype of the code in cudf that's using this, so I've been focusing on ironing out issues there in case any of them have an impact on this API/implementation. I'll probably be able to review this tomorrow. |
This PR added device
pair_retrievefunctions that don't use a shared memory output buffer. It included code cleanups like usingcount_least_significant_bits()as well.