diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h b/llvm/include/llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h new file mode 100644 index 0000000000000..5531a36accf44 --- /dev/null +++ b/llvm/include/llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h @@ -0,0 +1,28 @@ +//===---- SYCLPropagateAspectsUsage.cpp - SYCLPropagateAspectsUsage 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 +// +//===----------------------------------------------------------------------===// +// +// Pass propagates optional kernel features metadata through a module call graph +// +//===----------------------------------------------------------------------===// +// +#ifndef LLVM_SYCL_PROPAGATE_ASPECTS_USAGE_H +#define LLVM_SYCL_PROPAGATE_ASPECTS_USAGE_H + +#include "llvm/IR/PassManager.h" + +namespace llvm { + +class SYCLPropagateAspectsUsagePass + : public PassInfoMixin { +public: + PreservedAnalyses run(Module &M, ModuleAnalysisManager &); +}; + +} // namespace llvm + +#endif // LLVM_SYCL_PROPAGATE_ASPECTS_USAGE_H diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp index 4b0c72797f583..45fb62f29b8b4 100644 --- a/llvm/lib/Passes/PassBuilder.cpp +++ b/llvm/lib/Passes/PassBuilder.cpp @@ -84,6 +84,7 @@ #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h" #include "llvm/SYCLLowerIR/LowerWGScope.h" #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h" +#include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/Debug.h" #include "llvm/Support/ErrorHandling.h" diff --git a/llvm/lib/Passes/PassRegistry.def b/llvm/lib/Passes/PassRegistry.def index 0202b23c6aa92..4017d201c68f8 100644 --- a/llvm/lib/Passes/PassRegistry.def +++ b/llvm/lib/Passes/PassRegistry.def @@ -136,6 +136,7 @@ MODULE_PASS("SPIRITTAnnotations", SPIRITTAnnotationsPass()) MODULE_PASS("deadargelim-sycl", DeadArgumentEliminationSYCLPass()) MODULE_PASS("sycllowerwglocalmemory", SYCLLowerWGLocalMemoryPass()) MODULE_PASS("lower-esimd-kernel-attrs", SYCLFixupESIMDKernelWrapperMDPass()) +MODULE_PASS("sycl-propagate-aspects-usage", SYCLPropagateAspectsUsagePass()) #undef MODULE_PASS #ifndef MODULE_PASS_WITH_PARAMS diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index 21a5786bde3d7..973f2aa63b267 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt @@ -58,6 +58,7 @@ add_llvm_component_library(LLVMSYCLLowerIR LowerWGScope.cpp LowerWGLocalMemory.cpp MutatePrintfAddrspace.cpp + SYCLPropagateAspectsUsage.cpp LocalAccessorToSharedMemory.cpp GlobalOffset.cpp diff --git a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp new file mode 100644 index 0000000000000..629cad79a1b9a --- /dev/null +++ b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp @@ -0,0 +1,346 @@ +//===---- SYCLPropagateAspectsUsage.cpp - SYCLPropagateAspectsUsage 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 +// +//===----------------------------------------------------------------------===// +// +// Pass propagates optional kernel features metadata through a module call graph +// +// The pass consists of four main steps: +// +// I. It builds Type -> set of aspects mapping for usage in step II +// II. It builds Function -> set of aspects mapping to use in further steps +// III. FIXME: this step is not yet implemented +// Analyzes aspects usage and emit warnings if necessary +// IV. Generates metadata with information about aspects used by each function +// +// Note: step I is not strictly necessary, because we can simply check if a +// function actually uses one or another type to say whether or not it uses any +// aspects. However, from customers point of view it could be more transparent +// that if a function is declared accepting an optional type, then it means that +// it uses an associated aspect, regardless of whether or not compiler was able +// to optimize out that variable. +// +//===----------------------------------------------------------------------===// + +#include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h" + +#include "llvm/ADT/SetVector.h" +#include "llvm/ADT/SmallPtrSet.h" +#include "llvm/ADT/SmallSet.h" +#include "llvm/IR/Constants.h" +#include "llvm/IR/InstIterator.h" +#include "llvm/IR/Instructions.h" +#include "llvm/IR/IntrinsicInst.h" +#include "llvm/IR/Module.h" +#include "llvm/Pass.h" +#include "llvm/Support/Path.h" + +#include +#include +#include + +using namespace llvm; + +namespace { + +using AspectsSetTy = SmallSet; +using TypeToAspectsMapTy = std::unordered_map; + +/// Retrieves from metadata (intel_types_that_use_aspects) types +/// and aspects these types depend on. +TypeToAspectsMapTy getTypesThatUseAspectsFromMetadata(const Module &M) { + const NamedMDNode *Node = M.getNamedMetadata("intel_types_that_use_aspects"); + TypeToAspectsMapTy Result; + if (!Node) + return Result; + + LLVMContext &C = M.getContext(); + for (const auto OperandIt : Node->operands()) { + const MDNode &N = *OperandIt; + assert(N.getNumOperands() > 1 && "intel_types_that_use_aspect metadata " + "shouldn't contain empty metadata nodes"); + + const auto *TypeName = cast(N.getOperand(0)); + const Type *T = StructType::getTypeByName(C, TypeName->getString()); + assert(T && + "invalid type referenced by intel_types_that_use_aspect metadata"); + + AspectsSetTy &Aspects = Result[T]; + for (size_t I = 1; I != N.getNumOperands(); ++I) { + const auto *CAM = cast(N.getOperand(I)); + const Constant *C = CAM->getValue(); + Aspects.insert(cast(C)->getSExtValue()); + } + } + + return Result; +} + +using TypesEdgesTy = + std::unordered_map>; + +/// Propagates aspects from type @Start to all types which +/// are reachable by edges @Edges by BFS algorithm. +/// Result is recorded in @Aspects. +void propagateAspectsThroughTypes(const TypesEdgesTy &Edges, const Type *Start, + TypeToAspectsMapTy &Aspects) { + const AspectsSetTy &AspectsToPropagate = Aspects[Start]; + SmallSetVector TypesToPropagate; + TypesToPropagate.insert(Start); + for (size_t I = 0; I < TypesToPropagate.size(); ++I) { + const Type *T = TypesToPropagate[I]; + Aspects[T].insert(AspectsToPropagate.begin(), AspectsToPropagate.end()); + const auto It = Edges.find(T); + if (It != Edges.end()) + TypesToPropagate.insert(It->second.begin(), It->second.end()); + } +} + +/// Propagates given aspects to all types in module @M. Function accepts +/// aspects in @TypesWithAspects reference and writes a result in this +/// reference. +/// Type T in the result uses an aspect A if Type T is a composite +/// type (array, struct, vector) which contains elements/fields of +/// another type TT, which in turn uses the aspect A. +/// @TypesWithAspects argument consist of known types with aspects +/// from metadata information. +/// +/// The algorithm is the following: +/// 1) Make a list of all structure types from module @M. The list also +/// contains DoubleTy since it is optional as well. +/// 2) Make from list a type graph which consists of nodes corresponding to +/// types and directed edges between nodes. An edge from type A to type B +/// corresponds to the fact that A is contained within B. +/// Examples: B is a pointer to A, B is a struct containing field of type A. +/// 3) For every known type with aspects propagate it's aspects over graph. +/// Every propagation is a separate run of BFS algorithm. +/// +/// Time complexity: O((V + E) * T) where T is the number of input types +/// containing aspects. +void propagateAspectsToOtherTypesInModule( + const Module &M, TypeToAspectsMapTy &TypesWithAspects) { + std::unordered_set TypesToProcess; + const Type *DoubleTy = Type::getDoubleTy(M.getContext()); + + // 6 is taken from sycl/include/CL/sycl/aspects.hpp + // Note: that magic number must strictly correspond to the one assigned to + // 'fp64' value of 'aspect' enum. + // FIXME: we should develop some kind of mechanism which will allow us to + // avoid hardcoding this number here and having a build dependency between + // the compiler and the runtime. See intel/llvm#5892 + static constexpr int AspectFP64 = 6; + TypesWithAspects[DoubleTy].insert(AspectFP64); + + TypesToProcess.insert(DoubleTy); + for (const Type *T : M.getIdentifiedStructTypes()) + TypesToProcess.insert(T); + + TypesEdgesTy Edges; + for (const Type *T : TypesToProcess) { + for (const Type *TT : T->subtypes()) { + if (TT->isPointerTy()) + // We don't know the pointee type in opaque pointers world + continue; + + // If TT = [4 x [4 x [4 x %A]]] then we want to get TT = %A + // The same with vectors + while (TT->isArrayTy() || TT->isVectorTy()) { + TT = TT->getContainedType(0); + } + + // We are not interested in some types. For example, IntTy. + if (TypesToProcess.count(TT)) + Edges[TT].push_back(T); + } + } + + TypeToAspectsMapTy Result; + for (const Type *T : TypesToProcess) + propagateAspectsThroughTypes(Edges, T, TypesWithAspects); +} + +/// Returns all aspects which might be reached from type @T. +/// It encompases composite structures and pointers. +/// NB! This function inserts new records in @Types map for new discovered +/// types. For the best perfomance pass this map in the next invocations. +const AspectsSetTy &getAspectsFromType(const Type *T, + TypeToAspectsMapTy &Types) { + const auto It = Types.find(T); + if (It != Types.end()) + return It->second; + + // Empty value is inserted for absent key T. + // This is essential to no get into infinite recursive loops. + AspectsSetTy &Result = Types[T]; + + for (const Type *TT : T->subtypes()) { + const AspectsSetTy &Aspects = getAspectsFromType(TT, Types); + Result.insert(Aspects.begin(), Aspects.end()); + } + + return Result; +} + +/// Returns aspects which might be used in instruction @I. +/// Function inspects return type and all operand's types. +/// NB! This function inserts new records in @Types map for new discovered +/// types. For the best perfomance pass this map in the next invocations. +AspectsSetTy getAspectsUsedByInstruction(const Instruction &I, + TypeToAspectsMapTy &Types) { + const Type *ReturnType = I.getType(); + AspectsSetTy Result = getAspectsFromType(ReturnType, Types); + for (const auto &OperandIt : I.operands()) { + const AspectsSetTy &Aspects = + getAspectsFromType(OperandIt->getType(), Types); + Result.insert(Aspects.begin(), Aspects.end()); + } + + return Result; +} + +using FunctionToAspectsMapTy = DenseMap; +using CallGraphTy = DenseMap>; + +void createUsedAspectsMetadataForFunctions(FunctionToAspectsMapTy &Map) { + for (auto &It : Map) { + AspectsSetTy &Aspects = It.second; + if (Aspects.empty()) + continue; + + Function *F = It.first; + LLVMContext &C = F->getContext(); + + SmallVector AspectsMetadata; + for (const auto &A : Aspects) + AspectsMetadata.push_back(ConstantAsMetadata::get( + ConstantInt::getSigned(Type::getInt32Ty(C), A))); + + MDNode *MDN = MDNode::get(C, AspectsMetadata); + F->setMetadata("intel_used_aspects", MDN); + } +} + +/// Propagates aspects from leaves up to the top of call graph. +/// NB! Call graph corresponds to call graph of SYCL code which +/// can't contain recursive calls. So there can't be loops in +/// a call graph. But there can be path's intersections. +void propagateAspectsThroughCG(Function *F, CallGraphTy &CG, + FunctionToAspectsMapTy &AspectsMap, + SmallPtrSet &Visited) { + const auto It = CG.find(F); + if (It == CG.end()) + return; + + AspectsSetTy LocalAspects; + for (Function *Callee : It->second) { + if (Visited.insert(Callee).second) + propagateAspectsThroughCG(Callee, CG, AspectsMap, Visited); + + const auto &CalleeAspects = AspectsMap[Callee]; + LocalAspects.insert(CalleeAspects.begin(), CalleeAspects.end()); + } + + AspectsMap[F].insert(LocalAspects.begin(), LocalAspects.end()); +} + +/// Processes a function: +/// - checks if return and argument types are using any aspects +/// - checks if instructions are using any aspects +/// - updates call graph information +/// +void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToAspects, + TypeToAspectsMapTy &TypesWithAspects, CallGraphTy &CG) { + const AspectsSetTy RetTyAspects = + getAspectsFromType(F.getReturnType(), TypesWithAspects); + FunctionToAspects[&F].insert(RetTyAspects.begin(), RetTyAspects.end()); + for (Argument &Arg : F.args()) { + const AspectsSetTy ArgAspects = + getAspectsFromType(Arg.getType(), TypesWithAspects); + FunctionToAspects[&F].insert(ArgAspects.begin(), ArgAspects.end()); + } + + for (Instruction &I : instructions(F)) { + const AspectsSetTy Aspects = + getAspectsUsedByInstruction(I, TypesWithAspects); + FunctionToAspects[&F].insert(Aspects.begin(), Aspects.end()); + + if (const auto *CI = dyn_cast(&I)) { + if (!CI->isIndirectCall() && CI->getCalledFunction()) + CG[&F].insert(CI->getCalledFunction()); + } + } +} + +// Return true if the function is a SPIRV or SYCL builtin, e.g. +// _Z28__spirv_GlobalInvocationId_xv +// Note: this function was copied from sycl-post-link/ModuleSplitter.cpp and the +// definition of entry point (i.e. implementation of the function) should be in +// sync between those two. +bool isSpirvSyclBuiltin(StringRef FName) { + if (!FName.consume_front("_Z")) + return false; + // now skip the digits + FName = FName.drop_while([](char C) { return std::isdigit(C); }); + + return FName.startswith("__spirv_") || FName.startswith("__sycl_"); +} + +bool isEntryPoint(const Function &F) { + // Skip declarations, we can't analyze them + if (F.isDeclaration()) + return false; + + // Kernels are always considered to be entry points + if (CallingConv::SPIR_KERNEL == F.getCallingConv()) + return true; + + // FIXME: sycl-post-link allows to disable treating SYCL_EXTERNAL's as entry + // points - do we need similar flag here? + // SYCL_EXTERNAL functions with sycl-module-id attribute + // are also considered as entry points (except __spirv_* and __sycl_* + // functions) + return F.hasFnAttribute("sycl-module-id") && !isSpirvSyclBuiltin(F.getName()); +} + +/// Returns a map of functions with corresponding used aspects. +FunctionToAspectsMapTy +buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects) { + FunctionToAspectsMapTy FunctionToAspects; + CallGraphTy CG; + std::vector EntryPoints; + for (Function &F : M.functions()) { + if (F.isDeclaration()) + continue; + + if (isEntryPoint(F)) + EntryPoints.push_back(&F); + + processFunction(F, FunctionToAspects, TypesWithAspects, CG); + } + + SmallPtrSet Visited; + for (Function *F : EntryPoints) + propagateAspectsThroughCG(F, CG, FunctionToAspects, Visited); + + return FunctionToAspects; +} + +} // anonymous namespace + +PreservedAnalyses +SYCLPropagateAspectsUsagePass::run(Module &M, ModuleAnalysisManager &MAM) { + TypeToAspectsMapTy TypesWithAspects = getTypesThatUseAspectsFromMetadata(M); + propagateAspectsToOtherTypesInModule(M, TypesWithAspects); + + FunctionToAspectsMapTy FunctionToAspects = + buildFunctionsToAspectsMap(M, TypesWithAspects); + + createUsedAspectsMetadataForFunctions(FunctionToAspects); + // FIXME: check and diagnose if a function uses an aspect which was not + // declared through [[sycl::device_has()]] attribute + + return PreservedAnalyses::all(); +} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-1.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-1.ll new file mode 100644 index 0000000000000..bd193259b631a --- /dev/null +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-1.ll @@ -0,0 +1,55 @@ +; RUN: opt -passes=sycl-propagate-aspects-usage < %s -S | FileCheck %s +; +; Test checks that the pass is able to propagate information about used aspects +; through a call graph +; +; K1 K2 +; / \/ \ +; F1 F2 F3 +; +; F1 doesn't use optional type. +; F2 uses optional A. +; F3 uses optional B. + +%Optional.A = type { i32 } +%Optional.B = type { i32 } + +; CHECK: spir_kernel void @kernel1() !intel_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() !intel_used_aspects ![[#ID2:]] { +define spir_kernel void @kernel2() { + call spir_func void @func2() + call spir_func void @func3() + ret void +} + +; CHECK: spir_func void @func1() { +define spir_func void @func1() { + %tmp = alloca i32 + ret void +} + +; CHECK: spir_func void @func2() !intel_used_aspects ![[#ID1]] { +define spir_func void @func2() { + %tmp = alloca %Optional.A + ret void +} + +; CHECK: spir_func void @func3() !intel_used_aspects ![[#ID3:]] { +define spir_func void @func3() { + %tmp = alloca %Optional.B + ret void +} + +!intel_types_that_use_aspects = !{!0, !1} +!0 = !{!"Optional.A", i32 1} +!1 = !{!"Optional.B", i32 2} + +; CHECK: ![[#ID1]] = !{i32 1} +; CHECK: ![[#ID2]] = !{i32 1, i32 2} +; CHECK: ![[#ID3]] = !{i32 2} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-2.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-2.ll new file mode 100644 index 0000000000000..cf236861ad2ee --- /dev/null +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-2.ll @@ -0,0 +1,56 @@ +; RUN: opt -passes=sycl-propagate-aspects-usage < %s -S | FileCheck %s + +; Test checks that the pass is able to propagate information about used aspects +; through a call graph +; +; K +; / \ +; F1 F2 +; \ / \ +; F3 F4 +; +; F3 uses optional A. +; F4 uses optional B. + +%Optional.A = type { i32 } +%Optional.B = type { i32 } + +; CHECK: spir_kernel void @kernel() !intel_used_aspects ![[#ID1:]] { +define spir_kernel void @kernel() { + call spir_func void @func1() + call spir_func void @func2() + ret void +} + +; CHECK: spir_func void @func1() !intel_used_aspects ![[#ID2:]] { +define spir_func void @func1() { + call spir_func void @func3() + ret void +} + +; CHECK: spir_func void @func2() !intel_used_aspects ![[#ID1]] { +define spir_func void @func2() { + call spir_func void @func3() + call spir_func void @func4() + ret void +} + +; CHECK: spir_func void @func3() !intel_used_aspects ![[#ID2]] { +define spir_func void @func3() { + %tmp = alloca %Optional.A + ret void +} + +; CHECK: spir_func void @func4() !intel_used_aspects ![[#ID3:]] { +define spir_func void @func4() { + %tmp = alloca %Optional.B + ret void +} + +!intel_types_that_use_aspects = !{!0, !1} +!0 = !{!"Optional.A", i32 1} +!1 = !{!"Optional.B", i32 2} + +; CHECK: ![[#ID1]] = !{i32 1, i32 2} +; CHECK: ![[#ID2]] = !{i32 1} +; CHECK: ![[#ID3]] = !{i32 2} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/composite-types-1.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/composite-types-1.ll new file mode 100644 index 0000000000000..d0e92ebf5bf91 --- /dev/null +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/composite-types-1.ll @@ -0,0 +1,106 @@ +; RUN: opt -passes=sycl-propagate-aspects-usage < %s -S | FileCheck %s +; +; Test checks that the pass is able to propagate information about used aspects +; from simple composite types to functions and kernels which use them + +; Optional +%A.optional = type { i32 } + +; Not optional +%B.core = type { i32 } + +; Not optional +%C.core = type { i32 } + +%D1.contains.optional = type { %A.optional, %B.core, %C.core } + +%D2.does.not.contain.optional = type { %B.core, %C.core } + +%E.contains.optional = type { %B.core, %C.core, %D1.contains.optional } + +%F1.points.to.optional = type { %B.core, %C.core*, %D1.contains.optional* } + +%F2.does.not.contain.optional = type { %B.core, %C.core*, %D2.does.not.contain.optional* } + +; CHECK: spir_kernel void @kernelD1.uses.optional() !intel_used_aspects !1 { +define spir_kernel void @kernelD1.uses.optional() { + %tmp = alloca %D1.contains.optional + ret void +} + +; CHECK: spir_func void @funcD1.uses.optional() !intel_used_aspects !1 { +define spir_func void @funcD1.uses.optional() { + %tmp = alloca %D1.contains.optional + ret void +} + +; 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 +} + +; CHECK: spir_func void @funcD2.does.not.use.optional() { +define spir_func void @funcD2.does.not.use.optional() { + %tmp = alloca %D2.does.not.contain.optional + ret void +} + +; CHECK: spir_kernel void @kernelE.uses.optional() !intel_used_aspects !1 { +define spir_kernel void @kernelE.uses.optional() { + %tmp = alloca %E.contains.optional + ret void +} + +; CHECK: spir_func void @funcE.uses.optional() !intel_used_aspects !1 { +define spir_func void @funcE.uses.optional() { + %tmp = alloca %E.contains.optional + ret void +} + +; CHECK: spir_kernel void @kernelF1.points.to.optional() { +define spir_kernel void @kernelF1.points.to.optional() { + %tmp = alloca %F1.points.to.optional + ret void +} + +; CHECK: spir_func void @funcF1.points.to.optional() { +define spir_func void @funcF1.points.to.optional() { + %tmp = alloca %F1.points.to.optional + ret void +} + +; 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 +} + +; CHECK: spir_func void @funcF2.does.not.use.optional() { +define spir_func void @funcF2.does.not.use.optional() { + %tmp = alloca %F2.does.not.contain.optional + ret void +} + +; CHECK: spir_func %A.optional @funcA.returns.optional() !intel_used_aspects !1 { +define spir_func %A.optional @funcA.returns.optional() { + %tmp = alloca %A.optional + %ret = load %A.optional, %A.optional* %tmp + ret %A.optional %ret +} + +; CHECK: spir_func void @funcA.uses.array.of.optional() !intel_used_aspects !1 { +define spir_func void @funcA.uses.array.of.optional() { + %tmp = alloca [4 x %A.optional] + ret void +} + +; CHECK: spir_func void @funcA.assepts.optional(%A.optional %0) !intel_used_aspects !1 { +define spir_func void @funcA.assepts.optional(%A.optional %0) { + ret void +} + +!intel_types_that_use_aspects = !{!0} +!0 = !{!"A.optional", i32 1} + +; CHECK: !1 = !{i32 1} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/double.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/double.ll new file mode 100644 index 0000000000000..23bec94282ed7 --- /dev/null +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/double.ll @@ -0,0 +1,37 @@ +; RUN: opt -passes=sycl-propagate-aspects-usage < %s -S | FileCheck %s +; +; Test checks that double's aspect is spotted and propagated. + +%composite = type { double } + +; CHECK: spir_kernel void @kernel() !intel_used_aspects !0 { +define spir_kernel void @kernel() { + call spir_func void @func() + ret void +} + +; CHECK: spir_func void @func() !intel_used_aspects !0 { +define spir_func void @func() { + %tmp = alloca double + ret void +} + +; CHECK: spir_func void @func.array() !intel_used_aspects !0 { +define spir_func void @func.array() { + %tmp = alloca [4 x double] + ret void +} + +; CHECK: spir_func void @func.vector() !intel_used_aspects !0 { +define spir_func void @func.vector() { + %tmp = alloca <4 x double> + ret void +} + +; CHECK: spir_func void @func.composite() !intel_used_aspects !0 { +define spir_func void @func.composite() { + %tmp = alloca %composite + ret void +} + +; CHECK: !0 = !{i32 6} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/multiple-aspects.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/multiple-aspects.ll new file mode 100644 index 0000000000000..d7622f911b486 --- /dev/null +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/multiple-aspects.ll @@ -0,0 +1,52 @@ +; RUN: opt -passes=sycl-propagate-aspects-usage -S < %s | FileCheck %s +; +; Test checks that the pass is able to collect all aspects used in a function + +%A = type { i32 } +%B = type { i32 } +%C = type { i32 } +%D = type { i32 } + +; CHECK: define spir_func void @funcA() !intel_used_aspects ![[#ID0:]] { +define spir_func void @funcA() { + %tmp = alloca %A + ret void +} + +; CHECK: define spir_func void @funcB() !intel_used_aspects ![[#ID1:]] { +define spir_func void @funcB() { + %tmp = alloca %B + call spir_func void @funcA() + ret void +} + +; CHECK: define spir_func void @funcC() !intel_used_aspects ![[#ID2:]] { +define spir_func void @funcC() { + %tmp = alloca %C + call spir_func void @funcB() + ret void +} + +; CHECK: define spir_func void @funcD() !intel_used_aspects ![[#ID3:]] { +define spir_func void @funcD() { + %tmp = alloca %D + call spir_func void @funcC() + ret void +} + +; CHECK: define spir_kernel void @kernel() !intel_used_aspects ![[#ID3]] { +define spir_kernel void @kernel() { + call spir_func void @funcD() + ret void +} + +!intel_types_that_use_aspects = !{!0, !1, !2, !3} +!0 = !{!"A", i32 0} +!1 = !{!"B", i32 1} +!2 = !{!"C", i32 2} +!3 = !{!"D", i32 3, i32 4} + +; CHECK: ![[#ID0]] = !{i32 0} +; CHECK: ![[#ID1]] = !{i32 1, i32 0} +; CHECK: ![[#ID2]] = !{i32 2, i32 1, i32 0} +; CHECK: ![[#ID3]] = !{i32 0, i32 1, i32 2, i32 3, i32 4} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/no-uses-of-optional.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/no-uses-of-optional.ll new file mode 100644 index 0000000000000..967accbb3e3bc --- /dev/null +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/no-uses-of-optional.ll @@ -0,0 +1,20 @@ +; RUN: opt -passes=sycl-propagate-aspects-usage < %s -S | FileCheck %s --implicit-check-not "!intel_used_aspects" +; +; Test checks that no metadata propagates because MyStruct +; isn't used inside functions. + +%MyStruct = type { i32 } + +; CHECK: dso_local spir_kernel void @kernel() { +define dso_local spir_kernel void @kernel() { + call spir_func void @func() + ret void +} + +; CHECK: dso_local spir_func void @func() { +define weak dso_local spir_func void @func() { + ret void +} + +!intel_types_that_use_aspects = !{!0} +!0 = !{!"MyStruct", i32 1}