• HIP is more restrictive in enforcing restrictions

HIP is more restrictive in enforcing restrictions

The language specification for HIP and CUDA forbid calling a __device__ function in a __host__ context. In practice, you may observe differences in the strictness of this restriction, with HIP exhibiting a tighter adherence to the specification and thus less tolerant of infringing code. The solution is to ensure that all functions which are called in a __device__ context are correctly annotated to reflect it. An interesting case where these differences emerge is shown below. This relies on a the common C++ Member Detector idiom, as it would be implemented pre C++11):

```c++ #include <cassert> #include <type_traits>

struct aye { bool a1; }; struct nay { bool a[2]; };

// Dual restriction is necessary in HIP if the detector is to work for // device contexts as well as host ones. NVCC is less strict. template<typename T> host device const T& cref_t();

template<typename T> struct Has_call_operator { // Dual restriction is necessary in HIP if the detector is to work for // device contexts as well as host ones. NVCC is less strict. template<typename C> host device static aye test( C const *, typename std::enable_if< (sizeof(cref_t<C>().operator()()) > 0)>::type* = nullptr); static nay test(...);

enum { value = sizeof(test(static_cast<T*>(0))) == sizeof(aye) }; };

template<typename T, typename U, bool callable = has_call_operator<U>::value> struct Wrapper { template<typename V> V f() const { return T{1}; } };

template<typename T, typename U> struct Wrapper<T, U, true> { template<typename V> V f() const { return T{10}; } };

// This specialisation will yield a compile-time error, if selected. template<typename T, typename U> struct Wrapper<T, U, false> {};

template<typename T> struct Functor;

template<> struct Functor<float> { device float operator()() const { return 42.0f; } };

device void this_will_not_compile_if_detector_is_not_marked_device() { float f = Wrapper<float, Functor<float>>().f<float>(); }

host void this_will_not_compile_if_detector_is_marked_device_only() { float f = Wrapper<float, Functor<float>>().f<float>(); } ```