Template parameter dependent __device__ qualifier


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…


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

You can use #pragma hd_warning_disable on the function to disable this warning.

Is device_container a RAII wrapper for cudaMalloc* and passed to a kernel, or is it only constructed and used on the device?


Thanks for the hint about #pragma hd_warning_disable. (I didn’t know about this pragma so far.) I don’t want to use it in this particular case, as these headers are used in a lot of places in our code (and since this is templated code, I can’t re-enable the warning after the declaration of the templated class), and I do want to look out for these warnings in other places.

So I just bit the bullet after some further reflection, and split this container type into 3 separate classes. (A base class, and a “device” and “host” class that would both inherit from the base class. Each declaring its functions with the correct qualifiers.)

As for the “device container”, the situation is a bit more elaborate. ;-) It is actually the “host vector” that is an RAII type on top of CUDA memory management. The “device vector” type is a type that can operate on top of global memory that was allocated on the host with a “host vector”.

Since none of this code is secret, you can have a look at the full details of what I did in the end, and what all of these classes look like in reality, on: Container Re-design, main branch (2022.04.06.) by krasznaa · Pull Request #165 · acts-project/traccc · GitHub