Skip to content

Commit

Permalink
Remove cuda::proclaim_return_type from nested lambda (#14607)
Browse files Browse the repository at this point in the history
This removes `cuda::proclaim_return_type` from a device lambda because that lambda is going to be nested inside another device lambda, which is in turn enclosed by `cuda::proclaim_return_type`. 

This PR is to fix a compile issue that we encountered:
```
/usr/local/cuda/include/cuda/std/detail/libcxx/include/__functional/invoke.h(402): error: 
calling a __device__ function("cudf::tdigest::detail::_NV_ANON_NAMESPACE::build_output_column(int,   
  ::std::unique_ptr<   ::cudf::column,     ::std::default_delete<   ::cudf::column> >  &&,     ::std::unique_ptr<   ::cudf::column,     ::std::default_delete<   ::cudf::column> >  &&,     ::std::unique_ptr<   ::cudf::column,     ::std::default_delete<   ::cudf::column> >  &&,     ::std::unique_ptr<   ::cudf::column,     ::std::default_delete<   ::cudf::column> >  &&,     ::std::unique_ptr<   ::cudf::column,     ::std::default_delete<   ::cudf::column> >  &&, bool,  ::rmm::cuda_stream_view,  ::rmm::mr::device_memory_resource *)
::[lambda(int) (instance 2)]::operator ()(int) const") from a __host__ __device__ function("__invoke") is not allowed
```

Note: The issue is reproducible only in our build environment: ARM architecture, cuda 12 + rockylinux8.

Closes #14610.

Authors:
  - Nghia Truong (https://github.com/ttnghia)

Approvers:
  - Michael Schellenberger Costa (https://github.com/miscco)
  - Karthikeyan (https://github.com/karthikeyann)
  - https://github.com/nvdbaranec
  - Bradley Dice (https://github.com/bdice)

URL: #14607
  • Loading branch information
ttnghia authored Dec 11, 2023
1 parent fc0f181 commit fcaebeb
Showing 1 changed file with 7 additions and 6 deletions.
13 changes: 7 additions & 6 deletions cpp/src/quantiles/tdigest/tdigest_aggregation.cu
Original file line number Diff line number Diff line change
Expand Up @@ -587,15 +587,16 @@ std::unique_ptr<column> build_output_column(size_type num_rows,
auto is_stub_weight = [weights = weights->view().begin<double>()] __device__(size_type i) {
return weights[i] == 0;
};
// whether or not this particular tdigest is a stub
auto is_stub_digest = cuda::proclaim_return_type<size_type>(
[offsets = offsets->view().begin<size_type>(), is_stub_weight] __device__(size_type i) {
return is_stub_weight(offsets[i]) ? 1 : 0;
});
// Whether or not this particular tdigest is a stub.
// This should not be wrapped in `proclaim_return_type` as it will be used inside another
// device lambda.
auto is_stub_digest = [offsets = offsets->view().begin<size_type>(), is_stub_weight] __device__(
size_type i) { return is_stub_weight(offsets[i]) ? 1 : 0; };

size_type const num_stubs = [&]() {
if (!has_nulls) { return 0; }
auto iter = cudf::detail::make_counting_transform_iterator(0, is_stub_digest);
auto iter = cudf::detail::make_counting_transform_iterator(
0, cuda::proclaim_return_type<size_type>(is_stub_digest));
return thrust::reduce(rmm::exec_policy(stream), iter, iter + num_rows);
}();

Expand Down

0 comments on commit fcaebeb

Please sign in to comment.