iterator_category_to_system returns host_system_tag for input_device_iterator_tag instead of device_system_tag

[brunodawagne]

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.

[brycelelbach]

Uh... this seems like it might be bad.

[brycelelbach]

Got a local repro. Is this a regression?

[brycelelbach]

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.

[brycelelbach]

@jaredhoberock, any thoughts?

[brycelelbach]

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

*/

[brycelelbach]

Tracked internally by NVBug 2062266.

[alliepiper]

Just ran across this issue while reviewing NVIDIA/thrust#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.

[alliepiper]

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.