Add kernel property to propagate compile options to backend

Signed-off-by: Arvind Sudarsanam <[email protected]>

Author: asudarsa

clang/include/clang/Basic/Attr.td

@@ -1309,12 +1309,12 @@ def SYCLType: InheritableAttr {
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
                             "private_memory", "aspect", "annotated_ptr", "annotated_arg",
-                            "stream", "sampler"],
+                            "compile_options", "stream", "sampler"],
                            ["accessor", "local_accessor", "spec_constant",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
                             "private_memory", "aspect", "annotated_ptr", "annotated_arg",
-                            "stream", "sampler"]>];
+                            "compile_options", "stream", "sampler"]>];
   // Only used internally by SYCL implementation
   let Documentation = [InternalOnly];
 }
@@ -1328,6 +1328,15 @@ def SYCLDeviceHas : InheritableAttr {
   let SupportsNonconformingLambdaSyntax = 1;
 }
 
+def SYCLDeviceCompileOptions : InheritableAttr {
+  let Spellings = [CXX11<"sycl", "device_compile_options">];
+  let Subjects = SubjectList<[Function], ErrorDiag>;
+  let Args = [VariadicExprArgument<"DeviceCompileOptions">];
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  // Only used internally by SYCL implementation
+  let Documentation = [InternalOnly];
+}
+
 def SYCLUsesAspects : InheritableAttr {
   let Spellings = [CXX11<"__sycl_detail__", "__uses_aspects__">];
   let Subjects = SubjectList<[CXXRecord, Function], ErrorDiag>;

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -582,11 +582,23 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD,
     && !FD->hasAttr<SYCLDeviceAttr>())
     return;
 
-  // TODO Module identifier is not reliable for this purpose since two modules
-  // can have the same ID, needs improvement
-  if (getLangOpts().SYCLIsDevice)
+  
+  if (getLangOpts().SYCLIsDevice) {
+    // TODO Module identifier is not reliable for this purpose since two modules
+    // can have the same ID, needs improvement
     Fn->addFnAttr("sycl-module-id", Fn->getParent()->getModuleIdentifier());
-
+    int SYCLDeviceCompileOptLevel;
+    switch (CGM.getCodeGenOpts().OptimizationLevel) {
+    default:
+      llvm_unreachable("Invalid optimization level!");
+    case 0:
+    case 1:
+    case 2:
+    case 3:
+      SYCLDeviceCompileOptLevel = CGM.getCodeGenOpts().OptimizationLevel;
+    }
+    Fn->addFnAttr("sycl-device-compile-optlevel", std::to_string(SYCLDeviceCompileOptLevel));
+  }
   llvm::LLVMContext &Context = getLLVMContext();
 
   if (FD->hasAttr<OpenCLKernelAttr>() || FD->hasAttr<CUDAGlobalAttr>())

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

@@ -208,6 +208,10 @@ attributeToExecModeMetadata(Module &M, const Attribute &Attr) {
                                             MDNode::get(Ctx, MD));
   }
 
+  if (AttrKindStr == "sycl-device-compile-options") {
+    auto Opt = Attr.getValueAsString();
+    llvm::errs() << "ARV: Opt is -O" << Opt << "\n";
+  }
   return None;
 }
 

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

@@ -89,6 +89,7 @@ cl::OptionCategory PostLinkCat{"sycl-post-link options"};
 // clang/lib/Driver/Driver.cpp, sycl-post-link.cpp, ClangOffloadWrapper.cpp
 constexpr char COL_CODE[] = "Code";
 constexpr char COL_SYM[] = "Symbols";
+constexpr char COL_OPTS[] = "Options";
 constexpr char COL_PROPS[] = "Properties";
 
 // InputFilename - The filename to read from.
@@ -215,10 +216,11 @@ struct GlobalBinImageProps {
   bool EmitDeviceGlobalPropSet;
 };
 
-struct IrPropSymFilenameTriple {
+struct IrPropSymFilenameQuad {
   std::string Ir;
   std::string Prop;
   std::string Sym;
+  std::string Opt;
 };
 
 void writeToFile(const std::string &Filename, const std::string &Content) {
@@ -466,6 +468,44 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
   return SCFile;
 }
 
+std::string getOptString(module_split::ModuleDesc &MD) {
+  auto &M = MD.getModule();
+  // Process all properties on kernels.
+  for (Function &F : M) {
+    // Only consider kernels.
+    if (F.getCallingConv() != CallingConv::SPIR_KERNEL)
+      continue;
+
+    SmallVector<Metadata *, 8> MDOps;
+    SmallVector<std::pair<std::string, MDNode *>, 8> NamedMDOps;
+    for (const Attribute &Attr : F.getAttributes().getFnAttrs()) {
+      // Currently, only string attributes are supported
+      if (!Attr.isStringAttribute())
+        continue;
+      StringRef AttrKindStr = Attr.getKindAsString();
+      if (AttrKindStr == "sycl-device-compile-optlevel") {
+        auto Opt = "-O" + Attr.getValueAsString();
+        llvm::errs() << "ARV: Opt is " << Opt << "\n";
+        return Opt.str();
+      }
+    }
+  }
+  return "";
+}
+
+std::string saveModuleOptions(module_split::ModuleDesc &MD,
+                                 const std::string &Opts, int I,
+                                 StringRef Suff) {
+  std::error_code EC;
+  std::string SCFile = makeResultFileName(".opt", I, Suff);
+  raw_fd_ostream SCOut(SCFile, EC);
+  checkError(EC, "error opening file '" + SCFile + "'");
+  SCOut << Opts;
+
+  return SCFile;
+}
+
+
 // Saves specified collection of symbols to a file.
 std::string saveModuleSymbolTable(const module_split::EntryPointSet &Es, int I,
                                   StringRef Suffix) {
@@ -570,11 +610,11 @@ StringRef getModuleSuffix(const module_split::ModuleDesc &MD) {
 // @param IRFilename filename of already available IR component. If not empty,
 //   IR component saving is skipped, and this file name is recorded as such in
 //   the result.
-// @return a triple of files where IR, Property and Symbols components of the
-//   Module descriptor are written respectively.
-IrPropSymFilenameTriple saveModule(module_split::ModuleDesc &MD, int I,
+// @return a quadruple of files where IR, Property, Symbols and Opts components
+//   of the Module descriptor are written respectively.
+IrPropSymFilenameQuad saveModule(module_split::ModuleDesc &MD, int I,
                                    StringRef IRFilename = "") {
-  IrPropSymFilenameTriple Res;
+  IrPropSymFilenameQuad Res;
   StringRef Suffix = getModuleSuffix(MD);
 
   if (!IRFilename.empty()) {
@@ -587,6 +627,9 @@ IrPropSymFilenameTriple saveModule(module_split::ModuleDesc &MD, int I,
                                EmitExportedSymbols, DeviceGlobals};
   Res.Prop = saveModuleProperties(MD, Props, I, Suffix);
 
+  std::string Opts = getOptString(MD);
+  Res.Opt = saveModuleOptions(MD, Opts, I, Suffix);
+
   if (DoSymGen) {
     // save the names of the entry points - the symbol table
     Res.Sym = saveModuleSymbolTable(MD.entries(), I, Suffix);
@@ -631,20 +674,37 @@ bool processSpecConstants(module_split::ModuleDesc &MD) {
   return MD.Props.SpecConstsMet;
 }
 
-constexpr int MAX_COLUMNS_IN_FILE_TABLE = 3;
+constexpr int MAX_COLUMNS_IN_FILE_TABLE = 4;
 
+#if 0
 void addTableRow(util::SimpleTable &Table,
-                 const IrPropSymFilenameTriple &RowData) {
+                 const IrPropSymFilenameQuad &RowData) {
   SmallVector<StringRef, MAX_COLUMNS_IN_FILE_TABLE> Row;
 
+  for (const std::string *S : {&RowData.Ir, &RowData.Prop, &RowData.Opt, &RowData.Sym}) {
+    if (!S->empty()) {
+      Row.push_back(StringRef(*S));
+    }
+  }
+  llvm::errs() << "ARV: " << static_cast<size_t>(Table.getNumColumns()) << "," << Row.size() <<"\n";
+  assert(static_cast<size_t>(Table.getNumColumns()) == Row.size());
+  Table.addRow(Row);
+}
+#else
+void addTableRow(util::SimpleTable &Table,
+                 const IrPropSymFilenameQuad &RowData) {
+  SmallVector<StringRef, 3> Row;
+
   for (const std::string *S : {&RowData.Ir, &RowData.Prop, &RowData.Sym}) {
     if (!S->empty()) {
       Row.push_back(StringRef(*S));
     }
   }
+  //llvm::errs() << "ARV: " << static_cast<size_t>(Table.getNumColumns()) << "," << Row.size() <<"\n";
   assert(static_cast<size_t>(Table.getNumColumns()) == Row.size());
   Table.addRow(Row);
 }
+#endif
 
 // Removes the global variable "llvm.used" and returns true on success.
 // "llvm.used" is a global constant array containing references to kernels
@@ -693,8 +753,13 @@ static bool removeSYCLKernelsConstRefArray(Module &M) {
 std::unique_ptr<util::SimpleTable>
 processInputModule(std::unique_ptr<Module> M) {
   // Construct the resulting table which will accumulate all the outputs.
+#if 0
+  SmallVector<StringRef, MAX_COLUMNS_IN_FILE_TABLE> ColumnTitles{
+      StringRef(COL_CODE), StringRef(COL_PROPS), StringRef(COL_OPTS)};
+#else
   SmallVector<StringRef, MAX_COLUMNS_IN_FILE_TABLE> ColumnTitles{
       StringRef(COL_CODE), StringRef(COL_PROPS)};
+#endif
 
   if (DoSymGen) {
     ColumnTitles.push_back(COL_SYM);
@@ -869,7 +934,7 @@ processInputModule(std::unique_ptr<Module> M) {
                   "have been made\n";
       }
       for (module_split::ModuleDesc &IrMD : MMs) {
-        IrPropSymFilenameTriple T = saveModule(IrMD, ID, OutIRFileName);
+        IrPropSymFilenameQuad T = saveModule(IrMD, ID, OutIRFileName);
         addTableRow(*Table, T);
       }
     }

sycl/include/sycl/compile_options.hpp

@@ -0,0 +1,22 @@
+//==---- compile_options.hpp - SYCL compile options Enums -----*- C++ -*---==//
+//
+// 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 <sycl/detail/defines.hpp>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+
+#define __SYCL_COMPILE_OPTION(OPTION, ID) OPTION = ID,
+enum class __SYCL_TYPE(compile_options) compile_options {
+#include <sycl/info/compile_options.def>
+};
+#undef __SYCL_COMPILE_OPTION
+
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl

sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp

@@ -9,6 +9,7 @@
 #pragma once
 
 #include <sycl/aspects.hpp>
+#include <sycl/compile_options.hpp>
 #include <sycl/ext/oneapi/properties/property.hpp>
 #include <sycl/ext/oneapi/properties/property_utils.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
@@ -56,6 +57,12 @@ struct device_has_key {
                                  std::integral_constant<aspect, Aspects>...>;
 };
 
+struct device_compile_options_key {
+  template <compile_options... CompileOptions>
+  using value_t = property_value<device_compile_options_key,
+                                 std::integral_constant<compile_options, CompileOptions>...>;
+};
+
 template <size_t Dim0, size_t... Dims>
 struct property_value<work_group_size_key, std::integral_constant<size_t, Dim0>,
                       std::integral_constant<size_t, Dims>...> {
@@ -108,6 +115,14 @@ struct property_value<device_has_key,
   static constexpr std::array<aspect, sizeof...(Aspects)> value{Aspects...};
 };
 
+template <compile_options... CompileOptions>
+struct property_value<device_compile_options_key,
+                      std::integral_constant<compile_options, CompileOptions>...> {
+  using key_t = device_compile_options_key;
+  static constexpr std::array<compile_options,
+                              sizeof...(CompileOptions)> value{CompileOptions...};
+};
+
 template <size_t Dim0, size_t... Dims>
 inline constexpr work_group_size_key::value_t<Dim0, Dims...> work_group_size;
 
@@ -121,11 +136,15 @@ inline constexpr sub_group_size_key::value_t<Size> sub_group_size;
 template <aspect... Aspects>
 inline constexpr device_has_key::value_t<Aspects...> device_has;
 
+template <compile_options... CompileOptions>
+inline constexpr device_compile_options_key::value_t<CompileOptions...> device_compile_options;
+
 template <> struct is_property_key<work_group_size_key> : std::true_type {};
 template <>
 struct is_property_key<work_group_size_hint_key> : std::true_type {};
 template <> struct is_property_key<sub_group_size_key> : std::true_type {};
 template <> struct is_property_key<device_has_key> : std::true_type {};
+template <> struct is_property_key<device_compile_options_key> : std::true_type {};
 
 namespace detail {
 template <> struct PropertyToKind<work_group_size_key> {
@@ -140,6 +159,9 @@ template <> struct PropertyToKind<sub_group_size_key> {
 template <> struct PropertyToKind<device_has_key> {
   static constexpr PropKind Kind = PropKind::DeviceHas;
 };
+template <> struct PropertyToKind<device_compile_options_key> {
+  static constexpr PropKind Kind = PropKind::DeviceCompileOptions;
+};
 
 template <>
 struct IsCompileTimeProperty<work_group_size_key> : std::true_type {};
@@ -148,6 +170,7 @@ struct IsCompileTimeProperty<work_group_size_hint_key> : std::true_type {};
 template <>
 struct IsCompileTimeProperty<sub_group_size_key> : std::true_type {};
 template <> struct IsCompileTimeProperty<device_has_key> : std::true_type {};
+template <> struct IsCompileTimeProperty<device_compile_options_key> : std::true_type {};
 
 template <size_t Dim0, size_t... Dims>
 struct PropertyMetaInfo<work_group_size_key::value_t<Dim0, Dims...>> {
@@ -170,6 +193,12 @@ struct PropertyMetaInfo<device_has_key::value_t<Aspects...>> {
   static constexpr const char *value =
       SizeListToStr<static_cast<size_t>(Aspects)...>::value;
 };
+template <compile_options... CompileOptions>
+struct PropertyMetaInfo<device_compile_options_key::value_t<CompileOptions...>> {
+  static constexpr const char *name = "sycl-device-compile-options";
+  static constexpr const char *value =
+      SizeListToStr<static_cast<size_t>(CompileOptions)...>::value;
+};
 
 template <typename T, typename = void>
 struct HasKernelPropertiesGetMethod : std::false_type {};

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -173,8 +173,9 @@ enum PropKind : uint32_t {
   WorkGroupSizeHint = 7,
   SubGroupSize = 8,
   DeviceHas = 9,
+  DeviceCompileOptions = 10,
   // PropKindSize must always be the last value.
-  PropKindSize = 10,
+  PropKindSize = 11,
 };
 
 // This trait must be specialized for all properties and must have a unique

sycl/include/sycl/info/compile_options.def

@@ -0,0 +1,4 @@
+__SYCL_COMPILE_OPTION(O0, 0)
+__SYCL_COMPILE_OPTION(O1, 1)
+__SYCL_COMPILE_OPTION(O2, 2)
+__SYCL_COMPILE_OPTION(O3, 3)

sycl/include/sycl/sycl.hpp

@@ -19,6 +19,7 @@
 #endif
 #include <sycl/buffer.hpp>
 #include <sycl/builtins.hpp>
+#include <sycl/compile_options.hpp>
 #include <sycl/context.hpp>
 #include <sycl/define_vendors.hpp>
 #include <sycl/device.hpp>

sycl/test/extensions/properties/prop_compile_options.cpp

@@ -0,0 +1,15 @@
+// RUN: %clangxx -O0 -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR
+// RUN: %clangxx -O0 -fsycl-device-only -Xclang -verify %s
+// expected-no-diagnostics
+// Tests for propagation of compile options
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi::experimental;
+
+int main() {
+  queue Q;
+  // CHECK-IR: spir_kernel void @{{.*}}Kernel0(){{.*}} #[[COAttr1:[0-9]+]]
+  Q.single_task<class Kernel0>([]() {});
+}