Adopt interface chance in HSAKernel. Note that nullptr is passed as the

owning HSAExecutable for the shared-object case - if at some point
HSAKernel will start using its _executable member this will create
issues.

hc2/headers/functions/integer_computational_basis.hpp

@@ -19,4 +19,15 @@ namespace hc2
     {
         return x > T{0};
     }
+
+    template<
+        typename T,
+        typename std::enable_if<std::is_integral<T>{}>::type* = nullptr>
+    inline
+    constexpr
+    bool zero(T x)
+    {
+        return x == T{0};
+    }
+
 }

lib/hsa/mcwamp_hsa.cpp

@@ -2518,7 +2518,8 @@ class HSADevice final : public KalmarDevice
                         if (hsa_executable_symbol_get_info(
                                 *k, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &h) ==
                             HSA_STATUS_SUCCESS) {
-                            kernel = new HSAKernel(str, shortName, *k, h);
+
+                            kernel = new HSAKernel(str, shortName, nullptr, *k, h);
                         }
                     }
                 }