diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 92ca5135ab4df..25be7a31b7fd3 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/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(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(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(InitNoImpCastE)) + C = static_cast(CharacterVal->getValue()); + else if (const auto *IntegerVal = + dyn_cast(InitNoImpCastE)) + C = static_cast(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(ValueE); if (!ValueCE) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 6459b78f03c2b..e15c2c4a349d5 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/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; // errors of expect.with.probability def err_probability_not_constant_float : Error< diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 1e36e5f668554..a7673ca2373d8 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/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, diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index cc7db7fbb450d..c5873abd1e892 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/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(); + 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{ + D->getAttr(), + D->getAttr(), + D->getAttr()}) + 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 = diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 73a61582f20f5..8777c6161987e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/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()) + 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()); } } diff --git a/clang/test/CodeGenSYCL/Inputs/mock_properties.hpp b/clang/test/CodeGenSYCL/Inputs/mock_properties.hpp index 5b3607ab16bae..b22c25c479ab8 100644 --- a/clang/test/CodeGenSYCL/Inputs/mock_properties.hpp +++ b/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; using prop2 = IntProperty; @@ -63,3 +65,4 @@ using prop4 = TestEnumProperty; using prop5 = StringProperty; using prop6 = NullptrProperty; using prop7 = ScopedTestEnumProperty; +using prop8 = StringProperty; diff --git a/clang/test/CodeGenSYCL/add_ir_annotations_member.cpp b/clang/test/CodeGenSYCL/add_ir_annotations_member.cpp index 9c1eccea0bf5b..f024e85d1d9ee 100644 --- a/clang/test/CodeGenSYCL/add_ir_annotations_member.cpp +++ b/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 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 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 class hg { int main() { sycl::queue q; - g a; + g a; q.submit([&](sycl::handler &h) { h.single_task( [=]() { @@ -87,14 +87,14 @@ int main() { (void)b.x; }); }); - gh c; + gh c; q.submit([&](sycl::handler &h) { h.single_task( [=]() { (void)c.x; }); }); - hg d; + hg d; q.submit([&](sycl::handler &h) { h.single_task( [=]() { @@ -112,6 +112,7 @@ 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" @@ -119,18 +120,20 @@ int main() { // 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]]) diff --git a/clang/test/CodeGenSYCL/add_ir_attributes_function.cpp b/clang/test/CodeGenSYCL/add_ir_attributes_function.cpp index a2c94e1eb3b0b..47c5c6e133444 100644 --- a/clang/test/CodeGenSYCL/add_ir_attributes_function.cpp +++ b/clang/test/CodeGenSYCL/add_ir_attributes_function.cpp @@ -19,8 +19,8 @@ free_func1() { #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( - "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 void free_func2() { @@ -29,8 +29,8 @@ free_func2() { template #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( - 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 void free_func3() { @@ -39,8 +39,8 @@ free_func3() { template #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( - "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 void free_func4() { @@ -49,8 +49,8 @@ free_func4() { template #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( - "", "", "", "", "", "", "", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] + "", "", "", "", "", "", "", "", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] #endif void free_func5() { @@ -59,8 +59,8 @@ free_func5() { template #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( - "", "Prop12", "", "", "", "Prop16", "Prop17", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] + "", "Prop12", "", "", "", "Prop16", "Prop17", "Prop18", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] #endif void free_func6() { @@ -83,8 +83,8 @@ class KernelFunctor2 { public: #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( - "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 void operator()() const { @@ -97,8 +97,8 @@ class KernelFunctor3 { public: #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( - 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 void operator()() const { @@ -111,8 +111,8 @@ class KernelFunctor4 { public: #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( - "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 void operator()() const { @@ -124,8 +124,8 @@ class KernelFunctor5 { public: #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( - "", "", "", "", "", "", "", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] + "", "", "", "", "", "", "", "", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] #endif void operator()() const { @@ -137,8 +137,8 @@ class KernelFunctor6 { public: #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_function( - "", "Prop12", "", "", "", "Prop16", "Prop17", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] + "", "Prop12", "", "", "", "Prop16", "Prop17", "Prop18", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] #endif void operator()() const { @@ -149,7 +149,7 @@ class KernelFunctor6 { int main() { sycl::queue q; q.submit([&](sycl::handler &h) { - KernelFunctor1 f{}; + KernelFunctor1 f{}; h.single_task(f); }); q.submit([&](sycl::handler &h) { @@ -157,11 +157,11 @@ int main() { h.single_task(f); }); q.submit([&](sycl::handler &h) { - KernelFunctor3 f{}; + KernelFunctor3 f{}; h.single_task(f); }); q.submit([&](sycl::handler &h) { - KernelFunctor4 f{}; + KernelFunctor4 f{}; h.single_task(f); }); q.submit([&](sycl::handler &h) { @@ -186,14 +186,14 @@ int main() { // CHECK-DAG: define {{.*}}spir_func void @{{.*}}free_func4{{.*}}() #[[Func3and4Attrs]] // CHECK-DAG: define {{.*}}spir_func void @{{.*}}free_func5{{.*}}() {{.*}} // CHECK-DAG: define {{.*}}spir_func void @{{.*}}free_func6{{.*}}() #[[Func6Attrs:[0-9]+]] -// CHECK-DAG: attributes #[[KernFunc1Attrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} } -// CHECK-DAG: attributes #[[KernFunc2Attrs]] = { {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} } -// CHECK-DAG: attributes #[[KernFunc3And4Attrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} } -// CHECK-DAG: attributes #[[KernFunc6Attrs]] = { {{.*}}"Prop12"="2"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} } -// CHECK-DAG: attributes #[[Func1Attrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} } -// CHECK-DAG: attributes #[[Func2Attrs]] = { {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} } -// CHECK-DAG: attributes #[[Func3and4Attrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} } -// CHECK-DAG: attributes #[[Func6Attrs]] = { {{.*}}"Prop12"="2"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} } +// CHECK-DAG: attributes #[[KernFunc1Attrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} } +// CHECK-DAG: attributes #[[KernFunc2Attrs]] = { {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} } +// CHECK-DAG: attributes #[[KernFunc3And4Attrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} } +// CHECK-DAG: attributes #[[KernFunc6Attrs]] = { {{.*}}"Prop12"="2"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} } +// CHECK-DAG: attributes #[[Func1Attrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} } +// CHECK-DAG: attributes #[[Func2Attrs]] = { {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} } +// CHECK-DAG: attributes #[[Func3and4Attrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} } +// CHECK-DAG: attributes #[[Func6Attrs]] = { {{.*}}"Prop12"="2"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} } // CHECK-NOT: ""="Another property string" // CHECK-NOT: ""="1" // CHECK-NOT: ""="2" diff --git a/clang/test/CodeGenSYCL/add_ir_attributes_global_variable.cpp b/clang/test/CodeGenSYCL/add_ir_attributes_global_variable.cpp index 38a9e579a8138..9361e5a92ae56 100644 --- a/clang/test/CodeGenSYCL/add_ir_attributes_global_variable.cpp +++ b/clang/test/CodeGenSYCL/add_ir_attributes_global_variable.cpp @@ -23,8 +23,8 @@ template struct struct #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_global_variable( - "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 h { int x; @@ -36,8 +36,8 @@ struct template struct #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_global_variable( - NameValues::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", - NameValues::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] + NameValues::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", + NameValues::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] #endif gh { int x; @@ -49,8 +49,8 @@ template struct template struct #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_global_variable( - "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", NameValues::name..., - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, NameValues::value...)]] + "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", NameValues::name..., + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8, NameValues::value...)]] #endif hg { int x; @@ -62,8 +62,8 @@ template struct struct #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_global_variable( - "", "", "", "", "", "", "", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] + "", "", "", "", "", "", "", "", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] #endif np { int x; @@ -75,8 +75,8 @@ struct struct #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_attributes_global_variable( - "", "Prop12", "", "", "", "Prop16", "Prop17", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] + "", "Prop12", "", "", "", "Prop16", "Prop17", "Prop18", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] #endif mp { int x; @@ -90,18 +90,18 @@ struct ih : public h {}; template struct igh : public gh {}; template struct ihg : public hg {}; -constexpr g g_v; +constexpr g g_v; constexpr h h_v; -constexpr gh gh_v; -constexpr hg hg_v; +constexpr gh gh_v; +constexpr hg hg_v; constexpr np np_v; constexpr mp mp_v; -constexpr ig ig_v; +constexpr ig ig_v; constexpr ih ih_v; -constexpr igh igh_v; -constexpr ihg ihg_v; +constexpr igh igh_v; +constexpr ihg ihg_v; int main() { sycl::queue q; @@ -132,7 +132,7 @@ int main() { // CHECK-DAG: @_ZL4ih_v = internal addrspace(1) constant {{.*}}, align 4{{$}} // CHECK-DAG: @_ZL5igh_v = internal addrspace(1) constant {{.*}}, align 4{{$}} // CHECK-DAG: @_ZL5ihg_v = internal addrspace(1) constant {{.*}}, align 4{{$}} -// CHECK-DAG: attributes #[[GlobalVarGAttrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} } -// CHECK-DAG: attributes #[[GlobalVarHAttrs]] = { {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} } -// CHECK-DAG: attributes #[[GlobalVarHGAndGHAttrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} } -// CHECK-DAG: attributes #[[GlobalVarMPAttrs]] = { {{.*}}"Prop12"="2"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} } +// CHECK-DAG: attributes #[[GlobalVarGAttrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} } +// CHECK-DAG: attributes #[[GlobalVarHAttrs]] = { {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} } +// CHECK-DAG: attributes #[[GlobalVarHGAndGHAttrs]] = { {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} } +// CHECK-DAG: attributes #[[GlobalVarMPAttrs]] = { {{.*}}"Prop12"="2"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} } diff --git a/clang/test/CodeGenSYCL/add_ir_attributes_kernel_parameter.cpp b/clang/test/CodeGenSYCL/add_ir_attributes_kernel_parameter.cpp index 4a02213b68d86..d5853518b96b7 100644 --- a/clang/test/CodeGenSYCL/add_ir_attributes_kernel_parameter.cpp +++ b/clang/test/CodeGenSYCL/add_ir_attributes_kernel_parameter.cpp @@ -37,8 +37,8 @@ class __attribute__((sycl_special_class)) h { #ifdef __SYCL_DEVICE_ONLY__ void __init( [[__sycl_detail__::add_ir_attributes_kernel_parameter( - "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] int *_x) { + "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] int *_x) { x = _x; } #endif @@ -54,8 +54,8 @@ template class __attribute__((sycl_special_class)) gh { #ifdef __SYCL_DEVICE_ONLY__ void __init( [[__sycl_detail__::add_ir_attributes_kernel_parameter( - Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", - Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] int *_x) { + Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", + Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] int *_x) { x = _x; } #endif @@ -71,8 +71,8 @@ template class __attribute__((sycl_special_class)) hg { #ifdef __SYCL_DEVICE_ONLY__ void __init( [[__sycl_detail__::add_ir_attributes_kernel_parameter( - "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", Properties::name..., - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, Properties::value...)]] int *_x) { + "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", Properties::name..., + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8, Properties::value...)]] int *_x) { x = _x; } #endif @@ -111,8 +111,8 @@ class __attribute__((sycl_special_class)) l { void __init( int *_x, [[__sycl_detail__::add_ir_attributes_kernel_parameter( - "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] float *_y) { + "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] float *_y) { x = _x; y = _y; } @@ -131,8 +131,8 @@ template class __attribute__((sycl_special_class)) kl { void __init( int *_x, [[__sycl_detail__::add_ir_attributes_kernel_parameter( - Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", - Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] float *_y) { + Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", + Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] float *_y) { x = _x; y = _y; } @@ -151,8 +151,8 @@ template class __attribute__((sycl_special_class)) lk { void __init( int *_x, [[__sycl_detail__::add_ir_attributes_kernel_parameter( - "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", Properties::name..., - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, Properties::value...)]] float *_y) { + "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", Properties::name..., + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8, Properties::value...)]] float *_y) { x = _x; y = _y; } @@ -192,11 +192,11 @@ class __attribute__((sycl_special_class)) n { #ifdef __SYCL_DEVICE_ONLY__ void __init( [[__sycl_detail__::add_ir_attributes_kernel_parameter( - "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] int *_x, + "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] int *_x, [[__sycl_detail__::add_ir_attributes_kernel_parameter( - "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] float *_y) { + "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] float *_y) { x = _x; y = _y; } @@ -214,11 +214,11 @@ template class __attribute__((sycl_special_class)) mn { #ifdef __SYCL_DEVICE_ONLY__ void __init( [[__sycl_detail__::add_ir_attributes_kernel_parameter( - Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", - Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] int *_x, + Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", + Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] int *_x, [[__sycl_detail__::add_ir_attributes_kernel_parameter( - Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", - Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] float *_y) { + Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", + Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] float *_y) { x = _x; y = _y; } @@ -236,11 +236,11 @@ template class __attribute__((sycl_special_class)) nm { #ifdef __SYCL_DEVICE_ONLY__ void __init( [[__sycl_detail__::add_ir_attributes_kernel_parameter( - "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", Properties::name..., - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, Properties::value...)]] int *_x, + "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", Properties::name..., + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8, Properties::value...)]] int *_x, [[__sycl_detail__::add_ir_attributes_kernel_parameter( - "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", Properties::name..., - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, Properties::value...)]] float *_y) { + "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", "Prop18", Properties::name..., + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8, Properties::value...)]] float *_y) { x = _x; y = _y; } @@ -259,8 +259,8 @@ class __attribute__((sycl_special_class)) np { #ifdef __SYCL_DEVICE_ONLY__ void __init( [[__sycl_detail__::add_ir_attributes_kernel_parameter( - "", "", "", "", "", "", "", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] int *_x) { + "", "", "", "", "", "", "", "", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] int *_x) { x = _x; } #endif @@ -276,8 +276,8 @@ class __attribute__((sycl_special_class)) mp { #ifdef __SYCL_DEVICE_ONLY__ void __init( [[__sycl_detail__::add_ir_attributes_kernel_parameter( - "", "Prop12", "", "", "", "Prop16", "Prop17", - "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]] int *_x) { + "", "Prop12", "", "", "", "Prop16", "Prop17", "Prop18", + "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, PropertyValue8)]] int *_x) { x = _x; } #endif @@ -285,7 +285,7 @@ class __attribute__((sycl_special_class)) mp { int main() { sycl::queue q; - g a1; + g a1; q.submit([&](sycl::handler &h) { h.single_task( [=]() { @@ -299,21 +299,21 @@ int main() { (void)b1; }); }); - gh c1; + gh c1; q.submit([&](sycl::handler &h) { h.single_task( [=]() { (void)c1; }); }); - hg d1; + hg d1; q.submit([&](sycl::handler &h) { h.single_task( [=]() { (void)d1; }); }); - k a2; + k a2; q.submit([&](sycl::handler &h) { h.single_task( [=]() { @@ -327,21 +327,21 @@ int main() { (void)b2; }); }); - kl c2; + kl c2; q.submit([&](sycl::handler &h) { h.single_task( [=]() { (void)c2; }); }); - lk d2; + lk d2; q.submit([&](sycl::handler &h) { h.single_task( [=]() { (void)d2; }); }); - m a3; + m a3; q.submit([&](sycl::handler &h) { h.single_task( [=]() { @@ -355,14 +355,14 @@ int main() { (void)b3; }); }); - mn c3; + mn c3; q.submit([&](sycl::handler &h) { h.single_task( [=]() { (void)c3; }); }); - nm d3; + nm d3; q.submit([&](sycl::handler &h) { h.single_task( [=]() { @@ -386,26 +386,26 @@ int main() { } // One __init parameter with add_ir_attributes_kernel_parameter attribute. -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel1({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}) -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel2({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} %{{.*}}) -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel3({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}) -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel4({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel1({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel2({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel3({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel4({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}) // Two __init parameters, one with add_ir_attributes_kernel_parameter attribute. -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel5({{.*}}ptr addrspace({{.*}}) {{[^"]*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}) -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel6({{.*}}ptr addrspace({{.*}}) {{[^"]*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} %{{.*}}) -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel7({{.*}}ptr addrspace({{.*}}) {{[^"]*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}) -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel8({{.*}}ptr addrspace({{.*}}) {{[^"]*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel5({{.*}}ptr addrspace({{.*}}) {{[^"]*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel6({{.*}}ptr addrspace({{.*}}) {{[^"]*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel7({{.*}}ptr addrspace({{.*}}) {{[^"]*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel8({{.*}}ptr addrspace({{.*}}) {{[^"]*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}) // Two __init parameters, both with add_ir_attributes_kernel_parameter attribute. -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel9({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}) -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel10({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} %{{.*}}) -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel11({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}) -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel12({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel9({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel10({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel11({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel12({{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}, {{.*}}ptr addrspace({{.*}}) {{.*}}"Prop1"="Property string"{{.*}}"Prop11"="Another property string"{{.*}}"Prop12"="2"{{.*}}"Prop13"="false"{{.*}}"Prop14"="1"{{.*}}"Prop15"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}}"Prop2"="1"{{.*}}"Prop3"="true"{{.*}}"Prop4"="2"{{.*}}"Prop5"{{.*}}"Prop6"{{.*}}"Prop7"="1"{{.*}}"Prop8"="Property"{{.*}} %{{.*}}) // Empty attribute names. // CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel13({{.*}}ptr addrspace({{.*}}) {{.*}} %{{.*}}) -// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel14({{.*}}) {{.*}}"Prop12"="2"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}} %{{.*}}) +// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}test_kernel14({{.*}}) {{.*}}"Prop12"="2"{{.*}}"Prop16"{{.*}}"Prop17"="2"{{.*}}"Prop18"="Property"{{.*}} %{{.*}}) // CHECK-NOT: ""="Another property string" // CHECK-NOT: ""="1" // CHECK-NOT: ""="2" diff --git a/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp b/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp new file mode 100644 index 0000000000000..2b94876cf1644 --- /dev/null +++ b/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp @@ -0,0 +1,42 @@ +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s + +// Tests that add_ir_attributes_function causes a warning when appearing with +// potentially conflicting SYCL attributes. + +constexpr const char AttrName1[] = "Attr1"; +constexpr const char AttrVal1[] = "Val1"; + +template struct Wrapper { + template + [[__sycl_detail__::add_ir_attributes_function(Strs...)]] __attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) { + kernelFunc(); + } +}; + +int main() { + Wrapper<> EmptyWrapper; + Wrapper NonemptyWrapper; + + EmptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); + + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); + // expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} + NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); +}