[SYCL]Link Fallback Device Libraries On Demand

libdevice/fallback-cmath-fp64.cpp

@@ -10,6 +10,9 @@
 
 #ifdef __SPIR__
 
+// To support fallback device libraries on-demand loading, please update the
+// DeviceLibFuncMap in llvm/tools/sycl-post-link/sycl-post-link.cpp if you add
+// or remove any item in this file.
 DEVICE_EXTERN_C
 double __devicelib_log(double x) { return __spirv_ocl_log(x); }
 

libdevice/fallback-cmath.cpp

@@ -10,6 +10,11 @@
 
 #ifdef __SPIR__
 
+// To support fallback device libraries on-demand loading, please update the
+// DeviceLibFuncMap in llvm/tools/sycl-post-link/sycl-post-link.cpp if you add
+// or remove any item in this file.
+// TODO: generate the DeviceLibFuncMap in sycl-post-link.cpp automatically
+// during the build based on libdevice to avoid manually sync.
 DEVICE_EXTERN_C
 float __devicelib_scalbnf(float x, int n) { return __spirv_ocl_ldexp(x, n); }
 

libdevice/fallback-complex-fp64.cpp

@@ -12,6 +12,9 @@
 #ifdef __SPIR__
 #include <cmath>
 
+// To support fallback device libraries on-demand loading, please update the
+// DeviceLibFuncMap in llvm/tools/sycl-post-link/sycl-post-link.cpp if you add
+// or remove any item in this file.
 DEVICE_EXTERN_C
 double __devicelib_creal(double __complex__ z) { return __real__(z); }
 

libdevice/fallback-complex.cpp

@@ -12,6 +12,9 @@
 #ifdef __SPIR__
 #include <cmath>
 
+// To support fallback device libraries on-demand loading, please update the
+// DeviceLibFuncMap in llvm/tools/sycl-post-link/sycl-post-link.cpp if you add
+// or remove any item in this file.
 DEVICE_EXTERN_C
 float __devicelib_crealf(float __complex__ z) { return __real__(z); }
 

llvm/include/llvm/Support/PropertySetIO.h

@@ -123,6 +123,7 @@ class PropertySetRegistry {
   // Specific property category names used by tools.
   static constexpr char SYCL_SPECIALIZATION_CONSTANTS[] =
       "SYCL/specialization constants";
+  static constexpr char SYCL_DEVICELIB_REQ_MASK[] = "SYCL/devicelib req mask";
 
   // Function for bulk addition of an entire property set under given category
   // (property set name).
@@ -160,4 +161,4 @@ class PropertySetRegistry {
 } // namespace util
 } // namespace llvm
 
-#endif // #define LLVM_SUPPORT_PROPERTYSETIO_H
+#endif // #define LLVM_SUPPORT_PROPERTYSETIO_H

llvm/include/llvm/Support/SYCLRTShared.h

@@ -0,0 +1,29 @@
+//=- SYCLRTShared.h - Shared definition between llvm tools and SYCL runtime -=//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+#ifndef LLVM_SUPPORT_SYCLRTSHARED_H
+#define LLVM_SUPPORT_SYCLRTSHARED_H
+/// Device binary image property set names recognized by the SYCL runtime.
+/// Name must be consistent with
+/// PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS defined in
+/// PropertySetIO.h
+#define PI_PROPERTY_SET_SPEC_CONST_MAP "SYCL/specialization constants"
+#define PI_PROPERTY_SET_DEVICELIB_REQ_MASK "SYCL/devicelib req mask"
+namespace llvm {
+namespace util {
+namespace sycl {
+enum DeviceLibExt {
+  cl_intel_devicelib_assert = 0,
+  cl_intel_devicelib_math,
+  cl_intel_devicelib_math_fp64,
+  cl_intel_devicelib_complex,
+  cl_intel_devicelib_complex_fp64
+};
+}
+} // namespace util
+} // namespace llvm
+#endif

llvm/lib/Support/PropertySetIO.cpp

@@ -112,5 +112,6 @@ template <> PropertyValue::Type PropertyValue::getTypeTag<uint32_t>() {
 }
 
 constexpr char PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS[];
+constexpr char PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK[];
 } // namespace util
 } // namespace llvm

llvm/lib/Support/SimpleTable.cpp

@@ -64,7 +64,7 @@ int SimpleTable::getColumnId(StringRef ColName) const {
 
 Error SimpleTable::addColumnName(StringRef ColName) {
   if (ColumnName2Num.find(ColName) != ColumnName2Num.end())
-    return makeError("column already exists" + ColName);
+    return makeError("column already exists " + ColName);
   ColumnNames.emplace_back(ColName.str());
   ColumnName2Num[ColumnNames.back()] = static_cast<int>(ColumnNames.size()) - 1;
   ColumnNum2Name.push_back(std::prev(ColumnNames.end()));

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

@@ -27,6 +27,7 @@
 #include "llvm/Support/InitLLVM.h"
 #include "llvm/Support/Path.h"
 #include "llvm/Support/PropertySetIO.h"
+#include "llvm/Support/SYCLRTShared.h"
 #include "llvm/Support/SimpleTable.h"
 #include "llvm/Support/SystemUtils.h"
 #include "llvm/Support/WithColor.h"
@@ -36,6 +37,11 @@
 #include <memory>
 
 using namespace llvm;
+using llvm::util::sycl::cl_intel_devicelib_assert;
+using llvm::util::sycl::cl_intel_devicelib_complex;
+using llvm::util::sycl::cl_intel_devicelib_complex_fp64;
+using llvm::util::sycl::cl_intel_devicelib_math;
+using llvm::util::sycl::cl_intel_devicelib_math_fp64;
 
 using string_vector = std::vector<std::string>;
 using SpecIDMapTy = std::map<StringRef, unsigned>;
@@ -47,7 +53,7 @@ cl::OptionCategory PostLinkCat{"sycl-post-link options"};
 static constexpr char COL_CODE[] = "Code";
 static constexpr char COL_SYM[] = "Symbols";
 static constexpr char COL_PROPS[] = "Properties";
-
+static constexpr char DEVICELIB_FUNC_PREFIX[] = "__devicelib_";
 // InputFilename - The filename to read from.
 static cl::opt<std::string> InputFilename{
     cl::Positional, cl::desc("<input bitcode file>"), cl::init("-"),
@@ -104,6 +110,143 @@ static cl::opt<SpecConstMode> SpecConstLower{
                    "set spec constants to C++ defaults")),
     cl::cat(PostLinkCat)};
 
+struct ImagePropSaveInfo {
+  bool NeedDeviceLibReqMask;
+  bool DoSpecConst;
+  bool SetSpecConstAtRT;
+  bool SpecConstsMet;
+};
+// Please update DeviceLibFuncMap if any item is added to or removed from
+// fallback device libraries in libdevice.
+static std::unordered_map<std::string, uint32_t> DeviceLibFuncMap = {
+    {"__devicelib_acosf", cl_intel_devicelib_math},
+    {"__devicelib_acoshf", cl_intel_devicelib_math},
+    {"__devicelib_asinf", cl_intel_devicelib_math},
+    {"__devicelib_asinhf", cl_intel_devicelib_math},
+    {"__devicelib_atan2f", cl_intel_devicelib_math},
+    {"__devicelib_atanf", cl_intel_devicelib_math},
+    {"__devicelib_atanhf", cl_intel_devicelib_math},
+    {"__devicelib_cbrtf", cl_intel_devicelib_math},
+    {"__devicelib_cosf", cl_intel_devicelib_math},
+    {"__devicelib_coshf", cl_intel_devicelib_math},
+    {"__devicelib_erfcf", cl_intel_devicelib_math},
+    {"__devicelib_erff", cl_intel_devicelib_math},
+    {"__devicelib_exp2f", cl_intel_devicelib_math},
+    {"__devicelib_expf", cl_intel_devicelib_math},
+    {"__devicelib_expm1f", cl_intel_devicelib_math},
+    {"__devicelib_fdimf", cl_intel_devicelib_math},
+    {"__devicelib_fmaf", cl_intel_devicelib_math},
+    {"__devicelib_fmodf", cl_intel_devicelib_math},
+    {"__devicelib_frexpf", cl_intel_devicelib_math},
+    {"__devicelib_hypotf", cl_intel_devicelib_math},
+    {"__devicelib_ilogbf", cl_intel_devicelib_math},
+    {"__devicelib_ldexpf", cl_intel_devicelib_math},
+    {"__devicelib_lgammaf", cl_intel_devicelib_math},
+    {"__devicelib_log10f", cl_intel_devicelib_math},
+    {"__devicelib_log1pf", cl_intel_devicelib_math},
+    {"__devicelib_log2f", cl_intel_devicelib_math},
+    {"__devicelib_logbf", cl_intel_devicelib_math},
+    {"__devicelib_logf", cl_intel_devicelib_math},
+    {"__devicelib_modff", cl_intel_devicelib_math},
+    {"__devicelib_nextafterf", cl_intel_devicelib_math},
+    {"__devicelib_powf", cl_intel_devicelib_math},
+    {"__devicelib_remainderf", cl_intel_devicelib_math},
+    {"__devicelib_remquof", cl_intel_devicelib_math},
+    {"__devicelib_sinf", cl_intel_devicelib_math},
+    {"__devicelib_sinhf", cl_intel_devicelib_math},
+    {"__devicelib_sqrtf", cl_intel_devicelib_math},
+    {"__devicelib_tanf", cl_intel_devicelib_math},
+    {"__devicelib_tanhf", cl_intel_devicelib_math},
+    {"__devicelib_tgammaf", cl_intel_devicelib_math},
+    {"__devicelib_acos", cl_intel_devicelib_math_fp64},
+    {"__devicelib_acosh", cl_intel_devicelib_math_fp64},
+    {"__devicelib_asin", cl_intel_devicelib_math_fp64},
+    {"__devicelib_asinh", cl_intel_devicelib_math_fp64},
+    {"__devicelib_atan", cl_intel_devicelib_math_fp64},
+    {"__devicelib_atan2", cl_intel_devicelib_math_fp64},
+    {"__devicelib_atanh", cl_intel_devicelib_math_fp64},
+    {"__devicelib_cbrt", cl_intel_devicelib_math_fp64},
+    {"__devicelib_cos", cl_intel_devicelib_math_fp64},
+    {"__devicelib_cosh", cl_intel_devicelib_math_fp64},
+    {"__devicelib_erf", cl_intel_devicelib_math_fp64},
+    {"__devicelib_erfc", cl_intel_devicelib_math_fp64},
+    {"__devicelib_exp", cl_intel_devicelib_math_fp64},
+    {"__devicelib_exp2", cl_intel_devicelib_math_fp64},
+    {"__devicelib_expm1", cl_intel_devicelib_math_fp64},
+    {"__devicelib_fdim", cl_intel_devicelib_math_fp64},
+    {"__devicelib_fma", cl_intel_devicelib_math_fp64},
+    {"__devicelib_fmod", cl_intel_devicelib_math_fp64},
+    {"__devicelib_frexp", cl_intel_devicelib_math_fp64},
+    {"__devicelib_hypot", cl_intel_devicelib_math_fp64},
+    {"__devicelib_ilogb", cl_intel_devicelib_math_fp64},
+    {"__devicelib_ldexp", cl_intel_devicelib_math_fp64},
+    {"__devicelib_lgamma", cl_intel_devicelib_math_fp64},
+    {"__devicelib_log", cl_intel_devicelib_math_fp64},
+    {"__devicelib_log10", cl_intel_devicelib_math_fp64},
+    {"__devicelib_log1p", cl_intel_devicelib_math_fp64},
+    {"__devicelib_log2", cl_intel_devicelib_math_fp64},
+    {"__devicelib_logb", cl_intel_devicelib_math_fp64},
+    {"__devicelib_modf", cl_intel_devicelib_math_fp64},
+    {"__devicelib_nextafter", cl_intel_devicelib_math_fp64},
+    {"__devicelib_pow", cl_intel_devicelib_math_fp64},
+    {"__devicelib_remainder", cl_intel_devicelib_math_fp64},
+    {"__devicelib_remquo", cl_intel_devicelib_math_fp64},
+    {"__devicelib_sin", cl_intel_devicelib_math_fp64},
+    {"__devicelib_sinh", cl_intel_devicelib_math_fp64},
+    {"__devicelib_sqrt", cl_intel_devicelib_math_fp64},
+    {"__devicelib_tan", cl_intel_devicelib_math_fp64},
+    {"__devicelib_tanh", cl_intel_devicelib_math_fp64},
+    {"__devicelib_tgamma", cl_intel_devicelib_math_fp64},
+    {"__devicelib___divsc3", cl_intel_devicelib_complex},
+    {"__devicelib___mulsc3", cl_intel_devicelib_complex},
+    {"__devicelib_cabsf", cl_intel_devicelib_complex},
+    {"__devicelib_cacosf", cl_intel_devicelib_complex},
+    {"__devicelib_cacoshf", cl_intel_devicelib_complex},
+    {"__devicelib_cargf", cl_intel_devicelib_complex},
+    {"__devicelib_casinf", cl_intel_devicelib_complex},
+    {"__devicelib_casinhf", cl_intel_devicelib_complex},
+    {"__devicelib_catanf", cl_intel_devicelib_complex},
+    {"__devicelib_catanhf", cl_intel_devicelib_complex},
+    {"__devicelib_ccosf", cl_intel_devicelib_complex},
+    {"__devicelib_ccoshf", cl_intel_devicelib_complex},
+    {"__devicelib_cexpf", cl_intel_devicelib_complex},
+    {"__devicelib_cimagf", cl_intel_devicelib_complex},
+    {"__devicelib_clogf", cl_intel_devicelib_complex},
+    {"__devicelib_cpolarf", cl_intel_devicelib_complex},
+    {"__devicelib_cpowf", cl_intel_devicelib_complex},
+    {"__devicelib_cprojf", cl_intel_devicelib_complex},
+    {"__devicelib_crealf", cl_intel_devicelib_complex},
+    {"__devicelib_csinf", cl_intel_devicelib_complex},
+    {"__devicelib_csinhf", cl_intel_devicelib_complex},
+    {"__devicelib_csqrtf", cl_intel_devicelib_complex},
+    {"__devicelib_ctanf", cl_intel_devicelib_complex},
+    {"__devicelib_ctanhf", cl_intel_devicelib_complex},
+    {"__devicelib___divdc3", cl_intel_devicelib_complex_fp64},
+    {"__devicelib___muldc3", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_cabs", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_cacos", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_cacosh", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_carg", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_casin", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_casinh", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_catan", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_catanh", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_ccos", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_ccosh", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_cexp", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_cimag", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_clog", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_cpolar", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_cpow", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_cproj", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_creal", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_csin", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_csinh", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_csqrt", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_ctan", cl_intel_devicelib_complex_fp64},
+    {"__devicelib_ctanh", cl_intel_devicelib_complex_fp64},
+};
+
 static void error(const Twine &Msg) {
   errs() << "sycl-post-link: " << Msg << '\n';
   exit(1);
@@ -295,20 +438,74 @@ saveResultModules(std::vector<std::unique_ptr<Module>> &ResModules) {
   return Res;
 }
 
-static string_vector
-saveSpecConstantIDMaps(const std::vector<SpecIDMapTy> &Maps) {
-  string_vector Res;
+// Each fallback device library corresponds to one bit in "require mask" which
+// is an unsigned int32. getDeviceLibBit checks which fallback device library
+// is required for FuncName and returns the corresponding bit. The corresponding
+// mask for each fallback device library is:
+// fallback-cassert:      0x1
+// fallback-cmath:        0x2
+// fallback-cmath-fp64:   0x4
+// fallback-complex:      0x8
+// fallback-complex-fp64: 0x10
+static uint32_t getDeviceLibBits(const std::string &FuncName) {
+  auto DeviceLibFuncIter = DeviceLibFuncMap.find(FuncName);
+  return ((DeviceLibFuncIter == DeviceLibFuncMap.end())
+              ? 0
+              : 0x1 << (DeviceLibFuncIter->second - cl_intel_devicelib_assert));
+}
+
+// For each device image module, we go through all functions which meets
+// 1. The function name has prefix "__devicelib_"
+// 2. The function has SPIR_FUNC calling convention
+// 3. The function is declaration which means it doesn't have function body
+static uint32_t getModuleReqMask(const Module &M) {
+  // Device libraries will be enabled only for spir-v module.
+  if (M.getTargetTriple().substr(0, 7) != "spir64-")
+    return 0;
+  // 0x1 means sycl runtime will link and load libsycl-fallback-assert.spv as
+  // default. In fact, default link assert spv is not necessary but dramatic
+  // perf regression is observed if we don't link any device library. The perf
+  // regression is caused by a clang issue.
+  uint32_t ReqMask = 0x1;
+  for (const Function &SF : M) {
+    if (SF.getName().startswith(DEVICELIB_FUNC_PREFIX) &&
+        (SF.getCallingConv() == CallingConv::SPIR_FUNC) && SF.isDeclaration()) {
+      uint32_t DeviceLibBits = getDeviceLibBits(SF.getName().str());
+      ReqMask |= DeviceLibBits;
+    }
+  }
+  return ReqMask;
+}
 
-  for (size_t I = 0; I < Maps.size(); ++I) {
+static string_vector saveDeviceImageProperty(
+    const std::vector<std::unique_ptr<Module>> &ResultModules,
+    const ImagePropSaveInfo &ImgPSInfo) {
+  string_vector Res;
+  for (size_t I = 0; I < ResultModules.size(); ++I) {
     std::string SCFile = makeResultFileName(".prop", I);
     llvm::util::PropertySetRegistry PropSet;
-    PropSet.add(llvm::util::PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS,
-                Maps[I]);
+    if (ImgPSInfo.NeedDeviceLibReqMask) {
+      uint32_t MRMask = getModuleReqMask(*ResultModules[I]);
+      std::map<StringRef, uint32_t> RMEntry = {{"DeviceLibReqMask", MRMask}};
+      PropSet.add(llvm::util::PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK,
+                  RMEntry);
+    }
+    if (ImgPSInfo.DoSpecConst && ImgPSInfo.SetSpecConstAtRT) {
+      // extract spec constant maps per each module
+      SpecIDMapTy TmpSpecIDMap;
+      if (ImgPSInfo.SpecConstsMet)
+        SpecConstantsPass::collectSpecConstantMetadata(*ResultModules[I].get(),
+                                                       TmpSpecIDMap);
+      PropSet.add(
+          llvm::util::PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS,
+          TmpSpecIDMap);
+    }
     std::error_code EC;
     raw_fd_ostream SCOut(SCFile, EC);
     PropSet.write(SCOut);
     Res.emplace_back(std::move(SCFile));
   }
+
   return Res;
 }
 
@@ -412,7 +609,6 @@ int main(int argc, char **argv) {
   }
 
   std::vector<std::unique_ptr<Module>> ResultModules;
-  std::vector<SpecIDMapTy> ResultSpecIDMaps;
   string_vector ResultSymbolsLists;
 
   util::SimpleTable Table;
@@ -456,18 +652,15 @@ int main(int argc, char **argv) {
     Error Err = Table.addColumn(COL_CODE, Files);
     CHECK_AND_EXIT(Err);
   }
-  if (DoSpecConst && SetSpecConstAtRT) {
-    // extract spec constant maps per each module
-    for (auto &MUptr : ResultModules) {
-      ResultSpecIDMaps.emplace_back(SpecIDMapTy());
-      if (SpecConstsMet)
-        SpecConstantsPass::collectSpecConstantMetadata(*MUptr.get(),
-                                                       ResultSpecIDMaps.back());
-    }
-    string_vector Files = saveSpecConstantIDMaps(ResultSpecIDMaps);
+
+  {
+    ImagePropSaveInfo ImgPSInfo = {true, DoSpecConst, SetSpecConstAtRT,
+                                   SpecConstsMet};
+    string_vector Files = saveDeviceImageProperty(ResultModules, ImgPSInfo);
     Error Err = Table.addColumn(COL_PROPS, Files);
     CHECK_AND_EXIT(Err);
   }
+
   if (DoSymGen) {
     // extract symbols per each module
     collectSymbolsLists(GlobalsSet, ResultSymbolsLists);