diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index c40010e2c1071..ccdede57104dc 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/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()); diff --git a/llvm/include/llvm/SYCLLowerIR/RecordSYCLAspectNames.h b/llvm/include/llvm/SYCLLowerIR/RecordSYCLAspectNames.h new file mode 100644 index 0000000000000..fa619136c2bb9 --- /dev/null +++ b/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 { +public: + PreservedAnalyses run(Module &M, ModuleAnalysisManager &); +}; + +} // namespace llvm + +#endif // LLVM_RECORD_SYCL_ASPECT_NAMES diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp index 5856d65175318..10dea4f0196b2 100644 --- a/llvm/lib/Passes/PassBuilder.cpp +++ b/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" diff --git a/llvm/lib/Passes/PassRegistry.def b/llvm/lib/Passes/PassRegistry.def index 758afe3697ea9..a39b521d446c7 100644 --- a/llvm/lib/Passes/PassRegistry.def +++ b/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 diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index 54946bea7da9e..7f2edaae323a9 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/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 diff --git a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp b/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp index 7b2e7ac0f978a..cf41aee46df28 100644 --- a/llvm/lib/SYCLLowerIR/ModuleSplitter.cpp +++ b/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 {!"", iN }. + void registerAspectListRule(StringRef MetadataName) { + registerRule([MetadataName](Function *F) { + SmallString<128> Result; + if (MDNode *UsedAspects = F->getMetadata(MetadataName)) { + SmallVector Values; + for (const MDOperand &MDOp : UsedAspects->operands()) { + if (auto MDN = dyn_cast(MDOp)) { + assert(MDN->getNumOperands() == 2); + Values.push_back(mdconst::extract(MDN->getOperand(1)) + ->getZExtValue()); + } else if (auto C = mdconst::dyn_extract(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( diff --git a/llvm/lib/SYCLLowerIR/RecordSYCLAspectNames.cpp b/llvm/lib/SYCLLowerIR/RecordSYCLAspectNames.cpp new file mode 100644 index 0000000000000..328560ab10fe3 --- /dev/null +++ b/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 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(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 AspectNameValuePairs; + for (const auto &MDOp : MDNode->operands()) { + auto *C = mdconst::extract(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(); +} \ No newline at end of file diff --git a/llvm/lib/SYCLLowerIR/SYCLDeviceRequirements.cpp b/llvm/lib/SYCLLowerIR/SYCLDeviceRequirements.cpp index 0b15feecca143..8ebec7f54013d 100644 --- a/llvm/lib/SYCLLowerIR/SYCLDeviceRequirements.cpp +++ b/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(MDOp)) { + assert(Pair->getNumOperands() == 2); + Val = mdconst::extract(Pair->getOperand(1)) + ->getZExtValue(); + } else { + Val = mdconst::extract(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) diff --git a/llvm/test/SYCLLowerIR/add-aspect-names.ll b/llvm/test/SYCLLowerIR/add-aspect-names.ll new file mode 100644 index 0000000000000..bc36d0938fd22 --- /dev/null +++ b/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 = !{} diff --git a/sycl/test/check_device_code/device_has.cpp b/sycl/test/check_device_code/device_has.cpp index cb56d38b91bec..92dbcb795a4b4 100644 --- a/sycl/test/check_device_code/device_has.cpp +++ b/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]+}}} diff --git a/sycl/test/optional_kernel_features/private_alloca.cpp b/sycl/test/optional_kernel_features/private_alloca.cpp index 3652f8c21d110..bf3e0a9902297 100644 --- a/sycl/test/optional_kernel_features/private_alloca.cpp +++ b/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 size(10);