[SYCL] Rename optional device feature metadata

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<SYCLUsesAspectsAttr>()) {
@@ -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));
     }
   }

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;

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::aspect Aspect>
 [[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<class kernel_name_1>(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<class kernel_name_2>([]() [[sycl::device_has(sycl::aspect::gpu)]] {});
   });
 }

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::aspect Aspect>
 [[__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<class kernel_name_1>(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]] = !{}

llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp

@@ -49,10 +49,10 @@ namespace {
 using AspectsSetTy = SmallSet<int, 4>;
 using TypeToAspectsMapTy = std::unordered_map<const Type *, AspectsSetTy>;
 
-/// 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);
   }
 }
 

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}
 

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}
 

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}

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

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}

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}