[SYCL] Clear usesAssert required by fallback assert

llvm/include/llvm/Support/PropertySetIO.h

@@ -216,7 +216,6 @@ class PropertySetRegistry {
   static constexpr char SYCL_KERNEL_PARAM_OPT_INFO[] = "SYCL/kernel param opt";
   static constexpr char SYCL_PROGRAM_METADATA[] = "SYCL/program metadata";
   static constexpr char SYCL_MISC_PROP[] = "SYCL/misc properties";
-  static constexpr char SYCL_ASSERT_USED[] = "SYCL/assert used";
   static constexpr char SYCL_KERNEL_NAMES[] = "SYCL/kernel names";
   static constexpr char SYCL_EXPORTED_SYMBOLS[] = "SYCL/exported symbols";
   static constexpr char SYCL_IMPORTED_SYMBOLS[] = "SYCL/imported symbols";

llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp

@@ -57,85 +57,6 @@ bool isModuleUsingTsan(const Module &M) {
   return M.getNamedGlobal("__TsanKernelMetadata");
 }
 
-// This function traverses over reversed call graph by BFS algorithm.
-// It means that an edge links some function @func with functions
-// which contain call of function @func. It starts from
-// @StartingFunction and lifts up until it reach all reachable functions,
-// or it reaches some function containing "referenced-indirectly" attribute.
-// If it reaches "referenced-indirectly" attribute than it returns an empty
-// Optional.
-// Otherwise, it returns an Optional containing a list of reached
-// SPIR kernel function's names.
-static std::optional<std::vector<StringRef>> traverseCGToFindSPIRKernels(
-    const std::vector<Function *> &StartingFunctionVec) {
-  std::queue<const Function *> FunctionsToVisit;
-  std::unordered_set<const Function *> VisitedFunctions;
-  for (const Function *FPtr : StartingFunctionVec)
-    FunctionsToVisit.push(FPtr);
-  std::vector<StringRef> KernelNames;
-
-  while (!FunctionsToVisit.empty()) {
-    const Function *F = FunctionsToVisit.front();
-    FunctionsToVisit.pop();
-
-    auto InsertionResult = VisitedFunctions.insert(F);
-    // It is possible that we insert some particular function several
-    // times in functionsToVisit queue.
-    if (!InsertionResult.second)
-      continue;
-
-    for (const auto *U : F->users()) {
-      const CallInst *CI = dyn_cast<const CallInst>(U);
-      if (!CI)
-        continue;
-
-      const Function *ParentF = CI->getFunction();
-
-      if (VisitedFunctions.count(ParentF))
-        continue;
-
-      if (ParentF->hasFnAttribute("referenced-indirectly"))
-        return {};
-
-      if (ParentF->getCallingConv() == CallingConv::SPIR_KERNEL)
-        KernelNames.push_back(ParentF->getName());
-
-      FunctionsToVisit.push(ParentF);
-    }
-  }
-
-  return {std::move(KernelNames)};
-}
-
-static std::vector<StringRef>
-getKernelNamesUsingSpecialFunctions(const Module &M,
-                                    const std::vector<StringRef> &FNames) {
-  std::vector<Function *> SpecialFunctionVec;
-  for (const auto Fn : FNames) {
-    Function *FPtr = M.getFunction(Fn);
-    if (FPtr)
-      SpecialFunctionVec.push_back(FPtr);
-  }
-
-  if (SpecialFunctionVec.size() == 0)
-    return {};
-
-  auto TraverseResult = traverseCGToFindSPIRKernels(SpecialFunctionVec);
-
-  if (TraverseResult.has_value())
-    return std::move(*TraverseResult);
-
-  // Here we reached "referenced-indirectly", so we need to find all kernels and
-  // return them.
-  std::vector<StringRef> SPIRKernelNames;
-  for (const Function &F : M) {
-    if (F.getCallingConv() == CallingConv::SPIR_KERNEL)
-      SPIRKernelNames.push_back(F.getName());
-  }
-
-  return SPIRKernelNames;
-}
-
 // Gets 1- to 3-dimension work-group related information for function Func.
 // Returns an empty vector if not present.
 template <typename T>
@@ -449,13 +370,6 @@ PropSetRegTy computeModuleProperties(const Module &M,
     if (OptLevel != -1)
       PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "optLevel", OptLevel);
   }
-  {
-    std::vector<StringRef> AssertFuncNames{"__devicelib_assert_fail"};
-    std::vector<StringRef> FuncNames =
-        getKernelNamesUsingSpecialFunctions(M, AssertFuncNames);
-    for (const StringRef &FName : FuncNames)
-      PropSet.add(PropSetRegTy::SYCL_ASSERT_USED, FName, true);
-  }
   {
     std::vector<std::pair<StringRef, int>> ArgPos =
         getKernelNamesUsingImplicitLocalMem(M);

llvm/lib/Support/PropertySetIO.cpp

@@ -200,7 +200,6 @@ constexpr char PropertySetRegistry::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES[];
 constexpr char PropertySetRegistry::SYCL_KERNEL_PARAM_OPT_INFO[];
 constexpr char PropertySetRegistry::SYCL_PROGRAM_METADATA[];
 constexpr char PropertySetRegistry::SYCL_MISC_PROP[];
-constexpr char PropertySetRegistry::SYCL_ASSERT_USED[];
 constexpr char PropertySetRegistry::SYCL_KERNEL_NAMES[];
 constexpr char PropertySetRegistry::SYCL_EXPORTED_SYMBOLS[];
 constexpr char PropertySetRegistry::SYCL_IMPORTED_SYMBOLS[];

sycl/include/sycl/queue.hpp

@@ -3817,13 +3817,6 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   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 &,
-                                           const detail::code_location &);
-#endif
-#endif
-
   template <typename CommandGroupFunc, typename PropertiesT>
   friend void ext::oneapi::experimental::detail::submit_impl(
       const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,

sycl/source/detail/compiler.hpp

@@ -55,8 +55,6 @@
 #define __SYCL_PROPERTY_SET_PROGRAM_METADATA "SYCL/program metadata"
 /// PropertySetRegistry::SYCL_MISC_PROP defined in PropertySetIO.h
 #define __SYCL_PROPERTY_SET_SYCL_MISC_PROP "SYCL/misc properties"
-/// PropertySetRegistry::SYCL_ASSERT_USED defined in PropertySetIO.h
-#define __SYCL_PROPERTY_SET_SYCL_ASSERT_USED "SYCL/assert used"
 /// PropertySetRegistry::SYCL_KERNEL_NAMES defined in PropertySetIO.h
 #define __SYCL_PROPERTY_SET_SYCL_KERNEL_NAMES "SYCL/kernel names"
 /// PropertySetRegistry::SYCL_EXPORTED_SYMBOLS defined in PropertySetIO.h

sycl/source/detail/device_binary_image.cpp

@@ -191,7 +191,6 @@ RTDeviceBinaryImage::RTDeviceBinaryImage(sycl_device_binary Bin) {
   DeviceLibReqMask.init(Bin, __SYCL_PROPERTY_SET_DEVICELIB_REQ_MASK);
   DeviceLibMetadata.init(Bin, __SYCL_PROPERTY_SET_DEVICELIB_METADATA);
   KernelParamOptInfo.init(Bin, __SYCL_PROPERTY_SET_KERNEL_PARAM_OPT_INFO);
-  AssertUsed.init(Bin, __SYCL_PROPERTY_SET_SYCL_ASSERT_USED);
   ImplicitLocalArg.init(Bin, __SYCL_PROPERTY_SET_SYCL_IMPLICIT_LOCAL_ARG);
   ProgramMetadata.init(Bin, __SYCL_PROPERTY_SET_PROGRAM_METADATA);
   // Convert ProgramMetadata into the UR format
@@ -517,8 +516,6 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
       naiveMergeBinaryProperties(Imgs, [](const RTDeviceBinaryImage &Img) {
         return Img.getKernelParamOptInfo();
       });
-  auto MergedAssertUsed = naiveMergeBinaryProperties(
-      Imgs, [](const RTDeviceBinaryImage &Img) { return Img.getAssertUsed(); });
   auto MergedDeviceGlobals =
       naiveMergeBinaryProperties(Imgs, [](const RTDeviceBinaryImage &Img) {
         return Img.getDeviceGlobals();
@@ -546,13 +543,12 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
         return Img.getRegisteredKernels();
       });
 
-  std::array<const std::vector<sycl_device_binary_property> *, 11> MergedVecs{
+  std::array<const std::vector<sycl_device_binary_property> *, 10> MergedVecs{
       &MergedSpecConstants,      &MergedSpecConstantsDefaultValues,
-      &MergedKernelParamOptInfo, &MergedAssertUsed,
-      &MergedDeviceGlobals,      &MergedHostPipes,
-      &MergedVirtualFunctions,   &MergedImplicitLocalArg,
-      &MergedKernelNames,        &MergedExportedSymbols,
-      &MergedRegisteredKernels};
+      &MergedKernelParamOptInfo, &MergedDeviceGlobals,
+      &MergedHostPipes,          &MergedVirtualFunctions,
+      &MergedImplicitLocalArg,   &MergedKernelNames,
+      &MergedExportedSymbols,    &MergedRegisteredKernels};
 
   // Exclusive merges.
   auto MergedDeviceLibReqMask =
@@ -672,7 +668,6 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
   CopyPropertiesVec(MergedSpecConstantsDefaultValues,
                     SpecConstDefaultValuesMap);
   CopyPropertiesVec(MergedKernelParamOptInfo, KernelParamOptInfo);
-  CopyPropertiesVec(MergedAssertUsed, AssertUsed);
   CopyPropertiesVec(MergedDeviceGlobals, DeviceGlobals);
   CopyPropertiesVec(MergedHostPipes, HostPipes);
   CopyPropertiesVec(MergedVirtualFunctions, VirtualFunctions);

sycl/source/detail/device_binary_image.hpp

@@ -223,7 +223,6 @@ class RTDeviceBinaryImage {
   const PropertyRange &getKernelParamOptInfo() const {
     return KernelParamOptInfo;
   }
-  const PropertyRange &getAssertUsed() const { return AssertUsed; }
   const PropertyRange &getProgramMetadata() const { return ProgramMetadata; }
   const std::vector<ur_program_metadata_t> &getProgramMetadataUR() const {
     return ProgramMetadataUR;
@@ -259,7 +258,6 @@ class RTDeviceBinaryImage {
   RTDeviceBinaryImage::PropertyRange DeviceLibReqMask;
   RTDeviceBinaryImage::PropertyRange DeviceLibMetadata;
   RTDeviceBinaryImage::PropertyRange KernelParamOptInfo;
-  RTDeviceBinaryImage::PropertyRange AssertUsed;
   RTDeviceBinaryImage::PropertyRange ProgramMetadata;
   RTDeviceBinaryImage::PropertyRange KernelNames;
   RTDeviceBinaryImage::PropertyRange ExportedSymbols;

sycl/source/detail/device_kernel_info.cpp

@@ -24,7 +24,6 @@ DeviceKernelInfo::DeviceKernelInfo(const CompileTimeKernelInfoTy &Info)
 
 void DeviceKernelInfo::init(KernelNameStrRefT KernelName) {
   auto &PM = detail::ProgramManager::getInstance();
-  MUsesAssert = PM.kernelUsesAssert(KernelName);
   MImplicitLocalArgPos = PM.kernelImplicitLocalArgPos(KernelName);
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   MInitialized.store(true);
@@ -78,10 +77,7 @@ FastKernelSubcacheT &DeviceKernelInfo::getKernelSubcache() {
   assertInitialized();
   return MFastKernelSubcache;
 }
-bool DeviceKernelInfo::usesAssert() {
-  assertInitialized();
-  return MUsesAssert;
-}
+
 const std::optional<int> &DeviceKernelInfo::getImplicitLocalArgPos() {
   assertInitialized();
   return MImplicitLocalArgPos;

sycl/source/detail/device_kernel_info.hpp

@@ -108,7 +108,6 @@ class DeviceKernelInfo : public CompileTimeKernelInfoTy {
   void setCompileTimeInfoIfNeeded(const CompileTimeKernelInfoTy &Info);
 
   FastKernelSubcacheT &getKernelSubcache();
-  bool usesAssert();
   const std::optional<int> &getImplicitLocalArgPos();
 
 private:
@@ -119,7 +118,6 @@ class DeviceKernelInfo : public CompileTimeKernelInfoTy {
   std::atomic<bool> MInitialized = false;
 #endif
   FastKernelSubcacheT MFastKernelSubcache;
-  bool MUsesAssert;
   std::optional<int> MImplicitLocalArgPos;
 };
 

sycl/source/detail/kernel_data.hpp

@@ -137,11 +137,6 @@ class KernelData {
 
   void setKernelFunc(void *KernelFuncPtr) { MKernelFuncPtr = KernelFuncPtr; }
 
-  bool usesAssert() const {
-    assert(MDeviceKernelInfoPtr);
-    return MDeviceKernelInfoPtr->usesAssert();
-  }
-
   // Kernel launch properties getter and setters.
   ur_kernel_cache_config_t getKernelCacheConfig() const {
     return MKernelCacheConfig;

sycl/source/detail/program_manager/program_manager.cpp

@@ -1794,14 +1794,6 @@ Managed<ur_program_handle_t> ProgramManager::build(
   return LinkedProg;
 }
 
-void ProgramManager::cacheKernelUsesAssertInfo(const RTDeviceBinaryImage &Img) {
-  const RTDeviceBinaryImage::PropertyRange &AssertUsedRange =
-      Img.getAssertUsed();
-  if (AssertUsedRange.isAvailable())
-    for (const auto &Prop : AssertUsedRange)
-      m_KernelUsesAssert.insert(Prop->Name);
-}
-
 void ProgramManager::cacheKernelImplicitLocalArg(
     const RTDeviceBinaryImage &Img) {
   const RTDeviceBinaryImage::PropertyRange &ImplicitLocalArgRange =
@@ -2044,8 +2036,6 @@ void ProgramManager::addImage(sycl_device_binary RawImg,
     m_KernelNameRefCount[name]++;
   }
 
-  cacheKernelUsesAssertInfo(*Img);
-
   // check if kernel uses sanitizer
   {
     sycl_device_binary_property SanProp = Img->getProperty("sanUsed");
@@ -2116,12 +2106,11 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
 }
 
 template <typename MultimapT, typename KeyT, typename ValT>
-void removeFromMultimapByVal(MultimapT &Map, const KeyT &Key, const ValT &Val,
-                             bool AssertContains = true) {
+void removeFromMultimapByVal(MultimapT &Map, const KeyT &Key, const ValT &Val) {
   auto [RangeBegin, RangeEnd] = Map.equal_range(Key);
   auto It = std::find_if(RangeBegin, RangeEnd,
                          [&](const auto &Pair) { return Pair.second == Val; });
-  if (!AssertContains && It == RangeEnd)
+  if (It == RangeEnd)
     return;
   assert(It != RangeEnd);
   Map.erase(It);
@@ -2233,7 +2222,6 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
       if (--RefCount == 0) {
         // TODO aggregate all these maps into a single one since their entries
         // share lifetime.
-        m_KernelUsesAssert.erase(Name);
         m_KernelImplicitLocalArgPos.erase(Name);
         m_DeviceKernelInfoMap.erase(Name);
         m_KernelNameRefCount.erase(RefCountIt);
@@ -2249,8 +2237,7 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     // unmap loop)
     for (const sycl_device_binary_property &ESProp :
          Img->getExportedSymbols()) {
-      removeFromMultimapByVal(m_ExportedSymbolImages, ESProp->Name, Img,
-                              /*AssertContains*/ false);
+      removeFromMultimapByVal(m_ExportedSymbolImages, ESProp->Name, Img);
     }
 
     m_DeviceImages.erase(DevImgIt);

sycl/source/detail/program_manager/program_manager.hpp

@@ -371,11 +371,6 @@ class ProgramManager {
   ProgramManager();
   ~ProgramManager() = default;
 
-  template <typename NameT>
-  bool kernelUsesAssert(const NameT &KernelName) const {
-    return m_KernelUsesAssert.find(KernelName) != m_KernelUsesAssert.end();
-  }
-
   SanitizerType kernelUsesSanitizer() const { return m_SanitizerFoundInImage; }
 
   std::optional<int>
@@ -412,9 +407,6 @@ class ProgramManager {
   /// Dumps image to current directory
   void dumpImage(const RTDeviceBinaryImage &Img, uint32_t SequenceID = 0) const;
 
-  /// Add info on kernels using assert into cache
-  void cacheKernelUsesAssertInfo(const RTDeviceBinaryImage &Img);
-
   /// Add info on kernels using local arg into cache
   void cacheKernelImplicitLocalArg(const RTDeviceBinaryImage &Img);
 
@@ -528,8 +520,6 @@ class ProgramManager {
   // different types without temporary key_type object creation. This includes
   // standard overloads, such as comparison between std::string and
   // std::string_view or just char*.
-  using KernelUsesAssertSet = std::set<KernelNameStrT, std::less<>>;
-  KernelUsesAssertSet m_KernelUsesAssert;
   std::unordered_map<KernelNameStrT, int> m_KernelImplicitLocalArgPos;
 
   // Map for storing device kernel information. Runtime lookup should be avoided

sycl/source/detail/queue_impl.cpp

@@ -435,13 +435,6 @@ EventImplPtr queue_impl::submit_kernel_scheduler_bypass(
   }
 
   bool DiscardEvent = !EventNeeded && supportsDiscardingPiEvents();
-  if (DiscardEvent) {
-    // Kernel only uses assert if it's non interop one
-    bool KernelUsesAssert =
-        !(KernelImplPtr && KernelImplPtr->isInterop()) && KData.usesAssert();
-    DiscardEvent = !KernelUsesAssert;
-  }
-
   std::shared_ptr<detail::event_impl> ResultEvent =
       DiscardEvent ? nullptr : detail::event_impl::create_device_event(*this);
 

sycl/source/detail/scheduler/commands.cpp

@@ -3346,16 +3346,6 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
     const std::shared_ptr<detail::kernel_impl> &SyclKernel =
         ExecKernel->MSyclKernel;
     KernelNameStrRefT KernelName = ExecKernel->MDeviceKernelInfo.Name;
-
-    if (!EventImpl) {
-      // Kernel only uses assert if it's non interop one
-      bool KernelUsesAssert = (!SyclKernel || SyclKernel->hasSYCLMetadata()) &&
-                              ExecKernel->MDeviceKernelInfo.usesAssert();
-      if (KernelUsesAssert) {
-        EventImpl = MEvent.get();
-      }
-    }
-
     const RTDeviceBinaryImage *BinImage = nullptr;
     if (detail::SYCLConfig<detail::SYCL_JIT_AMDGCN_PTX_KERNELS>::get()) {
       BinImage = retrieveKernelBinary(*MQueue, KernelName);

sycl/unittests/helpers/MockDeviceImage.hpp

@@ -487,15 +487,6 @@ inline MockProperty makeSpecConstant(std::vector<char> &ValData,
   return Prop;
 }
 
-/// Utility function to mark kernel as the one using assert
-inline void setKernelUsesAssert(const std::vector<std::string> &Names,
-                                MockPropertySet &Set) {
-  std::vector<MockProperty> Value;
-  for (const std::string &N : Names)
-    Value.push_back({N, {0, 0, 0, 0}, SYCL_PROPERTY_TYPE_UINT32});
-  Set.insert(__SYCL_PROPERTY_SET_SYCL_ASSERT_USED, std::move(Value));
-}
-
 /// Utility function to add specialization constants to property set.
 ///
 /// This function overrides the default spec constant values.

sycl/unittests/program_manager/Cleanup.cpp

@@ -184,8 +184,6 @@ sycl::unittest::MockDeviceImage generateImage(const std::string &ImageId,
 
   PropSet.insert(__SYCL_PROPERTY_SET_SYCL_VIRTUAL_FUNCTIONS,
                  createVFPropertySet(VirtualFunctions));
-  setKernelUsesAssert(std::vector<std::string>{KernelNames.begin()[0]},
-                      PropSet);
 
   PropSet.insert(__SYCL_PROPERTY_SET_SYCL_IMPLICIT_LOCAL_ARG,
                  createPropertySet(ImplicitLocalArg));