Hi,
I fear that this will be a bit too elaborate of an issue, but at this point I’m just generally looking for ideas on how we may want to re-design our code to work around this issue…
We have a “container type” in our code that itself is templated on the exact vector type that it would use. Like:
template <typename item_t, template <typename> class vector_t >
class container {
...
__host__ __device__
const item_t& at(size_type i) const { return m_items.at(i); }
...
};
We would then use it to create “host” and “device” containers of different types like:
template <typename item_t>
using host_container = container<item_t, host_vector>;
template <typename item_t>
using device_container = container<item_t, device_vector>;
struct foo;
using foo_host_container = host_container<foo>;
using foo_device_container = device_container<foo>;
Here host_vector
and device_vector
are their own stories, but in short host_vector
is practically just std::vector
and device_vector
is something that provides an std::vector
-like API in device code.
At first we felt pretty smart about this code organisation. But as soon as we try to use a “host container” in a .cu
file, we get a warning about making use of a __host__
-only function (std::vector<T>::at
) in a __host__ __device__
function.
So I’m trying to come up with a clever way of how we could specialise some of the functions of our container
type so that they would correctly appear with either just the __host__
or both the __host__
and __device__
qualifiers, depending on the “vector type” template parameter. But I fear that this may not be possible. :frowning:
Am I right in thinking that the following type of code is not allowed?
template <typename T>
struct TestStruct {
__host__ void foo() {...}
};
template <>
__host__ __device__ void TestStruct<int>::foo() {...}
I.e. the __host__
and __device__
qualifiers are part of the function signature, right? So here I’m basically implementing a different function than what the struct declares…
I don’t know if anybody is still with me at this point… But if so, is it indeed not possible to change the __host__
and/or __device__
qualifiers of template functions during template specialisation? It would’ve been pretty cool if these qualifiers could’ve been changed like this. But I guess I can understand if they can’t be…
Cheers,
Attila
P.S. Our code does actually work in its current form. We just get these very annoying warnings from it…
/__w/traccc/traccc/core/include/traccc/edm/container.hpp(209): warning #20014-D: calling a __host__ function from a __host__ __device__ function is not allowed
detected during instantiation of "const item_t traccc::container<header_t, item_t, vector_t, jagged_vector_t, pair_t>::at(const traccc::container<header_t, item_t, vector_t, jagged_vector_t, pair_t>::link_type &) const [with header_t=traccc::geometry_id, item_t=traccc::spacepoint, vector_t=vecmem::vector, jagged_vector_t=vecmem::jagged_vector, pair_t=std::pair]"
/__w/traccc/traccc/core/include/traccc/edm/seed.hpp(53): here