NVIDIA
`iterator_category_to_system` returns `host_system_tag` for `input_device_iterator_tag` instead of `device_system_tag`
Hello,
I was trying to design my own iterators to make them compatible with thrust. I stumbled across some problems and dug inside the code. Ultimately, it seems that the following statement fails to compile:
static_assert(std::is_same<thrust::detail::iterator_category_to_systemthrust::input_device_iterator_tag::type, thrust::device_system_tag>::value, "test");
While the following passes:
static_assert(std::is_same<thrust::detail::iterator_category_to_systemthrust::input_device_iterator_tag::type, thrust::host_system_tag>::value, "test");
That is not really surprising given the code of iterator_category_to_system and the fact that input_device_iterator_tag actually inherits from std::input_iterator_tag, which is just another name for input_host_iterator_tag (see iterator_categories.h).
Am I missing something or there is a bug there?
Note that all classes inheriting from thrust::iterator_adapter seem to be fine. Still trying to figure out why.
Yah, this is not good, at all. I'm not sure the impact but this could potentially be causing parts of Thrust to dispatch to host execution instead of device execution.
Minimal test case below. @dawagnbr, thanks for the bug report - in the future please follow https://github.com/brycelelbach/cpp_bug_reporting_guidelines , it makes life a little easier for me.
/* This content is CUDA C++ source code (a `.cu` file).
NVBug: 2062266
GH Issue: https://github.com/thrust/thrust/issues/902
Reporter: Bryce Adelstein Lelbach <[email protected]>
=========
TEST CASE
=========
`iterator_category_to_system` returns `host_system_tag` for
`input_device_iterator_tag` instead of `device_system_tag`. This could
potentially be causing host execution instead of device execution in some
places.
*/
#include <thrust/device_vector.h>
static_assert(
std::is_same<
thrust::detail::iterator_category_to_system<
thrust::input_device_iterator_tag
>::type
, thrust::device_system_tag
>::value
, "not device"
);
static_assert(
std::is_same<
thrust::detail::iterator_category_to_system<
thrust::input_device_iterator_tag
>::type
, thrust::host_system_tag
>::value
, "not host"
);
int main() {}
/*
==================
STEPS TO REPRODUCE
==================
Run the following sequence of shell commands in the environment specified
below:
TEST=iterator_category_to_system_returns_host_system_for_device_input_iterator_tag
NVCC=/path/to/nvcc
${NVCC} ${TEST}.cu -o ${TEST}
===============
EXPECTED OUTPUT
===============
Program compiles successfully.
===============
OBSERVED OUTPUT
===============
The "not device" static assertion fails and the "not host" static assertion
passes.
===========
ENVIRONMENT
===========
This bug can be reproduced on:
CUDA Toolkit: At least 9.0 and later.
Host Compiler: All supported GCC versions.
My local environment:
CPU: i7-5820K Haswell 6-core 3300MHz (OC to 4000MHz)
CPU Arch: x86-64
CPU Memory: 32 GB DDR4 2133MHz (OC to 2400Mhz)
GPU: Quadro K2000
GPU Arch: Kepler SM30 GK107
OS: Debian 9.2
Kernel: Linux 4.9.0
CUDA KMD: Local driver and compiler build from CL 23542456
CUDA Toolkit: //sw/gpgpu local build from CL 23542456
Host Compiler: GCC 6.4.0
C Standard Library: GNU C Library 2.24
C++ Standard Library: libstdc++ 6.0.24
*/
Just ran across this issue while reviewing #1619. Put together another repro -- this affects all device iterators, not just input.
Repro
#include <thrust/iterator/detail/iterator_category_with_system_and_traversal.h>
#include <thrust/iterator/iterator_categories.h>
#include <thrust/iterator/iterator_traits.h>
//using iterator_category_t = thrust::forward_device_iterator_tag;
using iterator_category_t = thrust::random_access_device_iterator_tag;
using expected_system_t = thrust::device_system_tag;
using trait_system_t = typename thrust::detail::iterator_category_to_system<iterator_category_t>::type;
static_assert(std::is_same<expected_system_t, trait_system_t>::value);
$ nvcc -c repro.cu
repro.cu(12): error: static assertion failed
1 error detected in the compilation of "repro.cu".
Root Cause
Looking into this, it's a bad specialization.
Thrust's system-aware categories are implemented by inheriting iterator_category_with_system_and_traversal, which in turn inherits from Category:
// thrust/iterator/iterator_categories.h
struct random_access_device_iterator_tag
: thrust::detail::iterator_category_with_system_and_traversal<
std::random_access_iterator_tag,
thrust::device_system_tag,
thrust::random_access_traversal_tag
>
{};
// thrust/iterator/detail/iterator_category_with_system_and_traversal.h
template<typename Category, typename System, typename Traversal>
struct iterator_category_with_system_and_traversal
: Category
{
}; // end iterator_category_with_system_and_traversal
iterator_category_to_system is specialized for iterator_category_with_system_and_traversal:
// thrust/iterator/detail/iterator_category_with_system_and_traversal.h
template <typename Category, typename System, typename Traversal>
struct iterator_category_to_system<
iterator_category_with_system_and_traversal<Category, System, Traversal>>
{
typedef System type;
}; // end iterator_category_to_system
However, this specialization is never picked up, since iterator_category_with_system_and_traversal is a base class of the category and thus the unspecialized iterator_category_to_system is used instead:
// thrust/iterator/detail/iterator_category_to_system.h
template <typename Category>
struct iterator_category_to_system
// convertible to host iterator?
: eval_if<
or_<is_convertible<Category, thrust::input_host_iterator_tag>,
is_convertible<Category, thrust::output_host_iterator_tag>>::value,
detail::identity_<thrust::host_system_tag>,
// convertible to device iterator?
eval_if<or_<is_convertible<Category, thrust::input_device_iterator_tag>,
is_convertible<Category,
thrust::output_device_iterator_tag>>::value,
detail::identity_<thrust::device_system_tag>,
// unknown system
detail::identity_<void>> // if device
> // if host
{}; // end iterator_category_to_system
Using thrust::random_access_device_iterator_tag as an example, the above iterator_category_to_system implementation is checking if thrust::random_access_device_iterator_tag is convertible to thrust::input_host_iterator_tag, which is just an alias to std::input_iterator_tag. Since thrust::random_access_device_iterator_tag inheritance looks like this:
thrust::random_access_device_iterator_tag
-> thrust::detail::iterator_category_with_system_and_traversal<
std::random_access_iterator_tag,
thrust::device_system_tag,
thrust::random_access_traversal_tag>
-> std::random_access_iterator_tag (`Category`)
the check for convertibility passes (std::random_access_iterator_tag can convert to std::input_iterator_tag), and iterator_category_to_system ultimately returns that thrust::random_access_device_iterator_tag's system is thrust::host_system_tag.
Possible solutions
It may be as simple as changing the iterator categories to alias:
// Replace:
struct random_access_device_iterator_tag
: thrust::detail::iterator_category_with_system_and_traversal<
std::random_access_iterator_tag,
thrust::device_system_tag,
thrust::random_access_traversal_tag
>
{};
// With:
using random_access_device_iterator_tag =
thrust::detail::iterator_category_with_system_and_traversal<
std::random_access_iterator_tag,
thrust::device_system_tag,
thrust::random_access_traversal_tag>
{};
This would make the specializations on iterator_category_with_system_and_traversal work, but it will change the type system and make symbol names even more unreadable.
Alternatively, we could use SFINAE or some such to fix the iterator_category_to_[system|traversal] specialization so they can work with iterator_category_with_system_and_traversal.
Not sure which is better, we should explore this some more.
Digging in some more, since I was curious how this manages to work for the common case of device_vector::iterator, device_ptr, etc.
Both of these end up defining their iterator_category using iterator_facade_category_impl:
// thrust/iterator/detail/iterator_facade_category.h
template <typename System,
typename Traversal,
typename ValueParam,
typename Reference>
struct iterator_facade_category_impl
{
typedef typename iterator_facade_default_category<System,
Traversal,
ValueParam,
Reference>::type category;
// we must be able to deduce both Traversal & System from category
// otherwise, munge them all together
typedef typename thrust::detail::eval_if<
thrust::detail::and_<
thrust::detail::is_same<
Traversal,
typename thrust::detail::iterator_category_to_traversal<category>::type>,
thrust::detail::is_same<
System,
typename thrust::detail::iterator_category_to_system<category>::type>>::value,
thrust::detail::identity_<category>,
thrust::detail::identity_<
thrust::detail::iterator_category_with_system_and_traversal<category,
System,
Traversal>>>::
type type;
}; // end iterator_facade_category_impl
This ends up directly defining the iterator category to iterator_category_with_system_and_traversal<C, S, T> instead of the predefined thrust::*_device_iterator_tag iterator categories, so the correct specializations of iterator_category_to_system etc get picked up. So thankfully this bug doesn't affect the common usecases of thrust pointers and vector iterators -- only custom iterators using the predefined tags in thrust/iterator/iterator_categories.h are broken.