Merge pull request #688 from illuhad/feature/sycl-2020/std-hash

[SYCL-2020] Add std::hash specialization for reference semantics classes

include/hipSYCL/glue/embedded_pointer.hpp

@@ -101,6 +101,10 @@ struct unique_id {
     return !(a == b);
   }
 
+  std::size_t hipSYCL_hash_code() const {
+    return id[0] ^ id[1];
+  }
+
   unique_id(const unique_id&) = default;
   unique_id& operator=(const unique_id&) = default;
 

include/hipSYCL/runtime/device_id.hpp

@@ -156,7 +156,17 @@ struct hash<hipsycl::rt::device_id>
   std::size_t operator()(const hipsycl::rt::device_id& k) const
   {
     return hash<int>()(static_cast<int>(k.get_backend())) ^ 
-          (hash<int>()(k.get_id()) << 1);
+          (hash<int>()(k.get_id()) << 8);
+  }
+};
+
+template <>
+struct hash<hipsycl::rt::platform_id>
+{
+  std::size_t operator()(const hipsycl::rt::platform_id& p) const
+  {
+    return hash<int>()(static_cast<int>(p.get_backend())) ^ 
+          (hash<int>()(p.get_platform()) << 8);
   }
 };
 

include/hipSYCL/sycl/buffer.hpp

@@ -764,6 +764,10 @@ class buffer : public detail::property_carrying_object
     return !(lhs == rhs);
   }
 
+  std::size_t hipSYCL_hash_code() const {
+    return std::hash<void*>{}(_impl.get());
+  }
+
   // --- The following methods are part the hipSYCL buffer introspection API
   // which is part of the hipSYCL buffer-USM interoperability framework.
 
@@ -1192,4 +1196,18 @@ extract_buffer_range(const buffer<T, dimensions, AllocatorT> &buff) {
 } // sycl
 } // hipsycl
 
+namespace std {
+
+template <typename T, int dimensions,
+          typename AllocatorT>
+struct hash<hipsycl::sycl::buffer<T, dimensions, AllocatorT>>
+{
+  std::size_t
+  operator()(const hipsycl::sycl::buffer<T, dimensions, AllocatorT> &b) const {
+    return b.hipSYCL_hash_code();
+  }
+};
+
+}
+
 #endif

include/hipSYCL/sycl/context.hpp

@@ -148,6 +148,15 @@ class context
     throw unimplemented{"context::get_info() is unimplemented"};
   }
 
+  std::size_t hipSYCL_hash_code() const {
+    return std::hash<void*>{}(_impl.get());
+  }
+
+  friend bool operator ==(const context& lhs, const context& rhs)
+  { return lhs._impl == rhs._impl; }
+
+  friend bool operator!=(const context& lhs, const context &rhs)
+  { return !(lhs == rhs); }
 private:
   void init(async_handler handler) {
     _impl = std::make_shared<context_impl>();
@@ -198,6 +207,17 @@ inline const rt::unique_device_list &extract_context_devices(const context &ctx)
 } // namespace sycl
 } // namespace hipsycl
 
+namespace std {
+
+template <>
+struct hash<hipsycl::sycl::context>
+{
+  std::size_t operator()(const hipsycl::sycl::context& c) const
+  {
+    return c.hipSYCL_hash_code();
+  }
+};
 
+}
 
 #endif

include/hipSYCL/sycl/device.hpp

@@ -260,6 +260,10 @@ class device {
   backend get_backend() const noexcept {
     return _device_id.get_backend();
   }
+
+  std::size_t hipSYCL_hash_code() const {
+    return std::hash<hipsycl::rt::device_id>{}(_device_id);
+  }
 private:
   rt::device_id _device_id;
 
@@ -724,6 +728,17 @@ inline rt::device_id extract_rt_device(const device &d) {
 } // namespace sycl
 } // namespace hipsycl
 
+namespace std {
+
+template <>
+struct hash<hipsycl::sycl::device>
+{
+  std::size_t operator()(const hipsycl::sycl::device& d) const
+  {
+    return d.hipSYCL_hash_code();
+  }
+};
 
+}
 
 #endif

include/hipSYCL/sycl/event.hpp

@@ -40,6 +40,7 @@
 #include "hipSYCL/runtime/dag_node.hpp"
 #include "hipSYCL/runtime/application.hpp"
 #include "hipSYCL/runtime/instrumentation.hpp"
+#include <cstddef>
 
 namespace hipsycl {
 namespace sycl {
@@ -212,6 +213,10 @@ class event {
   friend bool operator !=(const event& lhs, const event& rhs)
   { return !(lhs == rhs); }
 
+  std::size_t hipSYCL_hash_code() const {
+    return std::hash<void*>{}(_node.get());
+  }
+
 private:
 
   rt::dag_node_ptr _node;
@@ -238,4 +243,17 @@ HIPSYCL_SPECIALIZE_GET_INFO(event, reference_count)
 } // namespace sycl
 } // namespace hipsycl
 
+namespace std {
+
+template <>
+struct hash<hipsycl::sycl::event>
+{
+  std::size_t operator()(const hipsycl::sycl::event& e) const
+  {
+    return e.hipSYCL_hash_code();
+  }
+};
+
+}
+
 #endif

include/hipSYCL/sycl/libkernel/accessor.hpp

@@ -815,6 +815,14 @@ class HIPSYCL_EMPTY_BASES accessor
     return !(lhs == rhs);
   }
 
+  std::size_t hipSYCL_hash_code() const {
+    // TODO: This only really guarantees unique hash codes
+    // outside of kernels, on the host side.
+    // Once on device, the UID will contain the device pointer
+    // which might be aliased by multiple accessors.
+    return get_uid().hipSYCL_hash_code();
+  }
+
   HIPSYCL_UNIVERSAL_TARGET
   bool is_placeholder() const noexcept
   {
@@ -1570,6 +1578,10 @@ class host_accessor {
   // const_reverse_iterator crbegin() const noexcept;
   // const_reverse_iterator crend() const noexcept;
 
+  std::size_t hipSYCL_hash_code() const {
+    return _impl.hipSYCL_hash_code();
+  }
+
 private:
   accessor_type _impl;
 };
@@ -1678,6 +1690,10 @@ class accessor<
     return !(lhs == rhs);
   }
 
+  std::size_t hipSYCL_hash_code() const {
+    return _addr;
+  }
+
   [[deprecated("get_size() was removed for SYCL 2020, use byte_size() instead")]]
   HIPSYCL_KERNEL_TARGET
   size_t get_size() const
@@ -1826,4 +1842,30 @@ glue::unique_id get_unique_id(const AccessorType& acc) {
 } // sycl
 } // hipsycl
 
+namespace std {
+
+template <typename dataT, int dimensions, hipsycl::sycl::access_mode accessmode,
+          hipsycl::sycl::target accessTarget,
+          hipsycl::sycl::accessor_variant AccessorVariant>
+struct hash<hipsycl::sycl::accessor<dataT, dimensions, accessmode, accessTarget,
+                                    AccessorVariant>> {
+  std::size_t
+  operator()(const hipsycl::sycl::accessor<dataT, dimensions, accessmode, accessTarget,
+                                    AccessorVariant> &acc) const {
+    return acc.hipSYCL_hash_code();
+  }
+};
+// accessor also covers the local_accessor specialization
+
+template<typename dataT, int dimensions, hipsycl::sycl::access_mode A>
+struct hash<hipsycl::sycl::host_accessor<dataT, dimensions, A>> {
+
+  std::size_t operator()(
+      const hipsycl::sycl::host_accessor<dataT, dimensions, A> &acc) const {
+    return acc.hipSYCL_hash_code();
+  }
+};
+
+}
+
 #endif

include/hipSYCL/sycl/platform.hpp

@@ -131,6 +131,10 @@ class platform {
     return !(lhs == rhs);
   }
 
+  std::size_t hipSYCL_hash_code() const {
+    return std::hash<rt::platform_id>{}(_platform);
+  }
+
 private:
   rt::platform_id _platform;
 };
@@ -170,4 +174,17 @@ inline platform device::get_platform() const  {
 }// namespace sycl
 }// namespace hipsycl
 
+namespace std {
+
+template <>
+struct hash<hipsycl::sycl::platform>
+{
+  std::size_t operator()(const hipsycl::sycl::platform& p) const
+  {
+    return p.hipSYCL_hash_code();
+  }
+};
+
+}
+
 #endif

include/hipSYCL/sycl/queue.hpp

@@ -48,6 +48,7 @@
 #include "info/info.hpp"
 #include "detail/function_set.hpp"
 
+#include <cstddef>
 #include <exception>
 #include <memory>
 #include <mutex>
@@ -847,7 +848,9 @@ class queue : public detail::property_carrying_object
     });  
   }
 
-
+  std::size_t hipSYCL_hash_code() const {
+    return _node_group_id;
+  }
 private:
   template<int Dim>
   void apply_preferred_group_size(const property_list& prop_list, handler& cgh) {
@@ -1099,6 +1102,16 @@ inline auto automatic_require(queue &q,
 }// namespace sycl
 }// namespace hipsycl
 
+namespace std {
+template <>
+struct hash<hipsycl::sycl::queue>
+{
+  std::size_t operator()(const hipsycl::sycl::queue& q) const
+  {
+    return q.hipSYCL_hash_code();
+  }
+};
 
+}
 
 #endif

tests/CMakeLists.txt

@@ -71,6 +71,7 @@ add_executable(sycl_tests
   sycl/math.cpp
   sycl/profiler.cpp
   sycl/reduction.cpp
+  sycl/reference_semantics.cpp
   sycl/sub_group.cpp
   sycl/sycl_test_suite.cpp 
   sycl/usm.cpp

tests/sycl/reference_semantics.cpp

@@ -0,0 +1,46 @@
+/*
+ * This file is part of hipSYCL, a SYCL implementation based on CUDA/HIP
+ *
+ * Copyright (c) 2018-2022 Aksel Alpay and contributors
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice, this
+ *    list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ *    this list of conditions and the following disclaimer in the documentation
+ *    and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR
+ * ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+
+#include "sycl_test_suite.hpp"
+
+#include <unordered_map>
+
+BOOST_FIXTURE_TEST_SUITE(reference_semantics, reset_device_fixture)
+
+
+
+BOOST_AUTO_TEST_CASE(hash) {
+  // Currently only compile-testing
+  std::unordered_map<cl::sycl::device, int> device_map;
+  std::unordered_map<cl::sycl::context, int> context_map;
+  std::unordered_map<cl::sycl::platform, int> platform_map;
+  std::unordered_map<cl::sycl::buffer<int>, int> buffer_map;
+}
+
+
+BOOST_AUTO_TEST_SUITE_END()