Merge remote-tracking branch 'origin/sycl' into private/asachkov/drop-platform-util

Author: AlexeySachkov

sycl/include/sycl/detail/spirv.hpp

@@ -1225,22 +1225,26 @@ EnableIfGenericShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
 template <typename Group>
 typename std::enable_if_t<
     ext::oneapi::experimental::is_fixed_topology_group_v<Group>>
-ControlBarrier(Group, memory_scope FenceScope, memory_order Order) {
+ControlBarrier(Group, [[maybe_unused]] memory_scope FenceScope,
+               [[maybe_unused]] memory_order Order) {
+#ifdef __SYCL_DEVICE_ONLY__
   __spirv_ControlBarrier(group_scope<Group>::value, getScope(FenceScope),
                          getMemorySemanticsMask(Order) |
                              __spv::MemorySemanticsMask::SubgroupMemory |
                              __spv::MemorySemanticsMask::WorkgroupMemory |
                              __spv::MemorySemanticsMask::CrossWorkgroupMemory);
+#endif
 }
 
 template <typename Group>
 typename std::enable_if_t<
     ext::oneapi::experimental::is_user_constructed_group_v<Group>>
-ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
+ControlBarrier([[maybe_unused]] Group g,
+               [[maybe_unused]] memory_scope FenceScope,
+               [[maybe_unused]] memory_order Order) {
 #if defined(__NVPTX__)
   __nvvm_bar_warp_sync(detail::ExtractMask(detail::GetMask(g))[0]);
-#else
-  (void)g;
+#elif defined(__SYCL_DEVICE_ONLY__)
   // SPIR-V does not define an instruction to synchronize partial groups.
   // However, most (possibly all?) of the current SPIR-V targets execute
   // work-items in lockstep, so we can probably get away with a MemoryBarrier.

sycl/include/sycl/kernel.hpp

@@ -255,8 +255,10 @@ class __SYCL_EXPORT kernel : public detail::OwnerLessBase<kernel> {
 
   ur_native_handle_t getNative() const;
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   __SYCL_DEPRECATED("Use getNative() member function")
   ur_native_handle_t getNativeImpl() const;
+#endif
 
   std::shared_ptr<detail::kernel_impl> impl;
 

sycl/include/sycl/kernel_bundle.hpp

@@ -123,9 +123,14 @@ class __SYCL_EXPORT device_image_plain {
 
   bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept;
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   ur_native_handle_t getNative() const;
-
+#endif
 protected:
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+  ur_native_handle_t getNative() const;
+#endif
+
   std::shared_ptr<device_image_impl> impl;
 
   template <class Obj>
@@ -207,6 +212,9 @@ class device_image : public detail::device_image_plain,
   template <class T>
   friend T detail::createSyclObjFromImpl(
       std::add_lvalue_reference_t<const decltype(T::impl)> ImplObj);
+
+  // To allow calling device_image_plain::getNative()
+  template <bundle_state> friend class kernel_bundle;
 };
 
 namespace detail {

sycl/include/sycl/queue.hpp

@@ -3549,7 +3549,12 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   bool khr_empty() const;
 #endif
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // TODO: to be made private in the next ABI-breaking window
+  __SYCL_DEPRECATED(
+      "This is a non-standard method, use sycl::get_native instead")
   ur_native_handle_t getNative(int32_t &NativeHandleDesc) const;
+#endif
 
   std::optional<event> ext_oneapi_get_last_event() const {
     return static_cast<std::optional<event>>(ext_oneapi_get_last_event_impl());
@@ -3558,6 +3563,10 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   void ext_oneapi_set_external_event(const event &external_event);
 
 private:
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+  ur_native_handle_t getNative(int32_t &NativeHandleDesc) const;
+#endif
+
   std::shared_ptr<detail::queue_impl> impl;
   queue(std::shared_ptr<detail::queue_impl> impl) : impl(impl) {}
 
@@ -3575,6 +3584,10 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   friend auto get_native(const SyclObjectT &Obj)
       -> backend_return_t<BackendName, SyclObjectT>;
 
+  template <backend BackendName>
+  friend auto get_native(const queue &Obj)
+      -> backend_return_t<BackendName, queue>;
+
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 #if __SYCL_USE_FALLBACK_ASSERT
   friend event detail::submitAssertCapture(const queue &, event &,

sycl/source/CMakeLists.txt

@@ -329,7 +329,6 @@ set(SYCL_COMMON_SOURCES
     "queue.cpp"
     "sampler.cpp"
     "stream.cpp"
-    "spirv_ops.cpp"
     "virtual_mem.cpp"
     "detail/memory_pool_impl.cpp"
     "detail/async_alloc.cpp"
@@ -340,6 +339,7 @@ set(SYCL_COMMON_SOURCES
 )
 
 set(SYCL_NON_PREVIEW_SOURCES "${SYCL_COMMON_SOURCES}"
+  "spirv_ops.cpp"
 )
 
 

sycl/source/kernel.cpp

@@ -253,9 +253,9 @@ kernel::kernel(std::shared_ptr<detail::kernel_impl> Impl) : impl(Impl) {}
 
 ur_native_handle_t kernel::getNative() const { return impl->getNative(); }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 ur_native_handle_t kernel::getNativeImpl() const { return impl->getNative(); }
 
-#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 // The following query was deprecated since it doesn't include a way to specify
 // the invdividual dimensions of the work group. All of the contents of this
 // #ifndef block should be removed during the next ABI breaking window.

sycl/source/ld-version-script.txt

@@ -18,10 +18,13 @@
     _ZN10__host_std*;
 
     /* Export SPIR-V built-ins for host device */
+    /* #ifndef __INTEL_PREVIEW_BREAKING_CHANGES */
+    /* TODO: drop those in the next ABI-breaking window */
     _Z23__spirv_GroupWaitEvents*;
     _Z22__spirv_ControlBarrier*;
     _Z21__spirv_MemoryBarrier*;
     _Z20__spirv_ocl_prefetch*;
+    /* #endif // __INTEL_PREVIEW_BREAKING_CHANGES */
 
     /* Export offload image hooks */
     __sycl_register_lib;