-
Notifications
You must be signed in to change notification settings - Fork 247
[BUG] cccl_async_resource_ref conversion operator triggers nvcc error 20013 with C++20 constexpr std::vector #2255
Description
Describe the bug
When using RMM in cudf, compiling code that involves cccl_async_resource_ref without the relaxed constexpr build option can trigger nvcc error #20013-D:
rmm/detail/cccl_adaptors.hpp(560): error #20013-D: calling a constexpr __host__ function("std::_Vector_base<...>::_Vector_impl::_Vector_impl") from a __host__ __device__ function("std::_Vector_base<...>::_Vector_base") is not allowedWhen constructing cuda::mr::any_resource at
rmm/cpp/include/rmm/detail/cccl_adaptors.hpp
Line 560 in ecba059
| return cuda::mr::any_resource<Properties...>{ref_}; |
the type erasure mechanism internally involves std::vector; with C++20, NVCC infers __host__ __device__ for std::_Vector_base's constexpr constructor, but std::_Vector_impl's move constructor remains __host__ only, creating an illegal call chain that triggers error 20013.
Steps/Code to reproduce bug
Building cudf without using the relaxed constexpr flag
Expected behavior
The existing diagnostic suppression only suppresses error 20011:
rmm/cpp/include/rmm/detail/cccl_adaptors.hpp
Lines 285 to 287 in ecba059
| #pragma nv_diagnostic push | |
| #pragma nv_diag_suppress 20011 | |
| #endif |
We probably want to silence 20013 as well
#ifdef __CUDACC__
#pragma nv_diagnostic push
#pragma nv_diag_suppress 20011
#pragma nv_diag_suppress 20013
#endifI'm not sure if this silence makes general sense though.
Additional context
This is needed to unblock rapidsai/cudf#7795.
The error is triggered by thrust::tabulate calls at https://github.com/rapidsai/cudf/blob/1ee5878631d7f4f0c0e48275116be79573cee3ab/cpp/src/groupby/common/m2_var_std.cu#L46-L56 and https://github.com/rapidsai/cudf/blob/1ee5878631d7f4f0c0e48275116be79573cee3ab/cpp/src/groupby/common/m2_var_std.cu#L135. The execution policy internally uses cccl_async_resource_ref, and during template instantiation with iterators involving mutable_column_view from the zip_iterator at line 134 and type_dispatcher at line 97, CCCL templates instantiate std::vector<mutable_column_view>, which causes NVCC's C++20 constexpr inference to compile std::_Vector_base's constructor as __host__ __device__ while std::_Vector_impl's move constructor remains __host__ only.
Metadata
Metadata
Assignees
Labels
Type
Projects
Status