Restore sycl-opt-level attribute functionality after merge

Author: AlexeySachkov

llvm/include/llvm/SYCLLowerIR/SYCLUtils.h

@@ -22,6 +22,7 @@ namespace llvm {
 namespace sycl {
 namespace utils {
 constexpr char ATTR_SYCL_MODULE_ID[] = "sycl-module-id";
+constexpr char ATTR_SYCL_OPTLEVEL[] = "sycl-optlevel";
 
 using CallGraphNodeAction = ::std::function<void(Function *)>;
 using CallGraphFunctionFilter =

llvm/test/tools/sycl-post-link/sycl-opt-level.ll

@@ -8,13 +8,18 @@
 ; RUN: FileCheck %s -input-file=%t.table
 ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-OPT-LEVEL-PROP-0
 ; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-OPT-LEVEL-PROP-1
+; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYM-0
+; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYM-1
 
 ; CHECK: [Code|Properties|Symbols]
 ; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
 ; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym
+; CHECK-EMPTY:
 
-; CHECK-OPT-LEVEL-PROP-0: optLevel=1|0
-; CHECK-OPT-LEVEL-PROP-1: optLevel=1|2
+; CHECK-OPT-LEVEL-PROP-0: optLevel=1|2
+; CHECK-OPT-LEVEL-PROP-1: optLevel=1|0
+; CHECK-SYM-0: _Z3fooii
+; CHECK-SYM-1: _Z3booii
 
 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
 target triple = "spir64-unknown-unknown"

llvm/tools/sycl-post-link/ModuleSplitter.cpp

@@ -43,7 +43,6 @@ constexpr char GLOBAL_SCOPE_NAME[] = "<GLOBAL>";
 constexpr char SYCL_SCOPE_NAME[] = "<SYCL>";
 constexpr char ESIMD_SCOPE_NAME[] = "<ESIMD>";
 constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd";
-constexpr char ATTR_OPT_LEVEL[] = "sycl-optlevel";
 
 bool hasIndirectFunctionsOrCalls(const Module &M) {
   for (const auto &F : M.functions()) {
@@ -595,7 +594,7 @@ void ModuleDesc::dump() const {
   llvm::errs() << "split_module::ModuleDesc[" << Name << "] {\n";
   llvm::errs() << "  ESIMD:" << toString(EntryPoints.Props.HasESIMD)
                << ", SpecConstMet:" << (Props.SpecConstsMet ? "YES" : "NO")
-               << ", OptLevel:" << EntryPoints.getOptLevel() << "\n";
+               << "\n";
   dumpEntryPoints(entries(), EntryPoints.GroupId.c_str(), 1);
   llvm::errs() << "}\n";
 }
@@ -851,6 +850,8 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly,
         ::sycl::kernel_props::ATTR_LARGE_GRF, "large-grf");
     Categorizer.registerListOfIntegersInMetadataSortedRule("sycl_used_aspects");
     Categorizer.registerListOfIntegersInMetadataRule("reqd_work_group_size");
+    Categorizer.registerSimpleStringAttributeRule(
+        sycl::utils::ATTR_SYCL_OPTLEVEL);
     break;
   }
 

llvm/tools/sycl-post-link/ModuleSplitter.h

@@ -55,8 +55,6 @@ struct EntryPointGroup {
   struct Properties {
     // Whether all EPs are ESIMD, SYCL or there are both kinds.
     SyclEsimdSplitStatus HasESIMD = SyclEsimdSplitStatus::SYCL_AND_ESIMD;
-    // front-end opt level for kernel compilation
-    int OptLevel = -1;
     // Scope represented by EPs in a group
     EntryPointsGroupScope Scope = Scope_Global;
 
@@ -66,8 +64,6 @@ struct EntryPointGroup {
                          ? HasESIMD
                          : SyclEsimdSplitStatus::SYCL_AND_ESIMD;
       // Scope remains global
-      // OptLevel is expected to be the same for both merging EPGs
-      assert(OptLevel == Other.OptLevel && "OptLevels are not same");
       return Res;
     }
   };
@@ -92,9 +88,6 @@ struct EntryPointGroup {
     return Props.HasESIMD == SyclEsimdSplitStatus::SYCL_ONLY;
   }
 
-  // Returns opt level
-  int getOptLevel() const { return Props.OptLevel; }
-
   void saveNames(std::vector<std::string> &Dest) const;
   void rebuildFromNames(const std::vector<std::string> &Names, const Module &M);
   void rebuild(const Module &M);
@@ -148,7 +141,6 @@ class ModuleDesc {
 
   bool isESIMD() const { return EntryPoints.isEsimd(); }
   bool isSYCL() const { return EntryPoints.isSycl(); }
-  int getOptLevel() const { return EntryPoints.getOptLevel(); }
 
   const EntryPointSet &entries() const { return EntryPoints.Functions; }
   const EntryPointGroup &getEntryPointGroup() const { return EntryPoints; }

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -40,6 +40,7 @@
 #include "llvm/SYCLLowerIR/HostPipes.h"
 #include "llvm/SYCLLowerIR/LowerInvokeSimd.h"
 #include "llvm/SYCLLowerIR/LowerKernelProps.h"
+#include "llvm/SYCLLowerIR/SYCLUtils.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/FileSystem.h"
 #include "llvm/Support/InitLLVM.h"
@@ -462,9 +463,23 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
     if (HasLargeGRF)
       PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isLargeGRF", true});
   }
-  if (MD.getOptLevel() != -1)
-    PropSet[PropSetRegTy::SYCL_MISC_PROP].insert(
-        {"optLevel", MD.getOptLevel()});
+  {
+    // Handle sycl-optlevel property
+    int OptLevel = -1;
+    for (const Function &F : M.functions()) {
+      if (!F.hasFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL))
+        continue;
+
+      // getAsInteger returns true on error
+      if (!F.getFnAttribute(llvm::sycl::utils::ATTR_SYCL_OPTLEVEL)
+               .getValueAsString()
+               .getAsInteger(10, OptLevel))
+        break;
+    }
+
+    if (OptLevel != -1)
+      PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"optLevel", OptLevel});
+  }
   {
     std::vector<StringRef> FuncNames = getKernelNamesUsingAssert(M);
     for (const StringRef &FName : FuncNames)