[SYCL] CFE fixes for sycl_ext_oneapi_kernel_properties (#7016)

This commit makes the following changes to the Clang frontend in
preparation for the initial implementation of the
sycl_ext_oneapi_kernel_properties extension:
* Fixed an issue preventing initializer lists of characters to be used
as strings in add_ir_{attributes|annotations}_* attributes.
* The add_ir_attributes_function will now be copied onto the generated
kernel declaration from the recognized kernel_* functions.
* Added a warning diagnostic for when a kernel has both an
add_ir_attributes_function (with values) and one or more potentially
conflicting SYCL attributes.

This is split from #6941.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

clang/include/clang/Basic/Attr.td

@@ -1603,6 +1603,20 @@ def SYCLAddIRAttrCommonMembers : SYCLAddIRAttrMemberCodeHolder<[{
     if (ValueQType->isIntegralOrEnumerationType() ||
         ValueQType->isFloatingType())
       return Value.getAsString(Context, ValueQType);
+    if (ValueQType->isArrayType() &&
+        (ValueQType->getArrayElementTypeNoTypeQual()->isCharType() ||
+         ValueQType->getArrayElementTypeNoTypeQual()
+                   ->isIntegralOrEnumerationType())) {
+      SmallString<10> StrBuffer;
+      for (unsigned I = 0; I < Value.getArraySize(); ++I) {
+        const APValue &ArrayElem = Value.getArrayInitializedElt(I);
+        char C = static_cast<char>(ArrayElem.getInt().getExtValue());
+        if (C == 0)
+          break;
+        StrBuffer += C;
+      }
+      return std::string(StrBuffer);
+    }
     return None;
   }
 
@@ -1628,6 +1642,33 @@ def SYCLAddIRAttrCommonMembers : SYCLAddIRAttrMemberCodeHolder<[{
                                       ValueE->getType()->isSignedIntegerType());
       return std::string(IntegerStrBuffer);
     }
+    if (const auto *InitListE = dyn_cast<InitListExpr>(ValueE)) {
+      if (InitListE->isStringLiteralInit()) {
+        const Expr *StringInitE = InitListE->getInit(0)->IgnoreParenImpCasts();
+        return getValidAttributeValueAsString(StringInitE, Context);
+      }
+
+      SmallString<10> StrBuffer;
+      for (const auto *InitE : InitListE->inits()) {
+        const Expr *InitNoImpCastE = InitE->IgnoreParenImpCasts();
+        char C = 0;
+        if (const auto *CharacterVal =
+              dyn_cast<CharacterLiteral>(InitNoImpCastE))
+          C = static_cast<char>(CharacterVal->getValue());
+        else if (const auto *IntegerVal =
+              dyn_cast<IntegerLiteral>(InitNoImpCastE))
+          C = static_cast<char>(IntegerVal->getValue().getSExtValue());
+        else
+          return None;
+
+        // Null terminator will end the string reading.
+        if (C == 0)
+          break;
+
+        StrBuffer += C;
+      }
+      return std::string(StrBuffer);
+    }
 
     const auto *ValueCE = dyn_cast<ConstantExpr>(ValueE);
     if (!ValueCE)

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11887,6 +11887,9 @@ def err_sycl_add_ir_attribute_invalid_value : Error<
 def err_sycl_add_ir_attribute_invalid_filter : Error<
   "initializer list in the first argument of %0 must contain only string "
   "literals">;
+def warn_sycl_old_and_new_kernel_attributes : Warning<
+  "kernel has both attribute %0 and kernel properties; conflicting properties "
+  "are ignored">, InGroup<IgnoredAttributes>;
 
 // errors of expect.with.probability
 def err_probability_not_constant_float : Error<

clang/include/clang/Sema/Sema.h

@@ -10984,6 +10984,7 @@ class Sema final {
   SYCLIntelMaxWorkGroupSizeAttr *
   MergeSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
                                      const SYCLIntelMaxWorkGroupSizeAttr &A);
+  void CheckSYCLAddIRAttributesFunctionAttrConflicts(Decl *D);
   SYCLAddIRAttributesFunctionAttr *MergeSYCLAddIRAttributesFunctionAttr(
       Decl *D, const SYCLAddIRAttributesFunctionAttr &A);
   void AddSYCLAddIRAttributesFunctionAttr(Decl *D,

clang/lib/Sema/SemaDeclAttr.cpp

@@ -7808,6 +7808,22 @@ static bool checkSYCLAddIRAttributesMergeability(const AddIRAttrT &NewAttr,
   return false;
 }
 
+void Sema::CheckSYCLAddIRAttributesFunctionAttrConflicts(Decl *D) {
+  const auto *AddIRFuncAttr = D->getAttr<SYCLAddIRAttributesFunctionAttr>();
+  if (!AddIRFuncAttr || AddIRFuncAttr->args_size() == 0 ||
+      hasDependentExpr(AddIRFuncAttr->args_begin(), AddIRFuncAttr->args_size()))
+    return;
+
+  // If there are potentially conflicting attributes, we issue a warning.
+  for (const auto *Attr : std::vector<AttributeCommonInfo *>{
+           D->getAttr<ReqdWorkGroupSizeAttr>(),
+           D->getAttr<IntelReqdSubGroupSizeAttr>(),
+           D->getAttr<WorkGroupSizeHintAttr>()})
+    if (Attr)
+      Diag(Attr->getLoc(), diag::warn_sycl_old_and_new_kernel_attributes)
+          << Attr;
+}
+
 SYCLAddIRAttributesFunctionAttr *Sema::MergeSYCLAddIRAttributesFunctionAttr(
     Decl *D, const SYCLAddIRAttributesFunctionAttr &A) {
   if (const auto *ExistingAttr =

clang/lib/Sema/SemaSYCL.cpp

@@ -2249,6 +2249,10 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
             createKernelDecl(S.getASTContext(), Loc, IsInline, IsSIMDKernel)),
         FuncContext(SemaRef, KernelDecl) {
     S.addSyclOpenCLKernel(SYCLKernel, KernelDecl);
+
+    if (const auto *AddIRAttrFunc =
+            SYCLKernel->getAttr<SYCLAddIRAttributesFunctionAttr>())
+      KernelDecl->addAttr(AddIRAttrFunc->clone(SemaRef.getASTContext()));
   }
 
   ~SyclKernelDeclCreator() {
@@ -4301,6 +4305,7 @@ void Sema::MarkDevices() {
     for (auto *A : T.GetCollectedAttributes())
       PropagateAndDiagnoseDeviceAttr(*this, T, A, T.GetSYCLKernel(),
                                      T.GetKernelBody());
+    CheckSYCLAddIRAttributesFunctionAttrConflicts(T.GetSYCLKernel());
   }
 }
 

clang/test/CodeGenSYCL/Inputs/mock_properties.hpp

@@ -55,6 +55,8 @@ const char PropertyName6[] = "Prop6";
 constexpr decltype(nullptr) PropertyValue6 = nullptr;
 const char PropertyName7[] = "Prop7";
 constexpr ScopedTestEnum PropertyValue7 = ScopedTestEnum::ScopedEnum1;
+const char PropertyName8[] = "Prop8";
+constexpr char PropertyValue8[] = {'P', 114, 'o', 'p', 0x65, 'r', 't', 'y', 0};
 
 using prop1 = StringProperty<PropertyName1, PropertyValue1>;
 using prop2 = IntProperty<PropertyName2, PropertyValue2>;
@@ -63,3 +65,4 @@ using prop4 = TestEnumProperty<PropertyName4, PropertyValue4>;
 using prop5 = StringProperty<PropertyName5, PropertyValue5>;
 using prop6 = NullptrProperty<PropertyName6, PropertyValue6>;
 using prop7 = ScopedTestEnumProperty<PropertyName7, PropertyValue7>;
+using prop8 = StringProperty<PropertyName8, PropertyValue8>;

clang/test/CodeGenSYCL/add_ir_annotations_member.cpp

@@ -34,8 +34,8 @@ class h {
   TEST_T x
 #ifdef __SYCL_DEVICE_ONLY__
       [[__sycl_detail__::add_ir_annotations_member(
-          "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17",
-          "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]]
+          "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18",
+          "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]]
 #endif
       ;
 
@@ -48,8 +48,8 @@ template <typename... Properties> class gh {
   TEST_T x
 #ifdef __SYCL_DEVICE_ONLY__
       [[__sycl_detail__::add_ir_annotations_member(
-          Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17",
-          Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]]
+          Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18",
+          Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]]
 #endif
       ;
 
@@ -62,8 +62,8 @@ template <typename... Properties> class hg {
   TEST_T x
 #ifdef __SYCL_DEVICE_ONLY__
       [[__sycl_detail__::add_ir_annotations_member(
-          "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", Properties::name...,
-          "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, Properties::value...)]]
+          "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", Properties::name...,
+          "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8, Properties::value...)]]
 #endif
       ;
 
@@ -73,7 +73,7 @@ template <typename... Properties> class hg {
 
 int main() {
   sycl::queue q;
-  g<prop1, prop2, prop3, prop4, prop5, prop6, prop7> a;
+  g<prop1, prop2, prop3, prop4, prop5, prop6, prop7, prop8> a;
   q.submit([&](sycl::handler &h) {
     h.single_task<class test_kernel1>(
         [=]() {
@@ -87,14 +87,14 @@ int main() {
           (void)b.x;
         });
   });
-  gh<prop1, prop2, prop3, prop4, prop5, prop6, prop7> c;
+  gh<prop1, prop2, prop3, prop4, prop5, prop6, prop7, prop8> c;
   q.submit([&](sycl::handler &h) {
     h.single_task<class test_kernel3>(
         [=]() {
           (void)c.x;
         });
   });
-  hg<prop1, prop2, prop3, prop4, prop5, prop6, prop7> d;
+  hg<prop1, prop2, prop3, prop4, prop5, prop6, prop7, prop8> d;
   q.submit([&](sycl::handler &h) {
     h.single_task<class test_kernel4>(
         [=]() {
@@ -112,25 +112,28 @@ int main() {
 // CHECK-DAG: @[[Prop5Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop5\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop6Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop6\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop7Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop7\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop8Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop8\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop11Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop11\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop12Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop12\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop13Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop13\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop14Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop14\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop15Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop15\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop16Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop16\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop17Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop17\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop18Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop18\00", section "llvm.metadata"
 
 // CHECK-DAG: @[[Prop1Value:.*]] = private unnamed_addr constant [16 x i8] c"Property string\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop2_7_14Value:.*]] = private unnamed_addr constant [2 x i8] c"1\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop3Value:.*]] = private unnamed_addr constant [5 x i8] c"true\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop4_12_17Value:.*]] = private unnamed_addr constant [2 x i8] c"2\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop8_18Value:.*]] = private unnamed_addr constant [9 x i8] c"Property\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop11Value:.*]] = private unnamed_addr constant [24 x i8] c"Another property string\00", section "llvm.metadata"
 // CHECK-DAG: @[[Prop13Value:.*]] = private unnamed_addr constant [6 x i8] c"false\00", section "llvm.metadata"
 
-// CHECK-DAG: @[[GArgs:.*]] = private unnamed_addr constant { ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr } { ptr @[[Prop1Name]], ptr @[[Prop1Value]], ptr @[[Prop2Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop3Name]], ptr @[[Prop3Value]], ptr @[[Prop4Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop5Name]], ptr null, ptr @[[Prop6Name]], ptr null, ptr @[[Prop7Name]], ptr @[[Prop2_7_14Value]] }, section "llvm.metadata"
-// CHECK-DAG: @[[HArgs:.*]] = private unnamed_addr constant { ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr } { ptr @[[Prop11Name]], ptr @[[Prop11Value]], ptr @[[Prop12Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop13Name]], ptr @[[Prop13Value]], ptr @[[Prop14Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop15Name]], ptr null, ptr @[[Prop16Name]], ptr null, ptr @[[Prop17Name]], ptr @[[Prop4_12_17Value]] }, section "llvm.metadata"
-// CHECK-DAG: @[[GHArgs:.*]] = private unnamed_addr constant { ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr } { ptr @[[Prop1Name]], ptr @[[Prop1Value]], ptr @[[Prop2Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop3Name]], ptr @[[Prop3Value]], ptr @[[Prop4Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop5Name]], ptr null, ptr @[[Prop6Name]], ptr null, ptr @[[Prop7Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop11Name]], ptr @[[Prop11Value]], ptr @[[Prop12Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop13Name]], ptr @[[Prop13Value]], ptr @[[Prop14Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop15Name]], ptr null, ptr @[[Prop16Name]], ptr null, ptr @[[Prop17Name]], ptr @[[Prop4_12_17Value]] }, section "llvm.metadata"
-// CHECK-DAG: @[[HGArgs:.*]] = private unnamed_addr constant { ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr } { ptr @[[Prop11Name]], ptr @[[Prop11Value]], ptr @[[Prop12Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop13Name]], ptr @[[Prop13Value]], ptr @[[Prop14Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop15Name]], ptr null, ptr @[[Prop16Name]], ptr null, ptr @[[Prop17Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop1Name]], ptr @[[Prop1Value]], ptr @[[Prop2Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop3Name]], ptr @[[Prop3Value]], ptr @[[Prop4Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop5Name]], ptr null, ptr @[[Prop6Name]], ptr null, ptr @[[Prop7Name]], ptr @[[Prop2_7_14Value]] }, section "llvm.metadata"
+// CHECK-DAG: @[[GArgs:.*]] = private unnamed_addr constant { ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr } { ptr @[[Prop1Name]], ptr @[[Prop1Value]], ptr @[[Prop2Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop3Name]], ptr @[[Prop3Value]], ptr @[[Prop4Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop5Name]], ptr null, ptr @[[Prop6Name]], ptr null, ptr @[[Prop7Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop8Name]], ptr @[[Prop8_18Value]] }, section "llvm.metadata"
+// CHECK-DAG: @[[HArgs:.*]] = private unnamed_addr constant { ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr } { ptr @[[Prop11Name]], ptr @[[Prop11Value]], ptr @[[Prop12Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop13Name]], ptr @[[Prop13Value]], ptr @[[Prop14Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop15Name]], ptr null, ptr @[[Prop16Name]], ptr null, ptr @[[Prop17Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop18Name]], ptr @[[Prop8_18Value]] }, section "llvm.metadata"
+// CHECK-DAG: @[[GHArgs:.*]] = private unnamed_addr constant { ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr } { ptr @[[Prop1Name]], ptr @[[Prop1Value]], ptr @[[Prop2Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop3Name]], ptr @[[Prop3Value]], ptr @[[Prop4Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop5Name]], ptr null, ptr @[[Prop6Name]], ptr null, ptr @[[Prop7Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop8Name]], ptr @[[Prop8_18Value]], ptr @[[Prop11Name]], ptr @[[Prop11Value]], ptr @[[Prop12Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop13Name]], ptr @[[Prop13Value]], ptr @[[Prop14Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop15Name]], ptr null, ptr @[[Prop16Name]], ptr null, ptr @[[Prop17Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop18Name]], ptr @[[Prop8_18Value]] }, section "llvm.metadata"
+// CHECK-DAG: @[[HGArgs:.*]] = private unnamed_addr constant { ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr } { ptr @[[Prop11Name]], ptr @[[Prop11Value]], ptr @[[Prop12Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop13Name]], ptr @[[Prop13Value]], ptr @[[Prop14Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop15Name]], ptr null, ptr @[[Prop16Name]], ptr null, ptr @[[Prop17Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop18Name]], ptr @[[Prop8_18Value]], ptr @[[Prop1Name]], ptr @[[Prop1Value]], ptr @[[Prop2Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop3Name]], ptr @[[Prop3Value]], ptr @[[Prop4Name]], ptr @[[Prop4_12_17Value]], ptr @[[Prop5Name]], ptr null, ptr @[[Prop6Name]], ptr null, ptr @[[Prop7Name]], ptr @[[Prop2_7_14Value]], ptr @[[Prop8Name]], ptr @[[Prop8_18Value]] }, section "llvm.metadata"
 
 // CHECK-DAG: %{{.*}} = call ptr addrspace(4) @llvm.ptr.annotation.p4(ptr {{.*}}, ptr @[[AnnotName]], {{.*}}, i32 {{.*}}, ptr @[[GArgs]])
 // CHECK-DAG: %{{.*}} = call ptr addrspace(4) @llvm.ptr.annotation.p4(ptr {{.*}}, ptr @[[AnnotName]], {{.*}}, i32 {{.*}}, ptr @[[HArgs]])