[SYCL] Add e2e test for optional kernel features and virtual functions

llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h

@@ -130,6 +130,7 @@ class ModuleDesc {
   EntryPointGroup EntryPoints;
   bool IsTopLevel = false;
   mutable std::optional<SYCLDeviceRequirements> Reqs;
+  bool IsDummyImage = false;
 
 public:
   struct Properties {
@@ -225,6 +226,9 @@ class ModuleDesc {
 
   void saveSplitInformationAsMetadata();
 
+  ModuleDesc makeDummy() const;
+  bool isDummyImage() { return IsDummyImage; }
+
 #ifndef NDEBUG
   void verifyESIMDProperty() const;
   void dump() const;

llvm/lib/SYCLLowerIR/ModuleSplitter.cpp

@@ -816,6 +816,15 @@ void ModuleDesc::saveSplitInformationAsMetadata() {
         SpecConstantsPass::SPEC_CONST_DEFAULT_VAL_MODULE_MD_STRING);
 }
 
+ModuleDesc ModuleDesc::makeDummy() const {
+  ModuleDesc MD(CloneModule(getModule()));
+  MD.EntryPoints = EntryPoints;
+  MD.IsTopLevel = IsTopLevel;
+  MD.Reqs = Reqs;
+  MD.IsDummyImage = true;
+  return MD;
+}
+
 void EntryPointGroup::saveNames(std::vector<std::string> &Dest) const {
   Dest.reserve(Dest.size() + Functions.size());
   std::transform(Functions.begin(), Functions.end(),

llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp

@@ -665,6 +665,22 @@ void processDeclaredVirtualFunctionSets(
     StringMap<SmallVector<Function *, 4>> &VirtualFunctionSets) {
   if (!F->hasFnAttribute("calls-indirectly"))
     return;
+
+  // "Construction" kernels which reference the vtable
+  // can be marked with calls-indirectly attribute by SYCLVirtualFunctionAnalysis pass. 
+  bool hasVirtualCall = false;
+  for (const Instruction &I : instructions(F)) {
+    const auto *CI = dyn_cast<CallInst>(&I);
+    if (!CI)
+      continue;
+    if (CI->isIndirectCall() && CI->hasFnAttr("virtual-call")) {
+      hasVirtualCall = true;
+      break;
+    }
+  }
+  if (!hasVirtualCall)
+    return;
+
   Attribute CallsIndirectlyAttr = F->getFnAttribute("calls-indirectly");
   SmallVector<StringRef, 4> DeclaredVirtualFunctionSetNames;
   CallsIndirectlyAttr.getValueAsString().split(DeclaredVirtualFunctionSetNames,

llvm/test/tools/sycl-post-link/virtual-functions/dummy.ll

@@ -0,0 +1,37 @@
+; RUN: sycl-post-link -split=auto -properties -S < %s -o %t.table
+; RUN: FileCheck %s --input-file=%t.table --check-prefix=CHECK-TABLE
+; RUN: FileCheck %s --input-file=%t_0.ll --check-prefix=CHECK-FP64-SPLIT
+; RUN: FileCheck %s --input-file=%t_1.ll --check-prefix=CHECK-FP64-DUMMY
+; RUN: FileCheck %s --input-file=%t_1.prop --check-prefix=CHECK-FP64-DUMMY-PROPS
+; RUN: FileCheck %s --input-file=%t_2.ll --check-prefix=CHECK-FP32-SPLIT
+
+; CHECK-TABLE:      _0.prop
+; CHECK-TABLE-NEXT: _1.prop
+; CHECK-TABLE-NEXT: _2.prop
+
+; CHECK-FP64-SPLIT: define spir_func void @bar()
+; CHECK-FP32-SPLIT: define spir_func void @foo()
+
+; CHECK-FP64-DUMMY: define spir_func void @bar()
+; CHECK-FP64-DUMMY-NEXT: entry:
+; CHECK-FP64-DUMMY-NEXT: ret void
+
+; CHECK-FP64-DUMMY-PROPS: dummy-image=1
+
+define spir_func void @foo() #1 {
+  %x = alloca float
+  ret void
+}
+
+define spir_func void @bar() #1 !sycl_used_aspects !1 {
+  %x = alloca double
+  %d = load double, ptr %x
+  %res = fadd double %d, %d
+  ret void
+}
+
+attributes #1 = { "sycl-module-id"="v.cpp" "indirectly-callable"="setA" }
+
+!sycl_aspects = !{!0}
+!0 = !{!"fp64", i32 6}
+!1 = !{i32 6}

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

@@ -58,6 +58,7 @@
 #include "llvm/Transforms/Scalar/DCE.h"
 #include "llvm/Transforms/Scalar/EarlyCSE.h"
 #include "llvm/Transforms/Scalar/SROA.h"
+#include "llvm/Transforms/Utils/Cloning.h"
 #include "llvm/Transforms/Utils/GlobalStatus.h"
 
 #include <algorithm>
@@ -295,11 +296,35 @@ void saveModuleIR(Module &M, StringRef OutFilename) {
   MPM.run(M, MAM);
 }
 
-std::string saveModuleIR(Module &M, int I, StringRef Suff) {
-  DUMP_ENTRY_POINTS(M, EmitOnlyKernelsAsEntryPoints, "saving IR");
+std::unique_ptr<Module> makeDummyImageIR(const Module &M) {
+  auto MCopy = CloneModule(M);
+  for (Function &F : MCopy->functions()) {
+    if (!F.hasFnAttribute("indirectly-callable"))
+      continue;
+
+    F.erase(F.begin(), F.end());
+    BasicBlock *newBB = BasicBlock::Create(F.getContext(), "entry", &F);
+    IRBuilder<> builder(newBB);
+    if (F.getReturnType()->isVoidTy())
+      builder.CreateRetVoid();
+    else
+      builder.CreateRet(UndefValue::get(F.getReturnType()));
+  }
+  return MCopy;
+}
+
+std::string saveModuleIR(module_split::ModuleDesc &MD, int I, StringRef Suff) {
+  std::unique_ptr<Module> Storage;
+  Module *M = &MD.getModule();
+  if (MD.isDummyImage()) {
+    Storage = makeDummyImageIR(MD.getModule());
+    M = Storage.get();
+  }
+
+  DUMP_ENTRY_POINTS(*M, EmitOnlyKernelsAsEntryPoints, "saving IR");
   StringRef FileExt = (OutputAssembly) ? ".ll" : ".bc";
   std::string OutFilename = makeResultFileName(FileExt, I, Suff);
-  saveModuleIR(M, OutFilename);
+  saveModuleIR(*M, OutFilename);
   return OutFilename;
 }
 
@@ -317,6 +342,9 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
     NewSuff += Target;
   }
 
+  if (MD.isDummyImage())
+    PropSet.add(PropSetRegTy::SYCL_VIRTUAL_FUNCTIONS, "dummy-image", 1);
+
   std::error_code EC;
   std::string SCFile = makeResultFileName(".prop", I, NewSuff);
   raw_fd_ostream SCOut(SCFile, EC);
@@ -424,7 +452,7 @@ void saveModule(std::vector<std::unique_ptr<util::SimpleTable>> &OutTables,
     BaseTriple.Ir = IRFilename.str();
   } else {
     MD.cleanup();
-    BaseTriple.Ir = saveModuleIR(MD.getModule(), I, Suffix);
+    BaseTriple.Ir = saveModuleIR(MD, I, Suffix);
   }
   if (DoSymGen) {
     // save the names of the entry points - the symbol table
@@ -740,6 +768,53 @@ bool isTargetCompatibleWithModule(const std::string &Target,
   return true;
 }
 
+bool hasVirtualFunctionsAndOptionalKernelFeatures(const Module &M) {
+  bool hasVirtualFunctions = false;
+  bool hasOptionalKernelFeatures = false;
+  for (const Function &F : M.functions()) {
+    if (F.hasFnAttribute("indirectly-callable"))
+      hasVirtualFunctions = true;
+    if (F.getMetadata("sycl_used_aspects"))
+      hasOptionalKernelFeatures = true;
+    if (hasVirtualFunctions && hasOptionalKernelFeatures)
+      break;
+  }
+  return hasVirtualFunctions && hasOptionalKernelFeatures;
+}
+
+// std::optional<module_split::ModuleDesc>
+// makeDummy(module_split::ModuleDesc &MD) {
+//   bool hasVirtualFunctions = false;
+//   bool hasOptionalKernelFeatures = false;
+//   for (Function &F : M.functions()) {
+//     if (F.hasFnAttribute("indirectly-callable"))
+//       hasVirtualFunctions = true;
+//     if (F.getMetadata("sycl_used_aspects"))
+//       hasOptionalKernelFeatures = true;
+//     if (hasVirtualFunctions && hasOptionalKernelFeatures)
+//       break;
+//   }
+//   if (!hasVirtualFunctions || !hasOptionalKernelFeatures)
+//     return {};
+
+//   auto MDCopy = MD.clone();
+
+//   for (Function &F : MDCopy.getModule().functions()) {
+//     if (!F.hasFnAttribute("indirectly-callable"))
+//       continue;
+
+//     F.erase(F.begin(), F.end());
+//     BasicBlock *newBB = BasicBlock::Create(F.getContext(), "entry", &F);
+//     IRBuilder<> builder(newBB);
+//     if (F.getReturnType()->isVoidTy())
+//       builder.CreateRetVoid();
+//     else
+//       builder.CreateRet(UndefValue::get(F.getReturnType()));
+//   }
+
+//   return MDCopy;
+// }
+
 std::vector<std::unique_ptr<util::SimpleTable>>
 processInputModule(std::unique_ptr<Module> M) {
   // Construct the resulting table which will accumulate all the outputs.
@@ -887,6 +962,21 @@ processInputModule(std::unique_ptr<Module> M) {
 
       ++ID;
     }
+
+    // For kernels with virtual functions and optional kernel features, generate
+    // a dummy image to avoid link errors. A dummy image for a set of virtual
+    // functions is a module with the same set of virtual functions, but with
+    // those function bodies replaced with just a return.
+    bool dummyEmitted = false;
+    for (module_split::ModuleDesc &IrMD : MMs) {
+      if ((dummyEmitted = hasVirtualFunctionsAndOptionalKernelFeatures(
+               IrMD.getModule()))) {
+        auto DummyImage = IrMD.makeDummy();
+        saveModule(Tables, DummyImage, ID, OutIRFileName);
+      }
+    }
+    if (dummyEmitted)
+      ++ID;
   }
   return Tables;
 }

sycl/source/detail/device_binary_image.cpp

@@ -20,6 +20,35 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 
+void printAspects(std::ostream &Out, ByteArray BA) {
+  BA.dropBytes(8);
+  Out << "[";
+  for (int i = 0; !BA.empty(); ++i) {
+    auto Aspect = BA.consume<sycl::aspect>();
+    switch (Aspect) {
+#define __SYCL_ASPECT(ASPECT, ID)                                              \
+  case sycl::aspect::ASPECT:                                                   \
+    Out << #ASPECT;                                                            \
+    break;
+#define __SYCL_ASPECT_DEPRECATED(ASPECT, ID, MESSAGE)                          \
+  case sycl::aspect::ASPECT:                                                   \
+    Out << #ASPECT;                                                            \
+    break;
+#include <sycl/info/aspects.def>
+#include <sycl/info/aspects_deprecated.def>
+#undef __SYCL_ASPECT
+#undef __SYCL_ASPECT_DEPRECATED
+    default:
+      Out << "unknown (" << static_cast<int>(Aspect) << ")";
+      break;
+    }
+    if (i != 0)
+      Out << ", ";
+  }
+  Out << "]";
+  return;
+}
+
 std::ostream &operator<<(std::ostream &Out, const DeviceBinaryProperty &P) {
   switch (P.Prop->Type) {
   case SYCL_PROPERTY_TYPE_UINT32:
@@ -42,6 +71,20 @@ std::ostream &operator<<(std::ostream &Out, const DeviceBinaryProperty &P) {
     Out << P.asUint32();
     break;
   case SYCL_PROPERTY_TYPE_BYTE_ARRAY: {
+    // Special case for aspects, print out the aspect names
+    if (P.Prop->Name == std::string_view("aspects")) {
+      printAspects(Out, P.asByteArray());
+      break;
+    }
+
+    // Special case for these properties, print out the value as a string
+    if (P.Prop->Name == std::string_view("virtual-functions-set") ||
+        P.Prop->Name == std::string_view("uses-virtual-functions-set")) {
+      Out << P.asStringView();
+      break;
+    }
+
+    // Otherwise, print out the byte array as hex
     ByteArray BA = P.asByteArray();
     std::ios_base::fmtflags FlagsBackup = Out.flags();
     Out << std::hex;
@@ -52,7 +95,7 @@ std::ostream &operator<<(std::ostream &Out, const DeviceBinaryProperty &P) {
     break;
   }
   case SYCL_PROPERTY_TYPE_STRING:
-    Out << P.asCString();
+    Out << P.asStringView();
     break;
   default:
     assert(false && "Unsupported property");
@@ -77,14 +120,14 @@ ByteArray DeviceBinaryProperty::asByteArray() const {
   return {Data, Prop->ValSize};
 }
 
-const char *DeviceBinaryProperty::asCString() const {
+std::string_view DeviceBinaryProperty::asStringView() const {
   assert((Prop->Type == SYCL_PROPERTY_TYPE_STRING ||
           Prop->Type == SYCL_PROPERTY_TYPE_BYTE_ARRAY) &&
          "property type mismatch");
   assert(Prop->ValSize > 0 && "property size mismatch");
   // Byte array stores its size in first 8 bytes
   size_t Shift = Prop->Type == SYCL_PROPERTY_TYPE_BYTE_ARRAY ? 8 : 0;
-  return ur::cast<const char *>(Prop->ValAddr) + Shift;
+  return {ur::cast<const char *>(Prop->ValAddr) + Shift, Prop->ValSize - Shift};
 }
 
 void RTDeviceBinaryImage::PropertyRange::init(sycl_device_binary Bin,

sycl/source/detail/device_binary_image.hpp

@@ -73,7 +73,7 @@ class DeviceBinaryProperty {
 
   uint32_t asUint32() const;
   ByteArray asByteArray() const;
-  const char *asCString() const;
+  std::string_view asStringView() const;
 
 protected:
   friend std::ostream &operator<<(std::ostream &Out,