From 136bc192e622d1d80611a171292c393a37ddeebe Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 20 Sep 2022 04:18:35 -0700 Subject: [PATCH] [SYCL] Rename optional device feature metadata This commit renames metadata nodes introduced by the optional device features design document from using `intel_` prefix to use `sycl_` prefix. Signed-off-by: Larsen, Steffen --- clang/lib/CodeGen/CodeGenFunction.cpp | 4 +- clang/lib/CodeGen/CodeGenModule.cpp | 2 +- clang/test/CodeGenSYCL/device_has.cpp | 16 ++-- clang/test/CodeGenSYCL/uses_aspects.cpp | 16 ++-- .../SYCLLowerIR/SYCLPropagateAspectsUsage.cpp | 6 +- .../PropagateAspectsUsage/call-graph-1.ll | 10 +-- .../PropagateAspectsUsage/call-graph-2.ll | 12 +-- .../composite-types-1.ll | 16 ++-- .../PropagateAspectsUsage/double.ll | 10 +-- .../PropagateAspectsUsage/multiple-aspects.ll | 12 +-- .../no-uses-of-optional.ll | 4 +- sycl/doc/design/OptionalDeviceFeatures.md | 75 +++++++++---------- 12 files changed, 91 insertions(+), 92 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index db79eb3dae6b1..cdefb491a8d8b 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1073,7 +1073,7 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, AspectsMD.push_back(llvm::ConstantAsMetadata::get( Builder.getInt32(AspectInt.getZExtValue()))); } - Fn->setMetadata("intel_declared_aspects", + Fn->setMetadata("sycl_declared_aspects", llvm::MDNode::get(getLLVMContext(), AspectsMD)); } if (const auto *A = D->getAttr()) { @@ -1083,7 +1083,7 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy, AspectsMD.push_back(llvm::ConstantAsMetadata::get( Builder.getInt32(AspectInt.getZExtValue()))); } - Fn->setMetadata("intel_used_aspects", + Fn->setMetadata("sycl_used_aspects", llvm::MDNode::get(getLLVMContext(), AspectsMD)); } } diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 68bf8070e2320..60b30cd272603 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -915,7 +915,7 @@ void CodeGenModule::Release() { // Emit type name with list of associated device aspects. if (TypesWithAspects.size() > 0) { llvm::NamedMDNode *AspectsMD = - TheModule.getOrInsertNamedMetadata("intel_types_that_use_aspects"); + TheModule.getOrInsertNamedMetadata("sycl_types_that_use_aspects"); for (const auto &Type : TypesWithAspects) { StringRef Name = Type.first; const RecordDecl *RD = Type.second; diff --git a/clang/test/CodeGenSYCL/device_has.cpp b/clang/test/CodeGenSYCL/device_has.cpp index bc3d3ec2effee..ecfd5b5b09226 100644 --- a/clang/test/CodeGenSYCL/device_has.cpp +++ b/clang/test/CodeGenSYCL/device_has.cpp @@ -6,27 +6,27 @@ using namespace sycl; queue q; -// CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !intel_declared_aspects ![[ASPECTS1:[0-9]+]] +// CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] -// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !intel_declared_aspects ![[ASPECTS1]] { +// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] { [[sycl::device_has(sycl::aspect::cpu)]] void func1() {} -// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !intel_declared_aspects ![[ASPECTS2:[0-9]+]] { +// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] { [[sycl::device_has(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {} -// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !intel_declared_aspects ![[EMPTYASPECTS:[0-9]+]] { +// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] { [[sycl::device_has()]] void func3() {} -// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !intel_declared_aspects ![[ASPECTS3:[0-9]+]] { +// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] { template [[sycl::device_has(Aspect)]] void func4() {} -// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !intel_declared_aspects ![[ASPECTS1]] { +// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_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{{.*}} !intel_declared_aspects ![[ASPECTS1]] { +// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] { [[sycl::device_has(getAspect())]] void func6() {} class KernelFunctor { @@ -45,7 +45,7 @@ void foo() { q.submit([&](handler &h) { KernelFunctor f1; h.single_task(f1); - // CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !intel_declared_aspects ![[ASPECTS4:[0-9]+]] + // CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] h.single_task([]() [[sycl::device_has(sycl::aspect::gpu)]] {}); }); } diff --git a/clang/test/CodeGenSYCL/uses_aspects.cpp b/clang/test/CodeGenSYCL/uses_aspects.cpp index feb9b9b746a63..eb3ec5d815d79 100644 --- a/clang/test/CodeGenSYCL/uses_aspects.cpp +++ b/clang/test/CodeGenSYCL/uses_aspects.cpp @@ -10,32 +10,32 @@ class [[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] Type1WithAspect{} class [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp16, sycl::aspect::cpu)]] Type2WithAspect{}; class [[__sycl_detail__::__uses_aspects__(sycl::aspect::host)]] UnusedType3WithAspect{}; -// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !intel_used_aspects ![[ASPECTS1:[0-9]+]] { +// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_used_aspects ![[ASPECTS1:[0-9]+]] { [[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] void func1() {} -// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !intel_used_aspects ![[ASPECTS2:[0-9]+]] { +// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_used_aspects ![[ASPECTS2:[0-9]+]] { [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {} -// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !intel_used_aspects ![[EMPTYASPECTS:[0-9]+]] { +// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_used_aspects ![[EMPTYASPECTS:[0-9]+]] { [[__sycl_detail__::__uses_aspects__()]] void func3() {} -// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !intel_used_aspects ![[ASPECTS3:[0-9]+]] { +// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_used_aspects ![[ASPECTS3:[0-9]+]] { template [[__sycl_detail__::__uses_aspects__(Aspect)]] void func4() {} -// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !intel_used_aspects ![[ASPECTS1]] { +// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_used_aspects ![[ASPECTS1]] { [[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] void func5(); void func5() {} [[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] void func6(); -// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !intel_used_aspects ![[ASPECTS1]] { +// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_used_aspects ![[ASPECTS1]] { void func6() { Type1WithAspect TestObj1; Type2WithAspect TestObj2; } constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; } -// CHECK: define dso_local spir_func void @{{.*}}func7{{.*}} !intel_used_aspects ![[ASPECTS1]] { +// CHECK: define dso_local spir_func void @{{.*}}func7{{.*}} !sycl_used_aspects ![[ASPECTS1]] { [[__sycl_detail__::__uses_aspects__(getAspect())]] void func7() {} class KernelFunctor { @@ -57,7 +57,7 @@ void foo() { h.single_task(f1); }); } -// CHECK: !intel_types_that_use_aspects = !{![[TYPE1:[0-9]+]], ![[TYPE2:[0-9]+]]} +// CHECK: !sycl_types_that_use_aspects = !{![[TYPE1:[0-9]+]], ![[TYPE2:[0-9]+]]} // CHECK: [[TYPE1]] = !{!"class.Type1WithAspect", i32 1} // CHECK: [[TYPE2]] = !{!"class.Type2WithAspect", i32 5, i32 1} // CHECK: [[EMPTYASPECTS]] = !{} diff --git a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp index 629cad79a1b9a..743d65fa84605 100644 --- a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp @@ -49,10 +49,10 @@ namespace { using AspectsSetTy = SmallSet; using TypeToAspectsMapTy = std::unordered_map; -/// Retrieves from metadata (intel_types_that_use_aspects) types +/// Retrieves from metadata (sycl_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"); + const NamedMDNode *Node = M.getNamedMetadata("sycl_types_that_use_aspects"); TypeToAspectsMapTy Result; if (!Node) return Result; @@ -219,7 +219,7 @@ void createUsedAspectsMetadataForFunctions(FunctionToAspectsMapTy &Map) { ConstantInt::getSigned(Type::getInt32Ty(C), A))); MDNode *MDN = MDNode::get(C, AspectsMetadata); - F->setMetadata("intel_used_aspects", MDN); + F->setMetadata("sycl_used_aspects", MDN); } } diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-1.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-1.ll index bd193259b631a..fbfebc9c4c3c6 100644 --- a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-1.ll +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-1.ll @@ -14,14 +14,14 @@ %Optional.A = type { i32 } %Optional.B = type { i32 } -; CHECK: spir_kernel void @kernel1() !intel_used_aspects ![[#ID1:]] { +; CHECK: spir_kernel void @kernel1() !sycl_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:]] { +; CHECK: spir_kernel void @kernel2() !sycl_used_aspects ![[#ID2:]] { define spir_kernel void @kernel2() { call spir_func void @func2() call spir_func void @func3() @@ -34,19 +34,19 @@ define spir_func void @func1() { ret void } -; CHECK: spir_func void @func2() !intel_used_aspects ![[#ID1]] { +; CHECK: spir_func void @func2() !sycl_used_aspects ![[#ID1]] { define spir_func void @func2() { %tmp = alloca %Optional.A ret void } -; CHECK: spir_func void @func3() !intel_used_aspects ![[#ID3:]] { +; CHECK: spir_func void @func3() !sycl_used_aspects ![[#ID3:]] { define spir_func void @func3() { %tmp = alloca %Optional.B ret void } -!intel_types_that_use_aspects = !{!0, !1} +!sycl_types_that_use_aspects = !{!0, !1} !0 = !{!"Optional.A", i32 1} !1 = !{!"Optional.B", i32 2} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-2.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-2.ll index cf236861ad2ee..edbc5422cd5b1 100644 --- a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-2.ll +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-2.ll @@ -15,39 +15,39 @@ %Optional.A = type { i32 } %Optional.B = type { i32 } -; CHECK: spir_kernel void @kernel() !intel_used_aspects ![[#ID1:]] { +; 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() !intel_used_aspects ![[#ID2:]] { +; 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() !intel_used_aspects ![[#ID1]] { +; 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() !intel_used_aspects ![[#ID2]] { +; CHECK: spir_func void @func3() !sycl_used_aspects ![[#ID2]] { define spir_func void @func3() { %tmp = alloca %Optional.A ret void } -; CHECK: spir_func void @func4() !intel_used_aspects ![[#ID3:]] { +; CHECK: spir_func void @func4() !sycl_used_aspects ![[#ID3:]] { define spir_func void @func4() { %tmp = alloca %Optional.B ret void } -!intel_types_that_use_aspects = !{!0, !1} +!sycl_types_that_use_aspects = !{!0, !1} !0 = !{!"Optional.A", i32 1} !1 = !{!"Optional.B", i32 2} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/composite-types-1.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/composite-types-1.ll index d0e92ebf5bf91..97248a63a1637 100644 --- a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/composite-types-1.ll +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/composite-types-1.ll @@ -22,13 +22,13 @@ %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 { +; CHECK: spir_kernel void @kernelD1.uses.optional() !sycl_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 { +; CHECK: spir_func void @funcD1.uses.optional() !sycl_used_aspects !1 { define spir_func void @funcD1.uses.optional() { %tmp = alloca %D1.contains.optional ret void @@ -46,13 +46,13 @@ define spir_func void @funcD2.does.not.use.optional() { ret void } -; CHECK: spir_kernel void @kernelE.uses.optional() !intel_used_aspects !1 { +; CHECK: spir_kernel void @kernelE.uses.optional() !sycl_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 { +; CHECK: spir_func void @funcE.uses.optional() !sycl_used_aspects !1 { define spir_func void @funcE.uses.optional() { %tmp = alloca %E.contains.optional ret void @@ -82,25 +82,25 @@ define spir_func void @funcF2.does.not.use.optional() { ret void } -; CHECK: spir_func %A.optional @funcA.returns.optional() !intel_used_aspects !1 { +; CHECK: spir_func %A.optional @funcA.returns.optional() !sycl_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 { +; CHECK: spir_func void @funcA.uses.array.of.optional() !sycl_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 { +; CHECK: spir_func void @funcA.assepts.optional(%A.optional %0) !sycl_used_aspects !1 { define spir_func void @funcA.assepts.optional(%A.optional %0) { ret void } -!intel_types_that_use_aspects = !{!0} +!sycl_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 index 23bec94282ed7..92134dd63dd67 100644 --- a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/double.ll +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/double.ll @@ -4,31 +4,31 @@ %composite = type { double } -; CHECK: spir_kernel void @kernel() !intel_used_aspects !0 { +; CHECK: spir_kernel void @kernel() !sycl_used_aspects !0 { define spir_kernel void @kernel() { call spir_func void @func() ret void } -; CHECK: spir_func void @func() !intel_used_aspects !0 { +; CHECK: spir_func void @func() !sycl_used_aspects !0 { define spir_func void @func() { %tmp = alloca double ret void } -; CHECK: spir_func void @func.array() !intel_used_aspects !0 { +; CHECK: spir_func void @func.array() !sycl_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 { +; CHECK: spir_func void @func.vector() !sycl_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 { +; CHECK: spir_func void @func.composite() !sycl_used_aspects !0 { define spir_func void @func.composite() { %tmp = alloca %composite ret void diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/multiple-aspects.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/multiple-aspects.ll index d7622f911b486..2bba0e1040020 100644 --- a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/multiple-aspects.ll +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/multiple-aspects.ll @@ -7,40 +7,40 @@ %C = type { i32 } %D = type { i32 } -; CHECK: define spir_func void @funcA() !intel_used_aspects ![[#ID0:]] { +; CHECK: define spir_func void @funcA() !sycl_used_aspects ![[#ID0:]] { define spir_func void @funcA() { %tmp = alloca %A ret void } -; CHECK: define spir_func void @funcB() !intel_used_aspects ![[#ID1:]] { +; CHECK: define spir_func void @funcB() !sycl_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:]] { +; CHECK: define spir_func void @funcC() !sycl_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:]] { +; CHECK: define spir_func void @funcD() !sycl_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]] { +; CHECK: define spir_kernel void @kernel() !sycl_used_aspects ![[#ID3]] { define spir_kernel void @kernel() { call spir_func void @funcD() ret void } -!intel_types_that_use_aspects = !{!0, !1, !2, !3} +!sycl_types_that_use_aspects = !{!0, !1, !2, !3} !0 = !{!"A", i32 0} !1 = !{!"B", i32 1} !2 = !{!"C", i32 2} diff --git a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/no-uses-of-optional.ll b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/no-uses-of-optional.ll index 967accbb3e3bc..1883423c08813 100644 --- a/llvm/test/SYCLLowerIR/PropagateAspectsUsage/no-uses-of-optional.ll +++ b/llvm/test/SYCLLowerIR/PropagateAspectsUsage/no-uses-of-optional.ll @@ -1,4 +1,4 @@ -; RUN: opt -passes=sycl-propagate-aspects-usage < %s -S | FileCheck %s --implicit-check-not "!intel_used_aspects" +; RUN: opt -passes=sycl-propagate-aspects-usage < %s -S | FileCheck %s --implicit-check-not "!sycl_used_aspects" ; ; Test checks that no metadata propagates because MyStruct ; isn't used inside functions. @@ -16,5 +16,5 @@ define weak dso_local spir_func void @func() { ret void } -!intel_types_that_use_aspects = !{!0} +!sycl_types_that_use_aspects = !{!0} !0 = !{!"MyStruct", i32 1} diff --git a/sycl/doc/design/OptionalDeviceFeatures.md b/sycl/doc/design/OptionalDeviceFeatures.md index 02556c5c5a396..0cb3ec4afd9a2 100644 --- a/sycl/doc/design/OptionalDeviceFeatures.md +++ b/sycl/doc/design/OptionalDeviceFeatures.md @@ -396,19 +396,19 @@ In order to communicate the information from `[[sycl::device_has()]]` and `[[sycl_detail::uses_aspects()]]` attributes to the DPC++ post-link tool, we introduce several new LLVM IR metadata. -The named metadata `!intel_types_that_use_aspects` conveys information about +The named metadata `!sycl_types_that_use_aspects` conveys information about types that are decorated with `[[sycl_detail::uses_aspects()]]`. This metadata is not referenced by any instruction in the module, so it must be looked up by name. The format looks like this: ``` -!intel_types_that_use_aspects = !{!0, !1, !2} +!sycl_types_that_use_aspects = !{!0, !1, !2} !0 = !{!"class.sycl::detail::half_impl::half", i32 8} !1 = !{!"class.sycl::amx_type", i32 9} !2 = !{!"class.sycl::other_type", i32 8, i32 9} ``` -The value of the `!intel_types_that_use_aspects` metadata is a list of unnamed +The value of the `!sycl_types_that_use_aspects` metadata is a list of unnamed metadata nodes, each of which describes one type that is decorated with `[[sycl_detail::uses_aspects()]]`. The value of each unnamed metadata node starts with a string giving the name of the type which is followed by a list of @@ -423,21 +423,20 @@ allow metadata to be attached directly to types. This representation works around that limitation by creating global named metadata that references the type's name. -We also introduce three metadata that can be attached to a function definition -similar to the existing `!intel_reqd_sub_group_size`: +We also introduce three metadata that can be attached to a function definition: -* The `!intel_declared_aspects` metadata is used for functions that are +* The `!sycl_declared_aspects` metadata is used for functions that are decorated with `[[sycl::device_has()]]`. The value of the metadata node is a list of `i32` constants, where each constant is a value from `enum class aspect` representing the aspects listed in the attribute. -* The `!intel_used_aspects` metadata is used to store the propagated +* The `!sycl_used_aspects` metadata is used to store the propagated information about all aspects used by a kernel or exported device function. The value of this metadata node is also a list of `i32` constants, where each constant is a value from `enum class aspect` representing the aspects that are used by the kernel or exported device function. -* The `!intel_fixed_targets` metadata is used to decorate kernel functions and +* The `!sycl_fixed_targets` metadata is used to decorate kernel functions and `SYCL_EXTERNAL` functions, telling the value of the `-fsycl-fixed-targets` switch that was used to compile the translation unit. The value of this metadata node is a list of string literals corresponding to the list of @@ -458,7 +457,7 @@ numerical values `8` and `9`. In addition, the function uses an optional feature that corresponds to an aspect with numerical value `8`. ``` -define void @foo() !intel_declared_aspects !1 !intel_used_aspects !2 {} +define void @foo() !sycl_declared_aspects !1 !sycl_used_aspects !2 {} !1 = !{i32 8, i32 9} !2 = !{i32 8} ``` @@ -473,15 +472,15 @@ to the following rules: * If the translation unit contains any type definitions that are decorated with `[[sycl_detail::uses_aspects()]]`, the front-end creates an - `!intel_types_that_use_aspects` metadata describing the aspects used by all + `!sycl_types_that_use_aspects` metadata describing the aspects used by all such types. * If a function is decorated with `[[sycl_detail::uses_aspects()]]`, the - front-end adds an `!intel_used_aspects` metadata to the function's definition + front-end adds an `!sycl_used_aspects` metadata to the function's definition listing the aspects from that attribute. * If a function is decorated with `[[sycl::device_has()]]`, the front-end adds - an `!intel_declared_aspects` metadata to the function's definition listing + an `!sycl_declared_aspects` metadata to the function's definition listing the aspects from that attribute. @@ -489,18 +488,18 @@ to the following rules: We add a new IR phase to the device compiler which does the following: -* Creates (or augments) a function's `!intel_used_aspects` metadata with +* Creates (or augments) a function's `!sycl_used_aspects` metadata with aspects that come from references to types in the - `intel_types_that_use_aspects` list. + `sycl_types_that_use_aspects` list. -* Propagates each function's `!intel_used_aspects` metadata up the static call +* Propagates each function's `!sycl_used_aspects` metadata up the static call graph so that each function lists the aspects used by that function and by any functions it calls. -* Diagnoses a warning if any function that has `!intel_declared_aspects` uses +* Diagnoses a warning if any function that has `!sycl_declared_aspects` uses an aspect not listed in that declared set. -* Creates an `!intel_fixed_targets` metadata for each kernel function or +* Creates an `!sycl_fixed_targets` metadata for each kernel function or `SYCL_EXTERNAL` function that is defined. This is done regardless of whether the `-fsycl-fixed-targets` command line switch is specified. If the switch is not specified, the metadata has an empty list of targets. @@ -517,40 +516,40 @@ other IR phases. Implementing the first bullet point is straightforward. The implementation can scan the IR for each function looking for instructions that reference a type. -It can then see if that type is in the `!intel_types_that_use_aspects` set; if -so it adds the type's aspects to the function's `!intel_used_aspects` set. +It can then see if that type is in the `!sycl_types_that_use_aspects` set; if +so it adds the type's aspects to the function's `!sycl_used_aspects` set. While doing this, the implementation must have a special case for the `double` type because the front-end does not include that type in the -`!intel_types_that_use_aspects` set. If a function references the `double` +`!sycl_types_that_use_aspects` set. If a function references the `double` type, the implementation implicitly assumes that the function uses -`aspect::fp64` and adds that aspect to the function's `!intel_used_aspects` +`aspect::fp64` and adds that aspect to the function's `!sycl_used_aspects` set. **NOTE**: This scan of the IR will require comparing the type referenced by each IR instruction with the names of the types in the -`!intel_types_that_use_aspects` metadata. It would be very inefficient if we +`!sycl_types_that_use_aspects` metadata. It would be very inefficient if we did a string comparison each time. As an optimization, the implementation can -first lookup up each type name in the `!intel_types_that_use_aspects` metadata +first lookup up each type name in the `!sycl_types_that_use_aspects` metadata set, finding the "type pointer" that corresponds to each type name. Then the pass over the IR can compare the type pointer in each IR instruction with the -type pointers from the `!intel_types_that_use_aspects` metadata set. +type pointers from the `!sycl_types_that_use_aspects` metadata set. The second bullet point requires building the static call graph, but the implementation need not scan the instructions in each function. Instead, it -need only look at the `!intel_used_aspects` metadata for each function, +need only look at the `!sycl_used_aspects` metadata for each function, propagating the aspects used by each function up to it callers and augmenting -the caller's `!intel_used_aspects` set. +the caller's `!sycl_used_aspects` set. Diagnosing warnings for the third bullet point is then straightforward. The -implementation looks for functions that have `!intel_declared_aspects` and -compares that set with the `!intel_used_aspects` set (if any). If a function +implementation looks for functions that have `!sycl_declared_aspects` and +compares that set with the `!sycl_used_aspects` set (if any). If a function uses an aspect that is not in the declared set, the implementation issues a warning. Diagnosing warnings for the fifth bullet point requires the [device configuration file][7] which gives the set of allowed optional features for each target device. The implementation looks for functions that have either -`!intel_declared_aspects` or `!intel_used_aspects`, and it compares the aspects +`!sycl_declared_aspects` or `!sycl_used_aspects`, and it compares the aspects from these metadata to the allowed list in the configuration file. If any aspect is not on the allowed list, the implementation issues a warning. In addition, the implementation looks for device functions that have @@ -608,8 +607,8 @@ The downside, though, is that the warning message is less informative. ### Assumptions on other phases of clang -The post-link tool (described below) uses the `!intel_used_aspects` and -`!intel_declared_aspects` metadata, so this metadata must be retained by any +The post-link tool (described below) uses the `!sycl_used_aspects` and +`!sycl_declared_aspects` metadata, so this metadata must be retained by any other clang passes. However, post-link only uses this metadata when it decorates the definition of a kernel function or the definition of an exported device function, so it does not matter if intervening clang passes discard the @@ -622,7 +621,7 @@ always have external linkage, so there is no possibility that a clang phase will optimize them away. **NOTE**: Ideally, we would change the llvm-link tool to somehow preserve the -`!intel_declared_aspects` and `!intel_used_aspects` metadata for functions +`!sycl_declared_aspects` and `!sycl_used_aspects` metadata for functions marked `SYCL_EXTERNAL` so that we could compare the declared aspects (in the module that imports the function) with the used aspects (in the module the exports the function). This would allow us to diagnose errors where the @@ -663,7 +662,7 @@ of aspects. For the purposes of this analysis, the set of *UsedAspects* aspects is computed by taking the union of the aspects listed in the kernel's (or device -function's) `!intel_used_aspects` and `!intel_declared_aspects` sets. This is +function's) `!sycl_used_aspects` and `!sycl_declared_aspects` sets. This is consistent with the SYCL specification, which says that a kernel decorated with `[[sycl::device_has()]]` may only be submitted to a device that provides the listed aspects, regardless of whether the kernel actually uses those aspects. @@ -689,11 +688,11 @@ device image so long as the translation units were compiled with the same Therefore, two kernels or exported device functions can only be bundled together into the same device image if: -* They both have `!intel_fixed_targets` metadata with the same non-empty set of +* They both have `!sycl_fixed_targets` metadata with the same non-empty set of targets, or * All of the following are true: - - Both have an empty set of `!intel_fixed_targets` metadata, + - Both have an empty set of `!sycl_fixed_targets` metadata, - They share the same set of *UsedAspects* aspects, - They either both have no required work-group size or both have the same required work-group size, and @@ -740,8 +739,8 @@ If the image contains kernels that were *not* compiled with * Set *FinalWorkGroup* to the image's required work-group size (which could be the empty set if the image has no required work-group size). * Scan over all functions in the image and examine the function's metadata: - - If the function has either `!intel_used_aspects` or - `!intel_declared_aspects` metadata and one of the aspects in that metadata + - If the function has either `!sycl_used_aspects` or + `!sycl_declared_aspects` metadata and one of the aspects in that metadata is not in the image's *UsedAspects* set, issue a warning and add that aspect to the *FinalUsedAspects* set. - If the function has `!intel_reqd_sub_group_size` metadata and the size is @@ -775,7 +774,7 @@ the pass works as follows: * Set *FinalFixedTargets* to the image's set of fixed target devices. * Scan over all functions in the image looking for functions that have the - `!intel_fixed_targets` metadata. If the metadata exists and its set includes + `!sycl_fixed_targets` metadata. If the metadata exists and its set includes any target devices not in the image's set of fixed targets, issue a warning and set *FinalFixedTargets* to the intersection of the metadata's target set and the *FinalFixedTargets* set. (This may result in *FinalFixedTargets*