[SYCL] Handler exceptions on mutually exclusive operations

sycl/include/CL/sycl/detail/cg.hpp

@@ -94,6 +94,10 @@ namespace detail {
 enum class ExtendedMembersType : unsigned int {
   HANDLER_KERNEL_BUNDLE = 0,
   HANDLER_MEM_ADVICE,
+  // handler_impl is stored in the exended members to avoid breaking ABI.
+  // TODO: This should be made a member of the handler class once ABI can be
+  //       broken.
+  HANDLER_IMPL,
 };
 
 // Holds a pointer to an object of an arbitrary type and an ID value which

sycl/include/CL/sycl/handler.hpp

@@ -1116,6 +1116,10 @@ class __SYCL_EXPORT handler {
     kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc);
   }
 
+  bool setStateExplicitKernel();
+  bool setStateSpecConstSet();
+  bool isStateExplicitKernel() const;
+
   std::shared_ptr<detail::kernel_bundle_impl>
   getOrInsertHandlerKernelBundle(bool Insert) const;
 
@@ -1150,6 +1154,11 @@ class __SYCL_EXPORT handler {
   void set_specialization_constant(
       typename std::remove_reference_t<decltype(SpecName)>::value_type Value) {
 
+    if (!setStateSpecConstSet())
+      throw sycl::exception(make_error_code(errc::invalid),
+                            "Specialization constants cannot be set after "
+                            "explicitly setting the used kernel bundle");
+
     std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImplPtr =
         getOrInsertHandlerKernelBundle(/*Insert=*/true);
 
@@ -1162,6 +1171,11 @@ class __SYCL_EXPORT handler {
   typename std::remove_reference_t<decltype(SpecName)>::value_type
   get_specialization_constant() const {
 
+    if (isStateExplicitKernel())
+      throw sycl::exception(make_error_code(errc::invalid),
+                            "Specialization constants cannot be read after "
+                            "explicitly setting the used kernel bundle");
+
     std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImplPtr =
         getOrInsertHandlerKernelBundle(/*Insert=*/true);
 
@@ -1174,6 +1188,13 @@ class __SYCL_EXPORT handler {
 
   void
   use_kernel_bundle(const kernel_bundle<bundle_state::executable> &ExecBundle) {
+
+    if (!setStateExplicitKernel())
+      throw sycl::exception(
+          make_error_code(errc::invalid),
+          "Kernel bundle cannot be explicitly set after a specialization "
+          "constant has been set");
+
     setHandlerKernelBundle(detail::getSyclObjImpl(ExecBundle));
   }
 

sycl/source/detail/handler_impl.hpp

@@ -0,0 +1,36 @@
+//==---------------- handler_impl.hpp - SYCL handler -----------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <detail/kernel_bundle_impl.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+using KernelBundleImplPtr = std::shared_ptr<detail::kernel_bundle_impl>;
+
+enum class HandlerSubmissionState : std::uint8_t {
+  NO_STATE = 0,
+  EXPLICIT_KERNEL_BUNDLE_STATE,
+  SPEC_CONST_SET_STATE,
+};
+
+class handler_impl {
+public:
+  handler_impl() = default;
+
+  /// Registers mutually exclusive submission states.
+  HandlerSubmissionState MSubmissionState =
+      detail::HandlerSubmissionState::NO_STATE;
+};
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/handler.cpp

@@ -17,6 +17,7 @@
 #include <CL/sycl/stream.hpp>
 #include <detail/config.hpp>
 #include <detail/global_handler.hpp>
+#include <detail/handler_impl.hpp>
 #include <detail/kernel_bundle_impl.hpp>
 #include <detail/kernel_impl.hpp>
 #include <detail/queue_impl.hpp>
@@ -27,8 +28,124 @@ namespace sycl {
 
 handler::handler(std::shared_ptr<detail::queue_impl> Queue, bool IsHost)
     : MQueue(std::move(Queue)), MIsHost(IsHost) {
-  MSharedPtrStorage.emplace_back(
-      std::make_shared<std::vector<detail::ExtendedMemberT>>());
+  // Create extended members and insert handler_impl
+  // TODO: When allowed to break ABI the handler_impl should be made a member
+  //       of the handler class.
+  auto ExtendedMembers =
+      std::make_shared<std::vector<detail::ExtendedMemberT>>();
+  detail::ExtendedMemberT HandlerImplMember = {
+      detail::ExtendedMembersType::HANDLER_IMPL,
+      std::make_shared<detail::handler_impl>()};
+  ExtendedMembers->push_back(std::move(HandlerImplMember));
+  MSharedPtrStorage.push_back(std::move(ExtendedMembers));
+}
+
+/// Gets the handler_impl at the start of the extended members.
+// detail::GlobalHandler::instance().getHandlerExtendedMembersMutex() must
+// be held when calling this function.
+std::shared_ptr<detail::handler_impl>
+getHandlerImpl(const std::shared_ptr<std::vector<detail::ExtendedMemberT>>
+                   &ExtendedMembersVec) {
+  assert(ExtendedMembersVec->size() > 0);
+
+  auto HandlerImplMember = (*ExtendedMembersVec)[0];
+
+  assert(detail::ExtendedMembersType::HANDLER_IMPL == HandlerImplMember.MType);
+
+  return std::static_pointer_cast<detail::handler_impl>(
+      HandlerImplMember.MData);
+}
+
+// Common implementation for getting/inserting handler kernel bundle.
+// detail::GlobalHandler::instance().getHandlerExtendedMembersMutex() must
+// be held when calling this function.
+std::shared_ptr<detail::kernel_bundle_impl>
+getOrInsertHandlerKernelBundleCommon(
+    const std::shared_ptr<std::vector<detail::ExtendedMemberT>>
+        &ExtendedMembersVec,
+    const std::shared_ptr<detail::queue_impl> &Queue, bool Insert) {
+  // Look for the kernel bundle in extended members
+  std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImpPtr;
+  for (const detail::ExtendedMemberT &EMember : *ExtendedMembersVec)
+    if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE == EMember.MType) {
+      KernelBundleImpPtr =
+          std::static_pointer_cast<detail::kernel_bundle_impl>(EMember.MData);
+      break;
+    }
+
+  // No kernel bundle yet, create one
+  if (!KernelBundleImpPtr && Insert) {
+    KernelBundleImpPtr = detail::getSyclObjImpl(
+        get_kernel_bundle<bundle_state::input>(Queue->get_context()));
+    if (KernelBundleImpPtr->empty()) {
+      KernelBundleImpPtr = detail::getSyclObjImpl(
+          get_kernel_bundle<bundle_state::executable>(Queue->get_context()));
+    }
+
+    detail::ExtendedMemberT EMember = {
+        detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE, KernelBundleImpPtr};
+    ExtendedMembersVec->push_back(EMember);
+  }
+
+  return KernelBundleImpPtr;
+}
+
+// If the submission state is SPEC_CONST_SET_STATE this function returns false.
+// Otherwise it sets the submission state to EXPLICIT_KERNEL_BUNDLE_STATE and
+// returns true.
+bool handler::setStateExplicitKernel() {
+  std::lock_guard<std::mutex> Lock(
+      detail::GlobalHandler::instance().getHandlerExtendedMembersMutex());
+
+  assert(!MSharedPtrStorage.empty());
+
+  std::shared_ptr<std::vector<detail::ExtendedMemberT>> ExendedMembersVec =
+      detail::convertToExtendedMembers(MSharedPtrStorage[0]);
+
+  auto HandlerImpl = getHandlerImpl(ExendedMembersVec);
+  if (HandlerImpl->MSubmissionState ==
+      detail::HandlerSubmissionState::SPEC_CONST_SET_STATE)
+    return false;
+  HandlerImpl->MSubmissionState =
+      detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE;
+  return true;
+}
+
+// If the submission state is EXPLICIT_KERNEL_BUNDLE_STATE this function returns
+// false. Otherwise it sets the submission state to SPEC_CONST_SET_STATE and
+// returns true.
+bool handler::setStateSpecConstSet() {
+  std::lock_guard<std::mutex> Lock(
+      detail::GlobalHandler::instance().getHandlerExtendedMembersMutex());
+
+  assert(!MSharedPtrStorage.empty());
+
+  std::shared_ptr<std::vector<detail::ExtendedMemberT>> ExendedMembersVec =
+      detail::convertToExtendedMembers(MSharedPtrStorage[0]);
+
+  auto HandlerImpl = getHandlerImpl(ExendedMembersVec);
+  if (HandlerImpl->MSubmissionState ==
+      detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE)
+    return false;
+  HandlerImpl->MSubmissionState =
+      detail::HandlerSubmissionState::SPEC_CONST_SET_STATE;
+  return true;
+}
+
+// Returns true if the submission state is EXPLICIT_KERNEL_BUNDLE_STATE and
+// false otherwise.
+bool handler::isStateExplicitKernel() const {
+  std::lock_guard<std::mutex> Lock(
+      detail::GlobalHandler::instance().getHandlerExtendedMembersMutex());
+
+  assert(!MSharedPtrStorage.empty());
+
+  std::shared_ptr<std::vector<detail::ExtendedMemberT>> ExendedMembersVec =
+      detail::convertToExtendedMembers(MSharedPtrStorage[0]);
+
+  auto HandlerImpl = getHandlerImpl(ExendedMembersVec);
+  return HandlerImpl->MSubmissionState ==
+         detail::HandlerSubmissionState::EXPLICIT_KERNEL_BUNDLE_STATE;
 }
 
 // Returns a shared_ptr to kernel_bundle stored in the extended members vector.
@@ -43,12 +160,11 @@ handler::getOrInsertHandlerKernelBundle(bool Insert) const {
 
   assert(!MSharedPtrStorage.empty());
 
-  std::shared_ptr<std::vector<detail::ExtendedMemberT>> ExendedMembersVec =
+  std::shared_ptr<std::vector<detail::ExtendedMemberT>> ExtendedMembersVec =
       detail::convertToExtendedMembers(MSharedPtrStorage[0]);
-
   // Look for the kernel bundle in extended members
   std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImpPtr;
-  for (const detail::ExtendedMemberT &EMember : *ExendedMembersVec)
+  for (const detail::ExtendedMemberT &EMember : *ExtendedMembersVec)
     if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE == EMember.MType) {
       KernelBundleImpPtr =
           std::static_pointer_cast<detail::kernel_bundle_impl>(EMember.MData);
@@ -66,8 +182,7 @@ handler::getOrInsertHandlerKernelBundle(bool Insert) const {
 
     detail::ExtendedMemberT EMember = {
         detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE, KernelBundleImpPtr};
-
-    ExendedMembersVec->push_back(EMember);
+    ExtendedMembersVec->push_back(EMember);
   }
 
   return KernelBundleImpPtr;
@@ -85,16 +200,18 @@ void handler::setHandlerKernelBundle(
   std::shared_ptr<std::vector<detail::ExtendedMemberT>> ExendedMembersVec =
       detail::convertToExtendedMembers(MSharedPtrStorage[0]);
 
-  for (detail::ExtendedMemberT &EMember : *ExendedMembersVec)
+  // Look for kernel bundle in extended members and overwrite it.
+  for (detail::ExtendedMemberT &EMember : *ExendedMembersVec) {
     if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE == EMember.MType) {
       EMember.MData = NewKernelBundleImpPtr;
       return;
     }
+  }
 
+  // Kernel bundle was set found so we add it.
   detail::ExtendedMemberT EMember = {
       detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE,
       NewKernelBundleImpPtr};
-
   ExendedMembersVec->push_back(EMember);
 }
 

sycl/test/abi/sycl_symbols_linux.dump

@@ -3920,7 +3920,9 @@ _ZN2cl4sycl7handler18ext_oneapi_barrierERKSt6vectorINS0_5eventESaIS3_EE
 _ZN2cl4sycl7handler18extractArgsAndReqsEv
 _ZN2cl4sycl7handler20DisableRangeRoundingEv
 _ZN2cl4sycl7handler20associateWithHandlerEPNS0_6detail16AccessorBaseHostENS0_6access6targetE
+_ZN2cl4sycl7handler20setStateSpecConstSetEv
 _ZN2cl4sycl7handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE
+_ZN2cl4sycl7handler22setStateExplicitKernelEv
 _ZN2cl4sycl7handler24GetRangeRoundingSettingsERmS2_S2_
 _ZN2cl4sycl7handler28extractArgsAndReqsFromLambdaEPcmPKNS0_6detail19kernel_param_desc_tE
 _ZN2cl4sycl7handler28extractArgsAndReqsFromLambdaEPcmPKNS0_6detail19kernel_param_desc_tEb
@@ -4263,6 +4265,7 @@ _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE4225EEENS3_12param_traitsIS4_XT
 _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE4228EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE65552EEENS3_12param_traitsIS4_XT_EE11return_typeEv
 _ZNK2cl4sycl7context9getNativeEv
+_ZNK2cl4sycl7handler21isStateExplicitKernelEv
 _ZNK2cl4sycl7handler30getOrInsertHandlerKernelBundleEb
 _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
 _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEb

sycl/test/abi/sycl_symbols_windows.dump

@@ -2331,6 +2331,7 @@
 ?isInterop@SYCLMemObjT@detail@sycl@cl@@QEBA_NXZ
 ?isOutOfRange@detail@sycl@cl@@YA_NV?$vec@H$03@23@W4addressing_mode@23@V?$range@$02@23@@Z
 ?isPathPresent@OSUtil@detail@sycl@cl@@SA_NAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z
+?isStateExplicitKernel@handler@sycl@cl@@AEBA_NXZ
 ?isValidModeForDestinationAccessor@handler@sycl@cl@@CA_NW4mode@access@23@@Z
 ?isValidModeForSourceAccessor@handler@sycl@cl@@CA_NW4mode@access@23@@Z
 ?isValidTargetForExplicitOp@handler@sycl@cl@@CA_NW4target@access@23@@Z
@@ -3812,6 +3813,8 @@
 ?setPitches@?$image_impl@$00@detail@sycl@cl@@AEAAXXZ
 ?setPitches@?$image_impl@$01@detail@sycl@cl@@AEAAXXZ
 ?setPitches@?$image_impl@$02@detail@sycl@cl@@AEAAXXZ
+?setStateExplicitKernel@handler@sycl@cl@@AEAA_NXZ
+?setStateSpecConstSet@handler@sycl@cl@@AEAA_NXZ
 ?setType@handler@sycl@cl@@AEAAXW4CGTYPE@CG@detail@23@@Z
 ?set_final_data@SYCLMemObjT@detail@sycl@cl@@QEAAX$$T@Z
 ?set_final_data_from_storage@SYCLMemObjT@detail@sycl@cl@@QEAAXXZ

sycl/unittests/SYCL2020/CMakeLists.txt

@@ -4,7 +4,7 @@ set(CMAKE_CXX_EXTENSIONS OFF)
 set(LLVM_REQUIRES_EH 1)
 add_sycl_unittest(SYCL2020Tests OBJECT
   GetNativeOpenCL.cpp
-  SpecConstDefaultValues.cpp
+  SpecializationConstant.cpp
   KernelBundle.cpp
   KernelID.cpp
 )