[SYCL] Propagate explicitly declared aspects even if excluded (#10650)

This PR changes `SYCLPropagateAspectsPass` to propagate aspects that
come from `sycl_declared_aspects` even if they are excluded. The reason
for this change is because a test like
`no-fp64-optimization-declared-aspects.cpp` added in this PR would
failed before with higher optimization level because

- on the first aspect propagation pass, `fp64` is not propagated (to
allow for trivial uses of `float x = 1.5` to optimized out)
- the call to the function marked with `device_has(fp64)` is inlined on
higher optimizations
- that function does not actually use `double` in its body

which means no usage of double ends up in the optimized function,
leading the second aspect propagation pass to not attach `fp64` to its
used aspects metadata.

---------

Co-authored-by: Alexey Sachkov <[email protected]>
Co-authored-by: Marcos Maronas <[email protected]>

Author: 3 people

llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp

@@ -316,11 +316,10 @@ getAspectUsageChain(const Function *F, const FunctionToAspectsMapTy &AspectsMap,
 }
 
 void createUsedAspectsMetadataForFunctions(
-    FunctionToAspectsMapTy &Map, const AspectsSetTy &ExcludeAspectVals) {
-  for (auto &[F, Aspects] : Map) {
-    if (Aspects.empty())
-      continue;
-
+    FunctionToAspectsMapTy &FunctionToUsedAspects,
+    FunctionToAspectsMapTy &FunctionToDeclaredAspects,
+    const AspectsSetTy &ExcludeAspectVals) {
+  for (auto &[F, Aspects] : FunctionToUsedAspects) {
     LLVMContext &C = F->getContext();
 
     // Create a set of unique aspects. First we add the ones from the found
@@ -330,6 +329,11 @@ void createUsedAspectsMetadataForFunctions(
       if (!ExcludeAspectVals.contains(A))
         UniqueAspects.insert(A);
 
+    // The aspects that were propagated via declared aspects are always
+    // added to the metadata.
+    for (const int &A : FunctionToDeclaredAspects[F])
+      UniqueAspects.insert(A);
+
     // If there are no new aspects, we can just keep the old metadata.
     if (UniqueAspects.empty())
       continue;
@@ -547,7 +551,7 @@ void setSyclFixedTargetsMD(const std::vector<Function *> &EntryPoints,
 }
 
 /// Returns a map of functions with corresponding used aspects.
-FunctionToAspectsMapTy
+std::pair<FunctionToAspectsMapTy, FunctionToAspectsMapTy>
 buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects,
                            const AspectValueToNameMapTy &AspectValues,
                            const std::vector<Function *> &EntryPoints,
@@ -575,10 +579,9 @@ buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects,
   Visited.clear();
   for (Function *F : EntryPoints)
     propagateAspectsThroughCG(F, CG, FunctionToDeclaredAspects, Visited);
-  for (const auto &It : FunctionToDeclaredAspects)
-    FunctionToUsedAspects[It.first].insert(It.second.begin(), It.second.end());
 
-  return FunctionToUsedAspects;
+  return {std::move(FunctionToUsedAspects),
+          std::move(FunctionToDeclaredAspects)};
 }
 
 } // anonymous namespace
@@ -617,8 +620,9 @@ SYCLPropagateAspectsUsagePass::run(Module &M, ModuleAnalysisManager &MAM) {
 
   propagateAspectsToOtherTypesInModule(M, TypesWithAspects, AspectValues);
 
-  FunctionToAspectsMapTy FunctionToUsedAspects = buildFunctionsToAspectsMap(
-      M, TypesWithAspects, AspectValues, EntryPoints, ValidateAspectUsage);
+  auto [FunctionToUsedAspects, FunctionToDeclaredAspects] =
+      buildFunctionsToAspectsMap(M, TypesWithAspects, AspectValues, EntryPoints,
+                                 ValidateAspectUsage);
 
   // Create a set of excluded aspect values.
   AspectsSetTy ExcludedAspectVals;
@@ -629,8 +633,8 @@ SYCLPropagateAspectsUsagePass::run(Module &M, ModuleAnalysisManager &MAM) {
     ExcludedAspectVals.insert(AspectValIter->second);
   }
 
-  createUsedAspectsMetadataForFunctions(FunctionToUsedAspects,
-                                        ExcludedAspectVals);
+  createUsedAspectsMetadataForFunctions(
+      FunctionToUsedAspects, FunctionToDeclaredAspects, ExcludedAspectVals);
 
   setSyclFixedTargetsMD(EntryPoints, TargetFixedAspects, AspectValues);
 

llvm/test/SYCLLowerIR/PropagateAspectsUsage/exclude-aspect.ll

@@ -47,34 +47,38 @@ define spir_kernel void @kernel1() {
   ret void
 }
 
-; funcE should get none of its explicitly declared aspects in its
+; funcE should get its explicitly declared aspects even if excluded
 ; sycl_used_aspects
-; CHECK: define spir_func void @funcE() !sycl_declared_aspects ![[#DA1:]] {
+; CHECK:      define spir_func void @funcE() !sycl_declared_aspects ![[#DA1:]]
+; CHECK-SAME: !sycl_used_aspects ![[#DA1]] {
 define spir_func void @funcE() !sycl_declared_aspects !10 {
   ret void
 }
 
 ; funcF should have the same aspects as funcE
-; CHECK-NOT: define spir_func void @funcF() {{.*}} !sycl_used_aspects
+; CHECK: define spir_func void @funcF() !sycl_used_aspects ![[#DA1]] {
 define spir_func void @funcF() {
   call spir_func void @funcE()
   ret void
 }
 
-; funcG only keeps one aspect, the rest are excluded
-; CHECK: define spir_func void @funcG() !sycl_declared_aspects ![[#DA2:]] !sycl_used_aspects ![[#ID3:]]
+; aspect1 is used but excluded, aspect2 and aspect4 are declared, so
+; attached metadata is aspect2 and aspect4
+; CHECK:      define spir_func void @funcG() !sycl_declared_aspects ![[#DA2:]]
+; CHECK-SAME: !sycl_used_aspects ![[#DA2]] {
 define spir_func void @funcG() !sycl_declared_aspects !11 {
+  %tmp = alloca %B
   ret void
 }
 
 ; funcH should have the same aspects as funcG
-; CHECK: define spir_func void @funcH() !sycl_used_aspects ![[#ID3]]
+; CHECK: define spir_func void @funcH() !sycl_used_aspects ![[#DA2]]
 define spir_func void @funcH() {
   call spir_func void @funcG()
   ret void
 }
 
-; CHECK: define spir_kernel void @kernel2() !sycl_used_aspects ![[#ID3]]
+; CHECK: define spir_kernel void @kernel2() !sycl_used_aspects ![[#ID5:]]
 define spir_kernel void @kernel2() {
   call spir_func void @funcF()
   call spir_func void @funcH()
@@ -100,7 +104,7 @@ define spir_func void @funcK() !sycl_used_aspects !11 {
   ret void
 }
 
-; CHECK: define spir_func void @funcL() !sycl_used_aspects ![[#ID3]]
+; CHECK: define spir_func void @funcL() !sycl_used_aspects ![[#ID3:]]
 define spir_func void @funcL() {
   call spir_func void @funcK()
   ret void
@@ -128,12 +132,12 @@ define spir_kernel void @kernel3() {
 !9 = !{!"fp64", i32 5}
 
 !10 = !{i32 1}
-!11 = !{i32 4, i32 2, i32 1}
+!11 = !{i32 4, i32 2}
 ; CHECK-DAG: ![[#DA1]] = !{i32 1}
-; CHECK-DAG: ![[#DA2]] = !{i32 4, i32 2, i32 1}
+; CHECK-DAG: ![[#DA2]] = !{i32 4, i32 2}
 
 ; CHECK-DAG: ![[#ID0]] = !{i32 0}
 ; CHECK-DAG: ![[#ID1]] = !{i32 2, i32 0}
 ; CHECK-DAG: ![[#ID2]] = !{i32 0, i32 2, i32 3}
 ; CHECK-DAG: ![[#ID3]] = !{i32 2}
-; CHECK-DAG: ![[#ID4]] = !{i32 2, i32 4, i32 1}
+; CHECK-DAG: ![[#ID4]] = !{i32 2, i32 4}

sycl/test-e2e/OptionalKernelFeatures/no-fp64-optimization-declared-aspects.cpp

@@ -0,0 +1,35 @@
+// UNSUPPORTED: aspect-fp64
+// RUN: %{build} -o %t.out -O3
+// RUN: %{run} %t.out
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+template <aspect asp, typename T>
+[[sycl::device_has(asp)]] void dummy_function_decorated(const T &acc) {
+  acc[0] = true;
+}
+
+int main() {
+  queue q;
+  bool b = false;
+  assert(!q.get_device().has(aspect::fp64));
+
+  buffer<bool, 1> buf(&b, 1);
+  try {
+    q.submit([&](handler &cgh) {
+      accessor acc(buf, cgh);
+      cgh.single_task([=]() { dummy_function_decorated<aspect::fp64>(acc); });
+    });
+    std::cout << "Exception should have been thrown!\n";
+    return 1;
+  } catch (const sycl::exception &e) {
+    if (e.code() != errc::kernel_not_supported) {
+      std::cout << "Exception caught, but wrong error code!\n";
+      throw;
+    }
+    std::cout << "pass\n";
+    return 0;
+  }
+}