Merge branch 'intel:sycl' into int-buffers

Author: mateuszpn

.github/workflows/sycl-linux-build.yml

@@ -204,7 +204,8 @@ jobs:
           --ci-defaults --use-zstd ${{ inputs.build_configure_extra_args }} \
           -DCMAKE_C_COMPILER_LAUNCHER=ccache \
           -DCMAKE_CXX_COMPILER_LAUNCHER=ccache \
-          -DLLVM_INSTALL_UTILS=ON
+          -DLLVM_INSTALL_UTILS=ON \
+          -DSYCL_UR_FORCE_FETCH_LEVEL_ZERO=ON
     - name: Compile
       id: build
       # Emulate default value for manual dispatch as we've run out of available arguments.

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13425,8 +13425,9 @@ def err_sycl_add_ir_attribute_invalid_filter : Error<
   "initializer list in the first argument of %0 must contain only string "
   "literals">;
 def warn_sycl_old_and_new_kernel_attributes : Warning<
-  "kernel has both attribute %0 and kernel properties; conflicting properties "
-  "are ignored">, InGroup<IgnoredAttributes>;
+  "kernel has both attribute %0 and kernel properties; if the kernel "
+  "properties contains the property \"%1\" it will be ignored">,
+  InGroup<IgnoredAttributes>;
 def err_attribute_argument_is_not_valid : Error<
   "%0 attribute requires integer constant value 0 or 1">;
 

clang/lib/Sema/SemaSYCLDeclAttr.cpp

@@ -3256,14 +3256,19 @@ void SemaSYCL::checkSYCLAddIRAttributesFunctionAttrConflicts(Decl *D) {
   }
 
   // If there are potentially conflicting attributes, we issue a warning.
-  for (const auto *Attr : std::vector<AttributeCommonInfo *>{
-           D->getAttr<SYCLReqdWorkGroupSizeAttr>(),
-           D->getAttr<IntelReqdSubGroupSizeAttr>(),
-           D->getAttr<SYCLWorkGroupSizeHintAttr>(),
-           D->getAttr<SYCLDeviceHasAttr>()})
+  for (const auto [Attr, PotentialConflictProp] :
+       std::vector<std::pair<AttributeCommonInfo *, StringRef>>{
+           {D->getAttr<SYCLReqdWorkGroupSizeAttr>(),
+            "sycl::ext::oneapi::experimental::work_group_size"},
+           {D->getAttr<IntelReqdSubGroupSizeAttr>(),
+            "sycl::ext::oneapi::experimental::sub_group_size"},
+           {D->getAttr<SYCLWorkGroupSizeHintAttr>(),
+            "sycl::ext::oneapi::experimental::work_group_size_hint"},
+           {D->getAttr<SYCLDeviceHasAttr>(),
+            "sycl::ext::oneapi::experimental::device_has"}})
     if (Attr)
       Diag(Attr->getLoc(), diag::warn_sycl_old_and_new_kernel_attributes)
-          << Attr;
+          << Attr << PotentialConflictProp;
 }
 
 void SemaSYCL::handleSYCLRegisteredKernels(Decl *D, const ParsedAttr &A) {

clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp

@@ -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/include/llvm/Support/PropertySetIO.h

@@ -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/SYCLPostLink/ComputeModuleRuntimeInfo.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[];