[SYCL] Propagate sycl_fixed_targets metadata

llvm/include/llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h

@@ -20,7 +20,14 @@ namespace llvm {
 class SYCLPropagateAspectsUsagePass
     : public PassInfoMixin<SYCLPropagateAspectsUsagePass> {
 public:
+  SYCLPropagateAspectsUsagePass(StringRef OptionsString = {}) {
+    OptionsString.split(this->TargetFixedAspects, ',', /*MaxSplit=*/-1,
+                        /*KeepEmpty=*/false);
+  };
   PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
+
+private:
+  SmallVector<StringRef, 8> TargetFixedAspects;
 };
 
 } // namespace llvm

llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp

@@ -36,6 +36,7 @@
 #include "llvm/IR/IntrinsicInst.h"
 #include "llvm/IR/Module.h"
 #include "llvm/Pass.h"
+#include "llvm/Support/CommandLine.h"
 #include "llvm/Support/Path.h"
 
 #include <queue>
@@ -44,6 +45,12 @@
 
 using namespace llvm;
 
+static cl::opt<std::string> ClSyclFixedTargets(
+    "sycl-propagate-aspects-usage-fixed-targets",
+    cl::desc("Specify target device(s) all device code in the translation unit "
+             "is expected to be runnable on"),
+    cl::Hidden, cl::init(""));
+
 namespace {
 
 using AspectsSetTy = SmallSet<int, 4>;
@@ -352,19 +359,41 @@ bool isEntryPoint(const Function &F) {
   return F.hasFnAttribute("sycl-module-id") && !isSpirvSyclBuiltin(F.getName());
 }
 
+void setSyclFixedTargetsMD(const std::vector<Function *> &EntryPoints,
+                           const SmallVector<StringRef, 8> &Targets,
+                           AspectValueToNameMapTy &AspectValues) {
+  if (EntryPoints.empty())
+    return;
+
+  SmallVector<Metadata *, 8> TargetsMD;
+  LLVMContext &C = EntryPoints[0]->getContext();
+
+  for (const auto &Target : Targets) {
+    if (!Target.empty()) {
+      auto AspectIt = AspectValues.find(Target);
+      if (AspectIt != AspectValues.end()) {
+        auto ConstIntTarget =
+            ConstantInt::getSigned(Type::getInt32Ty(C), AspectIt->second);
+        TargetsMD.push_back(ConstantAsMetadata::get(ConstIntTarget));
+      }
+    }
+  }
+
+  MDNode *MDN = MDNode::get(C, TargetsMD);
+  for (Function *F : EntryPoints)
+    F->setMetadata("sycl_fixed_targets", MDN);
+}
+
 /// Returns a map of functions with corresponding used aspects.
 FunctionToAspectsMapTy
-buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects) {
+buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects,
+                           const std::vector<Function *> &EntryPoints) {
   FunctionToAspectsMapTy FunctionToAspects;
   CallGraphTy CG;
-  std::vector<Function *> EntryPoints;
+
   for (Function &F : M.functions()) {
     if (F.isDeclaration())
       continue;
-
-    if (isEntryPoint(F))
-      EntryPoints.push_back(&F);
-
     processFunction(F, FunctionToAspects, TypesWithAspects, CG);
   }
 
@@ -392,14 +421,25 @@ SYCLPropagateAspectsUsagePass::run(Module &M, ModuleAnalysisManager &MAM) {
     return PreservedAnalyses::all();
   }
 
+  if (ClSyclFixedTargets.getNumOccurrences() > 0)
+    StringRef(ClSyclFixedTargets)
+        .split(TargetFixedAspects, ',', /*MaxSplit=*/-1, /*KeepEmpty=*/false);
+
+  std::vector<Function *> EntryPoints;
+  for (Function &F : M.functions())
+    if (isEntryPoint(F))
+      EntryPoints.push_back(&F);
+
   propagateAspectsToOtherTypesInModule(M, TypesWithAspects, AspectValues);
 
   FunctionToAspectsMapTy FunctionToAspects =
-      buildFunctionsToAspectsMap(M, TypesWithAspects);
+      buildFunctionsToAspectsMap(M, TypesWithAspects, EntryPoints);
 
   createUsedAspectsMetadataForFunctions(FunctionToAspects);
   // FIXME: check and diagnose if a function uses an aspect which was not
   // declared through [[sycl::device_has()]] attribute
 
+  setSyclFixedTargetsMD(EntryPoints, TargetFixedAspects, AspectValues);
+
   return PreservedAnalyses::all();
 }

llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropagateFunctionWithSyclDetailAspect/propagate-aspects-from-function-with-sycl-detail-1.ll

@@ -14,8 +14,8 @@ source_filename = "main.cpp"
 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"
 
-; CHECK: void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_() !sycl_used_aspects ![[#ASPECT:]] {
-define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_() {
+; CHECK: void @kernel() !sycl_used_aspects ![[#ASPECT:]]
+define weak_odr dso_local spir_kernel void @kernel() {
 entry:
   call spir_func void @_Z3bazv()
   ret void

llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropagateFunctionWithSyclDetailAspect/propagate-aspects-from-function-with-sycl-detail-2.ll

@@ -11,7 +11,7 @@ source_filename = "main.cpp"
 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"
 
-; CHECK: void @_Z3bazv() !sycl_used_aspects ![[#ASPECT1:]] {
+; CHECK: void @_Z3bazv() !sycl_used_aspects ![[#ASPECT1:]]
 define dso_local spir_kernel void @_Z3bazv() {
 entry:
   call spir_func void @_Z3barv()

llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropagateFunctionWithSyclDetailAspect/propagate-aspects-from-function-with-sycl-detail-3.ll

@@ -15,7 +15,7 @@
 %Optional.A = type { i32 }
 %Optional.B = type { i32 }
 
-; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[#ID1:]] {
+; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[#ID1:]]
 define spir_kernel void @kernel() {
   call spir_func void @func1()
   call spir_func void @func2()

llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-1.ll

@@ -14,14 +14,14 @@
 %Optional.A = type { i32 }
 %Optional.B = type { i32 }
 
-; CHECK: spir_kernel void @kernel1() !sycl_used_aspects ![[#ID1:]] {
+; CHECK: spir_kernel void @kernel1() !sycl_used_aspects ![[#ID1:]]
 define spir_kernel void @kernel1() {
   call spir_func void @func1()
   call spir_func void @func2()
   ret void
 }
 
-; CHECK: spir_kernel void @kernel2() !sycl_used_aspects ![[#ID2:]] {
+; CHECK: spir_kernel void @kernel2() !sycl_used_aspects ![[#ID2:]]
 define spir_kernel void @kernel2() {
   call spir_func void @func2()
   call spir_func void @func3()

llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-2.ll

@@ -15,7 +15,7 @@
 %Optional.A = type { i32 }
 %Optional.B = type { i32 }
 
-; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[#ID1:]] {
+; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[#ID1:]]
 define spir_kernel void @kernel() {
   call spir_func void @func1()
   call spir_func void @func2()

llvm/test/SYCLLowerIR/PropagateAspectsUsage/composite-types-1.ll

@@ -22,7 +22,7 @@
 
 %F2.does.not.contain.optional = type { %B.core, %C.core*, %D2.does.not.contain.optional* }
 
-; CHECK: spir_kernel void @kernelD1.uses.optional() !sycl_used_aspects ![[MDID:[0-9]+]] {
+; CHECK: spir_kernel void @kernelD1.uses.optional() !sycl_used_aspects ![[MDID:[0-9]+]]
 define spir_kernel void @kernelD1.uses.optional() {
   %tmp = alloca %D1.contains.optional
   ret void
@@ -34,7 +34,7 @@ define spir_func void @funcD1.uses.optional() {
   ret void
 }
 
-; CHECK: spir_kernel void @kernelD2.does.not.use.optional() {
+; CHECK: spir_kernel void @kernelD2.does.not.use.optional()
 define spir_kernel void @kernelD2.does.not.use.optional() {
   %tmp = alloca %D2.does.not.contain.optional
   ret void
@@ -46,7 +46,7 @@ define spir_func void @funcD2.does.not.use.optional() {
   ret void
 }
 
-; CHECK: spir_kernel void @kernelE.uses.optional() !sycl_used_aspects ![[MDID]] {
+; CHECK: spir_kernel void @kernelE.uses.optional() !sycl_used_aspects ![[MDID]]
 define spir_kernel void @kernelE.uses.optional() {
   %tmp = alloca %E.contains.optional
   ret void
@@ -58,7 +58,7 @@ define spir_func void @funcE.uses.optional() {
   ret void
 }
 
-; CHECK: spir_kernel void @kernelF1.points.to.optional() {
+; CHECK: spir_kernel void @kernelF1.points.to.optional()
 define spir_kernel void @kernelF1.points.to.optional() {
   %tmp = alloca %F1.points.to.optional
   ret void
@@ -70,7 +70,7 @@ define spir_func void @funcF1.points.to.optional() {
   ret void
 }
 
-; CHECK: spir_kernel void @kernelF2.does.not.use.optional() {
+; CHECK: spir_kernel void @kernelF2.does.not.use.optional()
 define spir_kernel void @kernelF2.does.not.use.optional() {
   %tmp = alloca %F2.does.not.contain.optional
   ret void

llvm/test/SYCLLowerIR/PropagateAspectsUsage/double.ll

@@ -4,7 +4,7 @@
 
 %composite = type { double }
 
-; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[MDID:[0-9]+]] {
+; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[MDID:[0-9]+]]
 define spir_kernel void @kernel() {
   call spir_func void @func()
   ret void

llvm/test/SYCLLowerIR/PropagateAspectsUsage/multiple-aspects.ll

@@ -34,7 +34,7 @@ define spir_func void @funcD() {
   ret void
 }
 
-; CHECK: define spir_kernel void @kernel() !sycl_used_aspects ![[#ID3]] {
+; CHECK: define spir_kernel void @kernel() !sycl_used_aspects ![[#ID3]]
 define spir_kernel void @kernel() {
   call spir_func void @funcD()
   ret void

llvm/test/SYCLLowerIR/PropagateAspectsUsage/no-uses-of-optional.ll

@@ -5,7 +5,7 @@
 
 %MyStruct = type { i32 }
 
-; CHECK: dso_local spir_kernel void @kernel() {
+; CHECK: dso_local spir_kernel void @kernel()
 define dso_local spir_kernel void @kernel() {
   call spir_func void @func()
   ret void

llvm/test/SYCLLowerIR/PropagateAspectsUsage/sycl-fixed-targets.ll

@@ -0,0 +1,20 @@
+; RUN: opt -passes=sycl-propagate-aspects-usage -sycl-propagate-aspects-usage-fixed-targets=host,cpu,gpu %s -S | FileCheck %s
+
+source_filename = "main.cpp"
+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"
+
+; CHECK: void @kernel(){{.*}}!sycl_fixed_targets ![[#MDNUM:]]
+define weak_odr dso_local spir_kernel void @kernel() {
+entry:
+  ret void
+}
+
+!sycl_aspects = !{!0, !1, !2, !3}
+
+; CHECK: ![[#MDNUM]] = !{i32 0, i32 1, i32 2}
+
+!0 = !{!"host", i32 0}
+!1 = !{!"cpu", i32 1}
+!2 = !{!"gpu", i32 2}
+!3 = !{!"fp64", i32 6}