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

Author: jzc

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

@@ -16,7 +16,7 @@
 ; CHECK-FP64-DUMMY-NEXT: entry:
 ; CHECK-FP64-DUMMY-NEXT: ret void
 
-; CHECK-FP64-DUMMY-PROPS: dummy=1
+; CHECK-FP64-DUMMY-PROPS: dummy-image=1
 
 define spir_func void @foo() #1 {
   %x = alloca float

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,18 +296,41 @@ 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;
 }
 
 std::string saveModuleProperties(module_split::ModuleDesc &MD,
                                  const GlobalBinImageProps &GlobProps, int I,
-                                 StringRef Suff, StringRef Target = "",
-                                 bool IsDummy = false) {
+                                 StringRef Suff, StringRef Target = "") {
   auto PropSet =
       computeModuleProperties(MD.getModule(), MD.entries(), GlobProps);
 
@@ -318,9 +342,8 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
     NewSuff += Target;
   }
 
-  if (IsDummy) {
-    PropSet.add(PropSetRegTy::SYCL_VIRTUAL_FUNCTIONS, "dummy", 1);
-  }
+  if (MD.isDummyImage())
+    PropSet.add(PropSetRegTy::SYCL_VIRTUAL_FUNCTIONS, "dummy-image", 1);
 
   std::error_code EC;
   std::string SCFile = makeResultFileName(".prop", I, NewSuff);
@@ -420,8 +443,7 @@ void addTableRow(util::SimpleTable &Table,
 //   IR component saving is skipped, and this file name is recorded as such in
 //   the result.
 void saveModule(std::vector<std::unique_ptr<util::SimpleTable>> &OutTables,
-                module_split::ModuleDesc &MD, int I, StringRef IRFilename,
-                bool IsDummy = false) {
+                module_split::ModuleDesc &MD, int I, StringRef IRFilename) {
   IrPropSymFilenameTriple BaseTriple;
   StringRef Suffix = getModuleSuffix(MD);
   MD.saveSplitInformationAsMetadata();
@@ -430,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
@@ -445,8 +467,8 @@ void saveModule(std::vector<std::unique_ptr<util::SimpleTable>> &OutTables,
       GlobalBinImageProps Props = {EmitKernelParamInfo, EmitProgramMetadata,
                                    EmitExportedSymbols, EmitImportedSymbols,
                                    DeviceGlobals};
-      CopyTriple.Prop = saveModuleProperties(MD, Props, I, Suffix,
-                                             OutputFile.Target, IsDummy);
+      CopyTriple.Prop =
+          saveModuleProperties(MD, Props, I, Suffix, OutputFile.Target);
     }
     addTableRow(*Table, CopyTriple);
   }
@@ -746,36 +768,53 @@ bool isTargetCompatibleWithModule(const std::string &Target,
   return true;
 }
 
-std::optional<module_split::ModuleDesc>
-makeDummy(module_split::ModuleDesc &MD) {
+bool hasVirtualFunctionsAndOptionalKernelFeatures(const Module &M) {
   bool hasVirtualFunctions = false;
   bool hasOptionalKernelFeatures = false;
-  for (Function &F : MD.getModule().functions()) {
+  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;
   }
-  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);
-    builder.CreateRetVoid();
-  }
-
-  return MDCopy;
+  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.
@@ -924,11 +963,16 @@ 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 (auto Dummy = makeDummy(IrMD)) {
-        saveModule(Tables, *Dummy, ID, OutIRFileName, /*IsDummy*/ true);
-        dummyEmitted = true;
+      if ((dummyEmitted = hasVirtualFunctionsAndOptionalKernelFeatures(
+               IrMD.getModule()))) {
+        auto DummyImage = IrMD.makeDummy();
+        saveModule(Tables, DummyImage, ID, OutIRFileName);
       }
     }
     if (dummyEmitted)

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;
@@ -84,7 +127,7 @@ std::string_view DeviceBinaryProperty::asStringView() const {
   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, Prop->ValSize};
+  return {ur::cast<const char *>(Prop->ValAddr) + Shift, Prop->ValSize - Shift};
 }
 
 void RTDeviceBinaryImage::PropertyRange::init(sycl_device_binary Bin,

sycl/source/detail/program_manager/program_manager.cpp

@@ -732,8 +732,10 @@ ProgramManager::collectDependentDeviceImagesForVirtualFunctions(
       // when the dummy image does not support the device requirements, we
       // know the corresponding image providing virtual functions was not
       // linked and we must link the dummy image.
-      if (doesDevSupportDeviceRequirements(Dev, *BinImage) + isDummyImage == 1)
+      if (doesDevSupportDeviceRequirements(Dev, *BinImage) + isDummyImage ==
+          1) {
         DeviceImagesToLink.insert(BinImage);
+      }
     }
   }
 
@@ -1765,6 +1767,9 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
     else
       Img = std::make_unique<RTDeviceBinaryImage>(RawImg);
 
+    if (std::getenv("SYCL_PRINT_IMAGES"))
+      Img->print();
+
     static uint32_t SequenceID = 0;
 
     // Fill the kernel argument mask map