Skip to content
This repository was archived by the owner on Mar 21, 2024. It is now read-only.
This repository was archived by the owner on Mar 21, 2024. It is now read-only.

device_reference not stripped before passing values to placeholder actors in some cases #1178

Closed
@smitsyn

Description

@smitsyn

I've faced a bug that I believe is caused by thrust::device_reference<T> in some cases not being stripped to T before being passed to function objects that are composed using placeholders (actors?). The minimal example I could come up with:

#include <thrust/device_vector.h>
#include <thrust/inner_product.h>
#include <thrust/functional.h>

void foo()
{
    using namespace thrust::placeholders;

    thrust::device_vector<float> v1, v2;

    thrust::inner_product(v1.begin(), v1.end(), v2.begin()
        , 0.0f, thrust::plus<float>{}, _1 + 1.0f);
}

Here, the problem is related to the last argument (BinaryOp2). The same can be demonstrated with this function object:

struct erroneous_thing
{
    template <typename T>
    __host__ __device__
        T operator()(const T& lhs, const T& rhs) const {
        return lhs + rhs;
    }
};
//...........
    thrust::inner_product(v1.begin(), v1.end(), v2.begin()
        , 0.0f, thrust::plus<float>{}, erroneous_thing{});

Somehow erroneous_thing is instantiated with T being thrust::device_reference<float>, while it should be float, resulting in a compile error. I haven't faced the problem with any other algorithm I've used; thrust::transform properly strips reference, as well as reduction part of inner_product (e.g. passing erroneous_thing{} as a first parameter is fine).

CUDA toolkit 10.2.

Metadata

Metadata

Assignees

No one assigned

    Labels

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions