[SYCL] Add aspect names to sycl_used_aspects before cleaning up

clang/lib/CodeGen/BackendUtil.cpp

@@ -53,6 +53,7 @@
 #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"
+#include "llvm/SYCLLowerIR/RecordSYCLAspectNames.h"
 #include "llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h"
 #include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h"
 #include "llvm/SYCLLowerIR/SYCLPropagateJointMatrixUsage.h"
@@ -1153,6 +1154,10 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
       // Process properties and annotations
       MPM.addPass(CompileTimePropertiesPass());
 
+      // Record SYCL aspect names (this should come after propagating aspects
+      // and before cleaning up metadata)
+      MPM.addPass(RecordSYCLAspectNamesPass());
+
       // Remove SYCL metadata added by the frontend, like sycl_aspects
       // Note, this pass should be at the end of the pipeline
       MPM.addPass(CleanupSYCLMetadataPass());

llvm/include/llvm/SYCLLowerIR/RecordSYCLAspectNames.h

@@ -0,0 +1,36 @@
+//===-------- RecordSYCLAspectNames.h - RecordSYCLAspectNames Pass --------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// The !sycl_used_aspects metadata is populated from C++ attributes and
+// further populated by the SYCLPropagateAspectsPass that describes which
+// apsects a function uses. The format of this metadata initially just an
+// integer value corresponding to the enum value in C++. The !sycl_aspects
+// named metadata contains the associations from aspect values to aspect names.
+// These associations are needed later in sycl-post-link, but we drop
+// !sycl_aspects before that to avoid LLVM IR bloat, so this pass takes
+// the associations from !sycl_aspects and then updates all the
+// !sycl_used_aspects metadata to include the aspect names, which allows us
+// to preserve these associations.
+//===----------------------------------------------------------------------===//
+//
+#ifndef LLVM_RECORD_SYCL_ASPECT_NAMES
+#define LLVM_RECORD_SYCL_ASPECT_NAMES
+
+#include "llvm/IR/PassManager.h"
+
+namespace llvm {
+
+class RecordSYCLAspectNamesPass
+    : public PassInfoMixin<RecordSYCLAspectNamesPass> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
+};
+
+} // namespace llvm
+
+#endif // LLVM_RECORD_SYCL_ASPECT_NAMES

llvm/lib/Passes/PassBuilder.cpp

@@ -120,6 +120,7 @@
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/SYCLLowerIR/LowerWGScope.h"
 #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"
+#include "llvm/SYCLLowerIR/RecordSYCLAspectNames.h"
 #include "llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h"
 #include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h"
 #include "llvm/SYCLLowerIR/SYCLPropagateJointMatrixUsage.h"

llvm/lib/Passes/PassRegistry.def

@@ -160,6 +160,7 @@ MODULE_PASS("sycl-add-opt-level-attribute", SYCLAddOptLevelAttributePass())
 MODULE_PASS("compile-time-properties", CompileTimePropertiesPass())
 MODULE_PASS("cleanup-sycl-metadata", CleanupSYCLMetadataPass())
 MODULE_PASS("lower-slm-reservation-calls", ESIMDLowerSLMReservationCalls())
+MODULE_PASS("record-sycl-aspect-names", RecordSYCLAspectNamesPass())
 #undef MODULE_PASS
 
 #ifndef MODULE_PASS_WITH_PARAMS

llvm/lib/SYCLLowerIR/CMakeLists.txt

@@ -44,6 +44,7 @@ add_llvm_component_library(LLVMSYCLLowerIR
   ESIMD/ESIMDRemoveOptnoneNoinline.cpp
   ESIMD/LowerESIMD.cpp
   ESIMD/LowerESIMDKernelAttrs.cpp
+  RecordSYCLAspectNames.cpp
   CleanupSYCLMetadata.cpp
   CompileTimePropertiesPass.cpp
   DeviceGlobals.cpp

llvm/lib/SYCLLowerIR/ModuleSplitter.cpp

@@ -802,6 +802,36 @@ class FunctionsCategorizer {
     Rules.emplace_back(Rule::RKind::K_SortedIntegersListMetadata, MetadataName);
   }
 
+  // Creates a rule, which adds a list of sorted dash-separated integers from
+  // converted into strings listed in a metadata to a resulting identifier.
+  // The form of the metadata is expected to be a metadata node, with its
+  // operands being either an integer or another metadata node with the
+  // form of {!"<aspect_name>", iN <aspect_value>}.
+  void registerAspectListRule(StringRef MetadataName) {
+    registerRule([MetadataName](Function *F) {
+      SmallString<128> Result;
+      if (MDNode *UsedAspects = F->getMetadata(MetadataName)) {
+        SmallVector<std::uint64_t, 8> Values;
+        for (const MDOperand &MDOp : UsedAspects->operands()) {
+          if (auto MDN = dyn_cast<MDNode>(MDOp)) {
+            assert(MDN->getNumOperands() == 2);
+            Values.push_back(mdconst::extract<ConstantInt>(MDN->getOperand(1))
+                                 ->getZExtValue());
+          } else if (auto C = mdconst::dyn_extract<ConstantInt>(MDOp)) {
+            Values.push_back(C->getZExtValue());
+          }
+        }
+
+        llvm::sort(Values);
+
+        for (std::uint64_t V : Values)
+          Result += ("-" + Twine(V)).str();
+      }
+
+      return std::string(Result);
+    });
+  }
+
 private:
   struct Rule {
     struct FlagRuleData {
@@ -980,7 +1010,7 @@ getDeviceCodeSplitter(ModuleDesc &&MD, IRSplitMode Mode, bool IROutputOnly,
     // output files in existing tests.
     Categorizer.registerSimpleStringAttributeRule("sycl-register-alloc-mode");
     Categorizer.registerSimpleStringAttributeRule("sycl-grf-size");
-    Categorizer.registerListOfIntegersInMetadataSortedRule("sycl_used_aspects");
+    Categorizer.registerAspectListRule("sycl_used_aspects");
     Categorizer.registerListOfIntegersInMetadataRule("reqd_work_group_size");
     Categorizer.registerListOfIntegersInMetadataRule("work_group_num_dim");
     Categorizer.registerListOfIntegersInMetadataRule(

llvm/lib/SYCLLowerIR/RecordSYCLAspectNames.cpp

@@ -0,0 +1,70 @@
+//===-------- RecordSYCLAspectNames.cpp - RecordSYCLAspectNames Pass ------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// The !sycl_used_aspects metadata is populated from C++ attributes and
+// further populated by the SYCLPropagateAspectsPass that describes which
+// apsects a function uses. The format of this metadata initially just an
+// integer value corresponding to the enum value in C++. The !sycl_aspects
+// named metadata contains the associations from aspect values to aspect names.
+// These associations are needed later in sycl-post-link, but we drop
+// !sycl_aspects before that to avoid LLVM IR bloat, so this pass takes
+// the associations from !sycl_aspects and then updates all the
+// !sycl_used_aspects metadata to include the aspect names, which allows us
+// to preserve these associations.
+//===----------------------------------------------------------------------===//
+//
+
+#include "llvm/SYCLLowerIR/RecordSYCLAspectNames.h"
+
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/Module.h"
+
+using namespace llvm;
+
+PreservedAnalyses RecordSYCLAspectNamesPass::run(Module &M,
+                                                 ModuleAnalysisManager &MAM) {
+  SmallDenseMap<int64_t, Metadata *, 128> ValueToNameValuePairMD;
+  if (NamedMDNode *Node = M.getNamedMetadata("sycl_aspects")) {
+    for (MDNode *N : Node->operands()) {
+      assert(N->getNumOperands() == 2 &&
+             "Each operand of sycl_aspects must be a pair.");
+
+      // The aspect's integral value is the second operand.
+      auto *C = mdconst::extract<ConstantInt>(N->getOperand(1));
+      ValueToNameValuePairMD[C->getSExtValue()] = N;
+    }
+  }
+
+  auto &Ctx = M.getContext();
+  const char *MetadataToProcess[] = {"sycl_used_aspects",
+                                     "sycl_declared_aspects"};
+  for (Function &F : M.functions()) {
+    for (auto MetadataName : MetadataToProcess) {
+      auto *MDNode = F.getMetadata(MetadataName);
+      if (!MDNode)
+        continue;
+
+      // Change the metadata from {1, 2} to
+      // a format like {{"cpu", 1}, {"gpu", 2}}
+      SmallVector<Metadata *, 8> AspectNameValuePairs;
+      for (const auto &MDOp : MDNode->operands()) {
+        auto *C = mdconst::extract<ConstantInt>(MDOp);
+        int64_t AspectValue = C->getSExtValue();
+        if (auto It = ValueToNameValuePairMD.find(AspectValue);
+            It != ValueToNameValuePairMD.end())
+          AspectNameValuePairs.push_back(It->second);
+        else
+          AspectNameValuePairs.push_back(MDOp);
+      }
+
+      F.setMetadata(MetadataName, MDNode::get(Ctx, AspectNameValuePairs));
+    }
+  }
+
+  return PreservedAnalyses::all();
+}

llvm/lib/SYCLLowerIR/SYCLDeviceRequirements.cpp

@@ -10,6 +10,7 @@
 
 #include "llvm/ADT/SmallString.h"
 #include "llvm/ADT/StringRef.h"
+#include "llvm/IR/Constants.h"
 #include "llvm/IR/Module.h"
 #include "llvm/SYCLLowerIR/ModuleSplitter.h"
 #include "llvm/Support/PropertySetIO.h"
@@ -42,8 +43,15 @@ llvm::computeDeviceRequirements(const module_split::ModuleDesc &MD) {
   // Process all functions in the module
   for (const Function &F : MD.getModule()) {
     if (auto *MDN = F.getMetadata("sycl_used_aspects")) {
-      for (size_t I = 0, E = MDN->getNumOperands(); I < E; ++I) {
-        auto Val = ExtractSignedIntegerFromMDNodeOperand(MDN, I);
+      for (auto &MDOp : MDN->operands()) {
+        int64_t Val;
+        if (auto Pair = dyn_cast<MDNode>(MDOp)) {
+          assert(Pair->getNumOperands() == 2);
+          Val = mdconst::extract<ConstantInt>(Pair->getOperand(1))
+                    ->getZExtValue();
+        } else {
+          Val = mdconst::extract<ConstantInt>(MDOp)->getZExtValue();
+        }
         // Don't put internal aspects (with negative integer value) into the
         // requirements, they are used only for device image splitting.
         if (Val >= 0)

llvm/test/SYCLLowerIR/add-aspect-names.ll

@@ -0,0 +1,75 @@
+; RUN: opt -passes=record-sycl-aspect-names -S < %s | FileCheck %s
+;
+; Basic add-aspect-names functionality test. Checks that 
+; the !sycl_used_apsects metadata is updated from just being
+; integer values to aspect value/name pairs as determined
+; by !sycl_aspects.
+
+; e.g. !sycl_used_aspects = !{i32 1, i32 0}
+; =>   !sycl_used_aspects = !{!{!"B", i32 1}, !{!"A", i32 0}}
+
+; Additionally checks that when there is no association 
+; for a given aspect value, that metadata remains unchanged. 
+
+; e.g. !sycl_used_aspects = !{i32 4, i32 0}
+; =>   !sycl_used_aspects = !{i32 4, !{!"A", i32 0}} 
+
+%A = type { i32 }
+%B = type { i32 }
+%C = type { i32 }
+%D = type { i32 }
+
+; CHECK: funcA() !sycl_used_aspects ![[fA:[0-9]+]]
+define spir_func void @funcA() !sycl_used_aspects !5 {
+  %tmp = alloca %A, align 8
+  ret void
+}
+
+; CHECK: funcB() !sycl_used_aspects ![[fB:[0-9]+]]
+define spir_func void @funcB() !sycl_used_aspects !6 {
+  %tmp = alloca %B, align 8
+  call spir_func void @funcA()
+  ret void
+}
+
+; CHECK: funcC() !sycl_used_aspects ![[fC:[0-9]+]]
+define spir_func void @funcC() !sycl_used_aspects !7 {
+  %tmp = alloca %C, align 8
+  call spir_func void @funcB()
+  ret void
+}
+
+; CHECK: funcD() !sycl_used_aspects ![[fD:[0-9]+]]
+define spir_func void @funcD() !sycl_used_aspects !8 {
+  %tmp = alloca %D, align 8
+  call spir_func void @funcC()
+  ret void
+}
+
+define spir_kernel void @kernel() !sycl_used_aspects !8 !sycl_fixed_targets !9 {
+  call spir_func void @funcD()
+  ret void
+}
+
+!sycl_types_that_use_aspects = !{!0, !1, !2, !3}
+!sycl_aspects = !{!0, !1, !2, !3, !4}
+
+; CHECK-DAG: ![[mA:[0-9]+]] = !{!"A", i32 0}
+; CHECK-DAG: ![[mB:[0-9]+]] = !{!"B", i32 1}
+; CHECK-DAG: ![[mC:[0-9]+]] = !{!"C", i32 2}
+; CHECK-DAG: ![[mD:[0-9]+]] = !{!"D", i32 3}
+; CHECK-DAG: ![[fA]] = !{![[mA]]}
+; CHECK-DAG: ![[fB]] = !{![[mB]], ![[mA]]}
+; CHECK-DAG: ![[fC]] = !{![[mC]], ![[mB]], ![[mA]]}
+; CHECK-DAG: ![[fD]] = !{![[mA]], ![[mB]], ![[mC]], ![[mD]], i32 4}
+
+!0 = !{!"A", i32 0}
+!1 = !{!"B", i32 1}
+!2 = !{!"C", i32 2}
+!3 = !{!"D", i32 3}
+!4 = !{!"fp64", i32 6}
+!5 = !{i32 0}
+!6 = !{i32 1, i32 0}
+!7 = !{i32 2, i32 1, i32 0}
+!8 = !{i32 0, i32 1, i32 2, i32 3, i32 4}
+!9 = !{}

sycl/test/check_device_code/device_has.cpp

@@ -68,16 +68,20 @@ void foo() {
   });
 }
 
-// CHECK-ASPECTS-DAG: [[ASPECTS1]] = !{i32 1}
+// CHECK-ASPECTS-DAG: [[ASPECTS1]] = !{![[ASPECTCPU:[0-9]+]]}
+// CHECK-ASPECTS-DAG: [[ASPECTCPU]] = !{!"cpu", i32 1}
 // CHECK-SRCLOC-DAG: [[SRCLOC1]] = !{i32 {{[0-9]+}}}
 // CHECK-ASPECTS-DAG: [[EMPTYASPECTS]] = !{}
 // CHECK-SRCLOC-DAG: [[SRCLOC2]] = !{i32 {{[0-9]+}}}
-// CHECK-ASPECTS-DAG: [[ASPECTS2]] = !{i32 5, i32 2}
+// CHECK-ASPECTS-DAG: [[ASPECTS2]] = !{![[ASPECTFP16:[0-9]+]], ![[ASPECTGPU:[0-9]+]]}
+// CHECK-ASPECTS-DAG: [[ASPECTFP16]] = !{!"fp16", i32 5}
+// CHECK-ASPECTS-DAG: [[ASPECTGPU]] = !{!"gpu", i32 2}
 // CHECK-SRCLOC-DAG: [[SRCLOC3]] = !{i32 {{[0-9]+}}}
 // CHECK-SRCLOC-DAG: [[SRCLOC4]] = !{i32 {{[0-9]+}}}
-// CHECK-ASPECTS-DAG: [[ASPECTS3]] = !{i32 0}
+// CHECK-ASPECTS-DAG: [[ASPECTS3]] = !{![[ASPECTHOST:[0-9]+]]}
+// CHECK-ASPECTS-DAG: [[ASPECTHOST]] = !{!"host", i32 0}
 // CHECK-SRCLOC-DAG: [[SRCLOC5]] = !{i32 {{[0-9]+}}}
 // CHECK-SRCLOC-DAG: [[SRCLOC6]] = !{i32 {{[0-9]+}}}
 // CHECK-SRCLOC-DAG: [[SRCLOC7]] = !{i32 {{[0-9]+}}}
-// CHECK-ASPECTS-DAG: [[ASPECTS4]] = !{i32 2}
+// CHECK-ASPECTS-DAG: [[ASPECTS4]] = !{![[ASPECTGPU]]}
 // CHECK-SRCLOC-DAG: [[SRCLOC8]] = !{i32 {{[0-9]+}}}

sycl/test/optional_kernel_features/private_alloca.cpp

@@ -16,7 +16,8 @@ class Kernel1;
 // CHECK-LABEL: spir_kernel void @_ZTS7Kernel1
 // CHECK-SAME:      !sycl_used_aspects ![[#USED_ASPECTS:]]
 
-// CHECK:       ![[#USED_ASPECTS]] = !{i32 64}
+// CHECK: ![[#USED_ASPECTS]] = !{![[#ASPECT:]]}
+// CHECK: ![[#ASPECT]] = !{!"ext_oneapi_private_alloca", i32 64}
 
 constexpr static sycl::specialization_id<int> size(10);