From e8c3e3262c29641dd64c243f0648c711f50dd13c Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 16 Nov 2022 07:52:40 -0800 Subject: [PATCH 1/8] [SYCL] Merge sycl_declared_aspects to sycl_used_aspects To make RT be able to throw an exception if device doesn't support the feature pass should also propagate sycl_declared aspects by merging it with sycl_used_aspects. Spec: https://github.com/intel/llvm/pull/7415 --- .../SYCLLowerIR/SYCLPropagateAspectsUsage.cpp | 25 +++++++-- .../propagate-declared-1.ll | 49 ++++++++++++++++++ .../propagate-declared-2.ll | 43 ++++++++++++++++ .../propagate-declared-3.ll | 51 +++++++++++++++++++ 4 files changed, 163 insertions(+), 5 deletions(-) create mode 100644 llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-1.ll create mode 100644 llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-2.ll create mode 100644 llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-3.ll diff --git a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp index 134711b1279f5..f68778afac5c5 100644 --- a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp @@ -410,6 +410,7 @@ void propagateAspectsThroughCG(Function *F, CallGraphTy &CG, /// - checks if function has "!sycl_used_aspects" metadata /// void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToAspects, + FunctionToAspectsMapTy &FunctionsWithDeclaredAspects, TypeToAspectsMapTy &TypesWithAspects, CallGraphTy &CG) { const AspectsSetTy RetTyAspects = getAspectsFromType(F.getReturnType(), TypesWithAspects); @@ -431,15 +432,18 @@ void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToAspects, } } - if (F.hasMetadata("sycl_used_aspects")) { - const MDNode *MD = F.getMetadata("sycl_used_aspects"); + auto CollectAspectsFromMD = [&F](const char* MDName, FunctionToAspectsMapTy &Map) { + if (const MDNode *MD = F.getMetadata(MDName)) { AspectsSetTy Aspects; for (const MDOperand &Op : MD->operands()) { Constant *C = cast(Op.get())->getValue(); Aspects.insert(cast(C)->getSExtValue()); } - FunctionToAspects[&F].insert(Aspects.begin(), Aspects.end()); - } + Map[&F].insert(Aspects.begin(), Aspects.end()); + } + }; + CollectAspectsFromMD("sycl_used_aspects", FunctionToAspects); + CollectAspectsFromMD("sycl_declared_aspects", FunctionsWithDeclaredAspects); } // Return true if the function is a SPIRV or SYCL builtin, e.g. @@ -504,12 +508,17 @@ buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects, const AspectValueToNameMapTy &AspectValues, const std::vector &EntryPoints) { FunctionToAspectsMapTy FunctionToAspects; + // The set of spects from FunctionsWithDeclaredAspects should be merged to the set + // of FunctionToAspects after validateUsedAspectsForFunctions call to avoid + // errors during validation. + FunctionToAspectsMapTy FunctionsWithDeclaredAspects; CallGraphTy CG; for (Function &F : M.functions()) { if (F.isDeclaration()) continue; - processFunction(F, FunctionToAspects, TypesWithAspects, CG); + processFunction(F, FunctionToAspects, FunctionsWithDeclaredAspects, + TypesWithAspects, CG); } SmallPtrSet Visited; @@ -519,6 +528,12 @@ buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects, validateUsedAspectsForFunctions(FunctionToAspects, AspectValues, EntryPoints, CG); + Visited.clear(); + for (Function *F : EntryPoints) + propagateAspectsThroughCG(F, CG, FunctionsWithDeclaredAspects, Visited); + for (const auto &It : FunctionsWithDeclaredAspects) + FunctionToAspects[It.first].insert(It.second.begin(), It.second.end()); + return FunctionToAspects; } diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-1.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-1.ll new file mode 100644 index 0000000000000..2ae02f01baa89 --- /dev/null +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-1.ll @@ -0,0 +1,49 @@ +; RUN: opt -passes=sycl-propagate-aspects-usage %s -S | FileCheck %s + +; kernel() +; | +; v +; baz() +; | +; v +; bar() +; | +; v +; foo() + +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_used_aspects ![[#ASPECT:]] +define weak_odr dso_local spir_kernel void @kernel() { +entry: + call spir_func void @_Z3bazv() + ret void +} + +; CHECK: void @_Z3bazv() !sycl_used_aspects ![[#ASPECT]] { +define dso_local spir_func void @_Z3bazv() { +entry: + call spir_func void @_Z3barv() + ret void +} + +; CHECK: void @_Z3barv() !sycl_used_aspects ![[#ASPECT]] { +define dso_local spir_func void @_Z3barv() { +entry: + call spir_func void @_Z3foov() + ret void +} + +; CHECK: void @_Z3foov() !sycl_declared_aspects ![[#ASPECT]] !sycl_used_aspects ![[#ASPECT]] { +define dso_local spir_func void @_Z3foov() !sycl_declared_aspects !2 { +entry: + ret void +} + +!sycl_aspects = !{!0, !1} + +!0 = !{!"gpu", i32 2} +!1 = !{!"fp64", i32 6} +!2 = !{i32 2} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-2.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-2.ll new file mode 100644 index 0000000000000..331c708743288 --- /dev/null +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-2.ll @@ -0,0 +1,43 @@ +; RUN: opt -passes=sycl-propagate-aspects-usage %s -S | FileCheck %s + +; baz() +; / \ +; v v +; bar() foo() + +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:]] +define dso_local spir_kernel void @_Z3bazv() { +entry: + call spir_func void @_Z3barv() + call spir_func void @_Z3foov() + ret void +} + +; CHECK: void @_Z3barv() !sycl_used_aspects ![[#ASPECT2:]] { +define dso_local spir_func void @_Z3barv() !sycl_used_aspects !3 { +entry: + ret void +} + +; CHECK: void @_Z3foov() !sycl_used_aspects ![[#ASPECT3:]] +; CHECK-SAME: !sycl_declared_aspects ![[#ASPECT3]] { +define dso_local spir_func void @_Z3foov() !sycl_declared_aspects !4 { +entry: + ret void +} + +; CHECK: ![[#ASPECT1]] = !{i32 2, i32 1} +; CHECK: ![[#ASPECT2]] = !{i32 2} +; CHECK: ![[#ASPECT3]] = !{i32 1} + +!sycl_aspects = !{!0, !1, !2} + +!0 = !{!"cpu", i32 1} +!1 = !{!"gpu", i32 2} +!2 = !{!"fp64", i32 6} +!3 = !{i32 2} +!4 = !{i32 1} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-3.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-3.ll new file mode 100644 index 0000000000000..663d43ae15ced --- /dev/null +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-3.ll @@ -0,0 +1,51 @@ +; RUN: opt -passes=sycl-propagate-aspects-usage %s -S | FileCheck %s + +; K +; / \ +; F1 F2 +; \ / \ +; F3 F4 + +; CHECK: spir_kernel void @kernel() !sycl_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() !sycl_used_aspects ![[#ID2:]] { +define spir_func void @func1() { + call spir_func void @func3() + ret void +} + +; CHECK: spir_func void @func2() !sycl_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() !sycl_used_aspects ![[#ID2]] { +define spir_func void @func3() !sycl_used_aspects !4 { + ret void +} + +; CHECK: spir_func void @func4() !sycl_used_aspects ![[#ID3:]] +; CHECK-SAME: !sycl_declared_aspects ![[#ID3]] { +define spir_func void @func4() !sycl_declared_aspects !3 { + ret void +} + +!sycl_aspects = !{!0, !1, !2} + +!0 = !{!"host", i32 0} +!1 = !{!"cpu", i32 1} +!2 = !{!"fp64", i32 6} +!3 = !{i32 0} +!4 = !{i32 1} +!5 = !{i32 0, i32 1} + +; CHECK: ![[#ID1]] = !{i32 1, i32 0} +; CHECK: ![[#ID2]] = !{i32 1} +; CHECK: ![[#ID3]] = !{i32 0} From db974013384e404716f22f0f08c2309ba2d0d1f4 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 16 Nov 2022 08:37:24 -0800 Subject: [PATCH 2/8] Fix test --- clang/test/CodeGenSYCL/device_has.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/test/CodeGenSYCL/device_has.cpp b/clang/test/CodeGenSYCL/device_has.cpp index 2f7ea8f28e5eb..6fe5b0d030a6d 100644 --- a/clang/test/CodeGenSYCL/device_has.cpp +++ b/clang/test/CodeGenSYCL/device_has.cpp @@ -8,25 +8,25 @@ queue q; // CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]] -// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] { +// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] [[sycl::device_has(sycl::aspect::cpu)]] void func1() {} -// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] { +// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] [[sycl::device_has(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {} -// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] { +// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] [[sycl::device_has()]] void func3() {} -// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] { +// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] template [[sycl::device_has(Aspect)]] void func4() {} -// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] { +// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] [[sycl::device_has(sycl::aspect::cpu)]] void func5(); void func5() {} constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; } -// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] { +// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] [[sycl::device_has(getAspect())]] void func6() {} class KernelFunctor { From ae09ca51ce4ecb0249d0156c26b7c766cf03df27 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 16 Nov 2022 08:49:06 -0800 Subject: [PATCH 3/8] Apply suggestions --- .../SYCLLowerIR/SYCLPropagateAspectsUsage.cpp | 46 +++++++++---------- 1 file changed, 23 insertions(+), 23 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp index f68778afac5c5..cba626bd0e649 100644 --- a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp @@ -407,24 +407,24 @@ void propagateAspectsThroughCG(Function *F, CallGraphTy &CG, /// - checks if return and argument types are using any aspects /// - checks if instructions are using any aspects /// - updates call graph information -/// - checks if function has "!sycl_used_aspects" metadata -/// -void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToAspects, - FunctionToAspectsMapTy &FunctionsWithDeclaredAspects, +/// - checks if function has "!sycl_used_aspects" and "!sycl_declared_aspects" +/// metadata and if so collects aspects from this metadata +void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToUsedAspects, + FunctionToAspectsMapTy &FunctionToDeclaredAspects, TypeToAspectsMapTy &TypesWithAspects, CallGraphTy &CG) { const AspectsSetTy RetTyAspects = getAspectsFromType(F.getReturnType(), TypesWithAspects); - FunctionToAspects[&F].insert(RetTyAspects.begin(), RetTyAspects.end()); + FunctionToUsedAspects[&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()); + FunctionToUsedAspects[&F].insert(ArgAspects.begin(), ArgAspects.end()); } for (Instruction &I : instructions(F)) { const AspectsSetTy Aspects = getAspectsUsedByInstruction(I, TypesWithAspects); - FunctionToAspects[&F].insert(Aspects.begin(), Aspects.end()); + FunctionToUsedAspects[&F].insert(Aspects.begin(), Aspects.end()); if (const auto *CI = dyn_cast(&I)) { if (!CI->isIndirectCall() && CI->getCalledFunction()) @@ -442,8 +442,8 @@ void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToAspects, Map[&F].insert(Aspects.begin(), Aspects.end()); } }; - CollectAspectsFromMD("sycl_used_aspects", FunctionToAspects); - CollectAspectsFromMD("sycl_declared_aspects", FunctionsWithDeclaredAspects); + CollectAspectsFromMD("sycl_used_aspects", FunctionToUsedAspects); + CollectAspectsFromMD("sycl_declared_aspects", FunctionToDeclaredAspects); } // Return true if the function is a SPIRV or SYCL builtin, e.g. @@ -507,34 +507,34 @@ FunctionToAspectsMapTy buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects, const AspectValueToNameMapTy &AspectValues, const std::vector &EntryPoints) { - FunctionToAspectsMapTy FunctionToAspects; - // The set of spects from FunctionsWithDeclaredAspects should be merged to the set - // of FunctionToAspects after validateUsedAspectsForFunctions call to avoid - // errors during validation. - FunctionToAspectsMapTy FunctionsWithDeclaredAspects; + FunctionToAspectsMapTy FunctionToUsedAspects; + FunctionToAspectsMapTy FunctionToDeclaredAspects; CallGraphTy CG; for (Function &F : M.functions()) { if (F.isDeclaration()) continue; - processFunction(F, FunctionToAspects, FunctionsWithDeclaredAspects, + processFunction(F, FunctionToUsedAspects, FunctionToDeclaredAspects, TypesWithAspects, CG); } SmallPtrSet Visited; for (Function *F : EntryPoints) - propagateAspectsThroughCG(F, CG, FunctionToAspects, Visited); + propagateAspectsThroughCG(F, CG, FunctionToUsedAspects, Visited); - validateUsedAspectsForFunctions(FunctionToAspects, AspectValues, EntryPoints, + validateUsedAspectsForFunctions(FunctionToUsedAspects, AspectValues, EntryPoints, CG); + // The set of spects from FunctionToDeclaredAspects should be merged to the set + // of FunctionToUsedAspects after validateUsedAspectsForFunctions call to avoid + // errors during validation. Visited.clear(); for (Function *F : EntryPoints) - propagateAspectsThroughCG(F, CG, FunctionsWithDeclaredAspects, Visited); - for (const auto &It : FunctionsWithDeclaredAspects) - FunctionToAspects[It.first].insert(It.second.begin(), It.second.end()); + propagateAspectsThroughCG(F, CG, FunctionToDeclaredAspects, Visited); + for (const auto &It : FunctionToDeclaredAspects) + FunctionToUsedAspects[It.first].insert(It.second.begin(), It.second.end()); - return FunctionToAspects; + return FunctionToUsedAspects; } } // anonymous namespace @@ -565,10 +565,10 @@ SYCLPropagateAspectsUsagePass::run(Module &M, ModuleAnalysisManager &MAM) { propagateAspectsToOtherTypesInModule(M, TypesWithAspects, AspectValues); - FunctionToAspectsMapTy FunctionToAspects = buildFunctionsToAspectsMap( + FunctionToAspectsMapTy FunctionToUsedAspects = buildFunctionsToAspectsMap( M, TypesWithAspects, AspectValues, EntryPoints); - createUsedAspectsMetadataForFunctions(FunctionToAspects); + createUsedAspectsMetadataForFunctions(FunctionToUsedAspects); setSyclFixedTargetsMD(EntryPoints, TargetFixedAspects, AspectValues); From 4c9f301486b360a67fd6a0bf288113bd81ae458e Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 16 Nov 2022 08:51:11 -0800 Subject: [PATCH 4/8] Format --- llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp index cba626bd0e649..4ae72af877270 100644 --- a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp @@ -522,12 +522,12 @@ buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects, for (Function *F : EntryPoints) propagateAspectsThroughCG(F, CG, FunctionToUsedAspects, Visited); - validateUsedAspectsForFunctions(FunctionToUsedAspects, AspectValues, EntryPoints, - CG); + validateUsedAspectsForFunctions(FunctionToUsedAspects, AspectValues, + EntryPoints, CG); - // The set of spects from FunctionToDeclaredAspects should be merged to the set - // of FunctionToUsedAspects after validateUsedAspectsForFunctions call to avoid - // errors during validation. + // The set of spects from FunctionToDeclaredAspects should be merged to the + // set of FunctionToUsedAspects after validateUsedAspectsForFunctions call to + // avoid errors during validation. Visited.clear(); for (Function *F : EntryPoints) propagateAspectsThroughCG(F, CG, FunctionToDeclaredAspects, Visited); From 435c611240986b9f36aa15651a3b8b997038a2bd Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 16 Nov 2022 09:09:00 -0800 Subject: [PATCH 5/8] Adjust FE test --- clang/test/CodeGenSYCL/device_has.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/device_has.cpp b/clang/test/CodeGenSYCL/device_has.cpp index 6fe5b0d030a6d..499e907541276 100644 --- a/clang/test/CodeGenSYCL/device_has.cpp +++ b/clang/test/CodeGenSYCL/device_has.cpp @@ -9,24 +9,29 @@ queue q; // CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]] // CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] +// CHECK-SAME: !sycl_used_aspects [[sycl::device_has(sycl::aspect::cpu)]] void func1() {} // CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] +// CHECK-SAME: !sycl_used_aspects [[sycl::device_has(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {} -// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] +// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] { [[sycl::device_has()]] void func3() {} // CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] +// CHECK-SAME: !sycl_used_aspects template [[sycl::device_has(Aspect)]] void func4() {} // CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] +// CHECK-SAME: !sycl_used_aspects [[sycl::device_has(sycl::aspect::cpu)]] void func5(); void func5() {} constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; } // CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] +// CHECK-SAME: !sycl_used_aspects [[sycl::device_has(getAspect())]] void func6() {} class KernelFunctor { From 586281c43ceccbe4bb21c09437de47228035a6d3 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Thu, 17 Nov 2022 02:25:16 -0800 Subject: [PATCH 6/8] Adjust comment --- clang/test/CodeGenSYCL/device_has.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/device_has.cpp b/clang/test/CodeGenSYCL/device_has.cpp index 499e907541276..14b19f4524c9b 100644 --- a/clang/test/CodeGenSYCL/device_has.cpp +++ b/clang/test/CodeGenSYCL/device_has.cpp @@ -1,6 +1,7 @@ // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s -// Tests for IR of device_has(aspect, ...) attribute +// Tests for IR of device_has(aspect, ...) attribute and +// !sycl_used_aspects metadata #include "sycl.hpp" using namespace sycl; From e647f365aa434c9690cab26add5933ab8d860e91 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Thu, 17 Nov 2022 02:27:55 -0800 Subject: [PATCH 7/8] Adjust CHECK-SAME lines --- clang/test/CodeGenSYCL/device_has.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/clang/test/CodeGenSYCL/device_has.cpp b/clang/test/CodeGenSYCL/device_has.cpp index 14b19f4524c9b..3b626a3516c76 100644 --- a/clang/test/CodeGenSYCL/device_has.cpp +++ b/clang/test/CodeGenSYCL/device_has.cpp @@ -10,29 +10,29 @@ queue q; // CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]] // CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] -// CHECK-SAME: !sycl_used_aspects +// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]] [[sycl::device_has(sycl::aspect::cpu)]] void func1() {} // CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] -// CHECK-SAME: !sycl_used_aspects +// CHECK-SAME: !sycl_used_aspects ![[ASPECTS2]] [[sycl::device_has(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {} // CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] { [[sycl::device_has()]] void func3() {} // CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] -// CHECK-SAME: !sycl_used_aspects +// CHECK-SAME: !sycl_used_aspects ![[ASPECTS3]] template [[sycl::device_has(Aspect)]] void func4() {} // CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] -// CHECK-SAME: !sycl_used_aspects +// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]] [[sycl::device_has(sycl::aspect::cpu)]] void func5(); void func5() {} constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; } // CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] -// CHECK-SAME: !sycl_used_aspects +// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]] [[sycl::device_has(getAspect())]] void func6() {} class KernelFunctor { From 6ec407f7a9c82f820a66261d6da505b4451fef90 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Thu, 17 Nov 2022 07:52:55 -0800 Subject: [PATCH 8/8] Fix format --- llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp index 4ae72af877270..b13e480170fec 100644 --- a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp @@ -434,12 +434,12 @@ void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToUsedAspects, auto CollectAspectsFromMD = [&F](const char* MDName, FunctionToAspectsMapTy &Map) { if (const MDNode *MD = F.getMetadata(MDName)) { - AspectsSetTy Aspects; - for (const MDOperand &Op : MD->operands()) { - Constant *C = cast(Op.get())->getValue(); - Aspects.insert(cast(C)->getSExtValue()); - } - Map[&F].insert(Aspects.begin(), Aspects.end()); + AspectsSetTy Aspects; + for (const MDOperand &Op : MD->operands()) { + Constant *C = cast(Op.get())->getValue(); + Aspects.insert(cast(C)->getSExtValue()); + } + Map[&F].insert(Aspects.begin(), Aspects.end()); } }; CollectAspectsFromMD("sycl_used_aspects", FunctionToUsedAspects); @@ -525,7 +525,7 @@ buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects, validateUsedAspectsForFunctions(FunctionToUsedAspects, AspectValues, EntryPoints, CG); - // The set of spects from FunctionToDeclaredAspects should be merged to the + // The set of aspects from FunctionToDeclaredAspects should be merged to the // set of FunctionToUsedAspects after validateUsedAspectsForFunctions call to // avoid errors during validation. Visited.clear();