diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index bfafb3da5c30d..d1e2a22702bb5 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -3195,9 +3195,6 @@ def warn_dllimport_dropped_from_inline_function : Warning< InGroup; def warn_attribute_ignored : Warning<"%0 attribute ignored">, InGroup; -def warn_attribute_on_direct_kernel_callee_only : Warning<"%0 attribute allowed" - " only on a function directly called from a SYCL kernel function; attribute ignored">, - InGroup; def warn_nothrow_attribute_ignored : Warning<"'nothrow' attribute conflicts with" " exception specification; attribute ignored">, InGroup; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 32e116b563a3c..53eb54c75bff4 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13017,7 +13017,6 @@ class Sema final { bool isKnownGoodSYCLDecl(const Decl *D); void checkSYCLDeviceVarDecl(VarDecl *Var); - void copySYCLKernelAttrs(const CXXRecordDecl *KernelObj); void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC); void MarkDevice(); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ca1f2ed1dbcbe..d6a43df48e54b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -306,36 +306,6 @@ static int64_t getIntExprValue(const Expr *E, ASTContext &Ctx) { return E->getIntegerConstantExpr(Ctx)->getSExtValue(); } -// Collect function attributes related to SYCL -static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, - llvm::SmallVector &Attrs, - bool DirectlyCalled = true) { - if (!FD->hasAttrs()) - return; - - llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { - return isa(A); - }); - - // Allow the kernel attribute "use_stall_enable_clusters" only on lambda - // functions and function objects called directly from a kernel. - // For all other cases, emit a warning and ignore. - if (auto *A = FD->getAttr()) { - if (DirectlyCalled) { - Attrs.push_back(A); - } else { - S.Diag(A->getLocation(), - diag::warn_attribute_on_direct_kernel_callee_only) - << A; - FD->dropAttr(); - } - } -} - class MarkDeviceFunction : public RecursiveASTVisitor { // Used to keep track of the constexpr depth, so we know whether to skip // diagnostics. @@ -507,7 +477,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { // Returns the kernel body function found during traversal. FunctionDecl * CollectPossibleKernelAttributes(FunctionDecl *SYCLKernel, - llvm::SmallVector &Attrs) { + llvm::SmallPtrSet &Attrs) { typedef std::pair ChildParentPair; llvm::SmallPtrSet Visited; llvm::SmallVector WorkList; @@ -538,23 +508,55 @@ class MarkDeviceFunction : public RecursiveASTVisitor { "function can be called"); KernelBody = FD; } - WorkList.pop_back(); if (!Visited.insert(FD).second) continue; // We've already seen this Decl - // Gather all attributes of FD that are SYCL related. - // Some attributes are allowed only on lambda functions and function - // objects called directly from a kernel (i.e. the one passed to the - // single_task or parallel_for functions). - bool DirectlyCalled = (ParentFD == SYCLKernel); - collectSYCLAttributes(SemaRef, FD, Attrs, DirectlyCalled); + if (auto *A = FD->getAttr()) + Attrs.insert(A); + + if (auto *A = FD->getAttr()) + Attrs.insert(A); + + if (auto *A = FD->getAttr()) + Attrs.insert(A); + + if (auto *A = FD->getAttr()) + Attrs.insert(A); + + if (auto *A = FD->getAttr()) + Attrs.insert(A); + + if (auto *A = FD->getAttr()) + Attrs.insert(A); + + if (auto *A = FD->getAttr()) + Attrs.insert(A); + + if (auto *A = FD->getAttr()) + Attrs.insert(A); + + if (auto *A = FD->getAttr()) + Attrs.insert(A); + + // Allow the kernel attribute "use_stall_enable_clusters" only on lambda + // functions and function objects that are called directly from a kernel + // (i.e. the one passed to the single_task or parallel_for functions). + // For all other cases, emit a warning and ignore. + if (auto *A = FD->getAttr()) { + if (ParentFD == SYCLKernel) { + Attrs.insert(A); + } else { + SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A; + FD->dropAttr(); + } + } // Attribute "loop_fuse" can be applied explicitly on kernel function. // Attribute should not be propagated from device functions to kernel. if (auto *A = FD->getAttr()) { if (ParentFD == SYCLKernel) { - Attrs.push_back(A); + Attrs.insert(A); } } @@ -2056,8 +2058,8 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler { using SyclKernelFieldHandler::handleSyclHalfType; }; -static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) { - for (auto *MD : Rec->methods()) { +static const CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) { + for (const auto *MD : Rec->methods()) { if (MD->getOverloadedOperator() == OO_Call) return MD; } @@ -3147,56 +3149,6 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, KernelFunc->setInvalidDecl(); } -// For a wrapped parallel_for, copy attributes from original -// kernel to wrapped kernel. -void Sema::copySYCLKernelAttrs(const CXXRecordDecl *KernelObj) { - // Get the operator() function of the wrapper - CXXMethodDecl *OpParens = getOperatorParens(KernelObj); - assert(OpParens && "invalid kernel object"); - - typedef std::pair ChildParentPair; - llvm::SmallPtrSet Visited; - llvm::SmallVector WorkList; - WorkList.push_back({OpParens, nullptr}); - FunctionDecl *KernelBody = nullptr; - - CallGraph SYCLCG; - SYCLCG.addToCallGraph(getASTContext().getTranslationUnitDecl()); - while (!WorkList.empty()) { - FunctionDecl *FD = WorkList.back().first; - FunctionDecl *ParentFD = WorkList.back().second; - - if ((ParentFD == OpParens) && isSYCLKernelBodyFunction(FD)) { - KernelBody = FD; - break; - } - - WorkList.pop_back(); - if (!Visited.insert(FD).second) - continue; // We've already seen this Decl - - CallGraphNode *N = SYCLCG.getNode(FD); - if (!N) - continue; - - for (const CallGraphNode *CI : *N) { - if (auto *Callee = dyn_cast(CI->getDecl())) { - Callee = Callee->getMostRecentDecl(); - if (!Visited.count(Callee)) - WorkList.push_back({Callee, FD}); - } - } - } - - assert(KernelBody && "improper parallel_for wrap"); - if (KernelBody) { - llvm::SmallVector Attrs; - collectSYCLAttributes(*this, KernelBody, Attrs); - if (!Attrs.empty()) - llvm::for_each(Attrs, [OpParens](Attr *A) { OpParens->addAttr(A); }); - } -} - // Generates the OpenCL kernel using KernelCallerFunc (kernel caller // function) defined is SYCL headers. // Generated OpenCL kernel contains the body of the kernel caller function, @@ -3229,20 +3181,14 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, if (KernelObj->isInvalidDecl()) return; + bool IsSIMDKernel = isESIMDKernelType(KernelObj); + // Calculate both names, since Integration headers need both. std::string CalculatedName, StableName; std::tie(CalculatedName, StableName) = constructKernelName(*this, KernelCallerFunc, MC); StringRef KernelName(getLangOpts().SYCLUnnamedLambda ? StableName : CalculatedName); - - // Attributes of a user-written SYCL kernel must be copied to the internally - // generated alternative kernel, identified by a known string in its name. - if (StableName.find("__pf_kernel_wrapper") != std::string::npos) - copySYCLKernelAttrs(KernelObj); - - bool IsSIMDKernel = isESIMDKernelType(KernelObj); - SyclKernelDeclCreator kernel_decl(*this, KernelName, KernelObj->getLocation(), KernelCallerFunc->isInlined(), IsSIMDKernel); @@ -3280,7 +3226,7 @@ void Sema::MarkDevice(void) { Marker.CollectKernelSet(SYCLKernel, SYCLKernel, VisitedSet); // Let's propagate attributes from device functions to a SYCL kernels - llvm::SmallVector Attrs; + llvm::SmallPtrSet Attrs; // This function collects all kernel attributes which might be applied to // a device functions, but need to be propagated down to callers, i.e. // SYCL kernels diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index a1a8d626b7c40..6162171e84001 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -206,24 +206,11 @@ template struct get_kernel_name_t { using name = Type; }; - -// Used when parallel_for range is rounded-up. -template class __pf_kernel_wrapper; - -template struct get_kernel_wrapper_name_t { - using name = - __pf_kernel_wrapper::name>; -}; - #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) template ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } -template -ATTR_SYCL_KERNEL void kernel_parallel_for(const KernelType &kernelFunc) { - kernelFunc(); -} class handler { public: template @@ -233,16 +220,6 @@ class handler { kernel_single_task(kernelFunc); #else kernelFunc(); -#endif - } - template - void parallel_for(const KernelType &kernelObj) { - using NameT = typename get_kernel_name_t::name; - using NameWT = typename get_kernel_wrapper_name_t::name; -#ifdef __SYCL_DEVICE_ONLY__ - kernel_parallel_for(kernelObj); -#else - kernelObj(); #endif } }; diff --git a/clang/test/SemaSYCL/args-size-overflow.cpp b/clang/test/SemaSYCL/args-size-overflow.cpp index f94822f2be49e..3b442f583b415 100644 --- a/clang/test/SemaSYCL/args-size-overflow.cpp +++ b/clang/test/SemaSYCL/args-size-overflow.cpp @@ -11,9 +11,9 @@ queue q; using Accessor = accessor; #ifdef SPIR64 -// expected-warning@Inputs/sycl.hpp:233 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} +// expected-warning@Inputs/sycl.hpp:220 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} #elif SPIR32 -// expected-warning@Inputs/sycl.hpp:233 {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}} +// expected-warning@Inputs/sycl.hpp:220 {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}} #endif void use() { diff --git a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp index 23bbbf5c0d29e..904b28bda7c87 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp @@ -12,7 +12,7 @@ int main(int argc, char **argv) { _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([]() { _mm_prefetch("test", 4); // expected-error {{builtin is not supported on this target}} _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} expected-error {{builtin is not supported on this target}} diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 5a7c03f5e9546..f634e7bc6933c 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -64,7 +64,7 @@ template void setup_sycl_operation(const T VA[]) { deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([]() { // ======= Zero Length Arrays Not Allowed in Kernel ========== // expected-error@+1 {{zero-length arrays are not permitted in C++}} @@ -156,7 +156,7 @@ int main(int argc, char **argv) { // --- direct lambda testing --- deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 2 {{called by 'kernel_single_task([]() { // expected-error@+1 {{zero-length arrays are not permitted in C++}} int BadArray[0]; diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp index b3b5d06ed4cad..0f9d6180d5652 100644 --- a/clang/test/SemaSYCL/float128.cpp +++ b/clang/test/SemaSYCL/float128.cpp @@ -71,7 +71,7 @@ int main() { __float128 CapturedToDevice = 1; host_ok(); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([=]() { // expected-error@+1 {{'__float128' is not supported on this target}} decltype(CapturedToDevice) D; @@ -88,7 +88,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 4{{called by 'kernel_single_task([=]() { // expected-note@+1 2{{called by 'operator()'}} usage(); @@ -104,7 +104,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([=]() { // expected-note@+1 3{{used here}} Z<__float128> S; diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp index c458859ea1be5..3f4239855055d 100644 --- a/clang/test/SemaSYCL/implicit_kernel_type.cpp +++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp @@ -25,12 +25,12 @@ int main() { queue q; #if defined(WARN) - // expected-error@Inputs/sycl.hpp:233 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:220 {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+8 {{in instantiation of function template specialization}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:233 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:220 {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+4 {{in instantiation of function template specialization}} #endif class InvalidKernelName1 {}; @@ -39,9 +39,9 @@ int main() { }); #if defined(WARN) - // expected-warning@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-warning@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-error@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #endif q.submit([&](handler &h) { @@ -53,9 +53,9 @@ int main() { }); #if defined(WARN) - // expected-warning@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-warning@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-error@Inputs/sycl.hpp:220 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #endif q.submit([&](handler &h) { diff --git a/clang/test/SemaSYCL/int128.cpp b/clang/test/SemaSYCL/int128.cpp index ea9d708083ea4..2eea7b8734715 100644 --- a/clang/test/SemaSYCL/int128.cpp +++ b/clang/test/SemaSYCL/int128.cpp @@ -80,7 +80,7 @@ int main() { __int128 CapturedToDevice = 1; host_ok(); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([=]() { // expected-error@+1 {{'__int128' is not supported on this target}} decltype(CapturedToDevice) D; @@ -97,7 +97,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 5{{called by 'kernel_single_task([=]() { // expected-note@+1 2{{called by 'operator()'}} usage(); @@ -115,7 +115,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([=]() { // expected-note@+1 3{{used here}} Z<__int128> S; diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp index 3d8a4be3e64a3..48e9305d6ecbf 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -86,11 +86,7 @@ int main() { h.single_task(TRIFuncObjGood1()); // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 - // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 0 - // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} - // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + // CHECK: ReqdWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} @@ -100,7 +96,7 @@ int main() { // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} @@ -110,33 +106,37 @@ int main() { // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 0 + // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} h.single_task(TRIFuncObjGood2()); // CHECK-LABEL: FunctionDecl {{.*}}test_kernel5 - // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 3 - // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} - // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + // CHECK: ReqdWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: value: Int 4 + // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 3 + // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} #ifdef TRIGGER_ERROR [[intel::max_global_work_dim(1)]] int Var = 0; // expected-error{{'max_global_work_dim' attribute only applies to functions}} diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp index bb00330af2a05..6a7156c646bac 100644 --- a/clang/test/SemaSYCL/kernelname-enum.cpp +++ b/clang/test/SemaSYCL/kernelname-enum.cpp @@ -67,15 +67,15 @@ int main() { }); q.submit([&](cl::sycl::handler &cgh) { - // expected-error@Inputs/sycl.hpp:233 {{'dummy_functor_2' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} + // expected-error@Inputs/sycl.hpp:220 {{'dummy_functor_2' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task(f2); }); q.submit([&](cl::sycl::handler &cgh) { - // expected-error@Inputs/sycl.hpp:233 {{'templated_functor' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} + // expected-error@Inputs/sycl.hpp:220 {{'templated_functor' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task(f5); }); diff --git a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp deleted file mode 100755 index 79c1c4f6d92a1..0000000000000 --- a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp +++ /dev/null @@ -1,54 +0,0 @@ -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s - -#include "Inputs/sycl.hpp" - -template class Fobj { -public: - Fobj() {} - void operator()() const { - auto L0 = []() [[intel::reqd_sub_group_size(4)]]{}; - L0(); - } -}; - -void invoke() { - sycl::queue q; - q.submit([&](sycl::handler &h) { - Fobj fobj1; - h.parallel_for(fobj1); - }); - q.submit([&](sycl::handler &h) { - Fobj fobj2; - h.parallel_for(fobj2); - }); -} - -// CHECK-LABEL: ClassTemplateSpecializationDecl {{.*}} class Fobj definition -// CHECK: TemplateArgument type 'int' -// CHECK: CXXMethodDecl {{.*}} used operator() 'void () const' -// CHECK: CXXMethodDecl {{.*}} used operator() 'void () const' inline -// CHECK-NEXT: CompoundStmt -// CHECK-NEXT: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 4 -// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} -// CHECK: CXXOperatorCallExpr {{.*}} 'void':'void' '()' -// CHECK: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 4 -// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} -// CHECK: CXXConstructorDecl -// CHECK: CXXConstructorDecl - -// CHECK-LABEL: ClassTemplateSpecializationDecl {{.*}} class Fobj definition -// CHECK: TemplateArgument type 'short' -// CHECK: CXXMethodDecl {{.*}} used operator() 'void () const' -// CHECK: CXXMethodDecl {{.*}} used operator() 'void () const' inline -// CHECK-NEXT: CompoundStmt -// CHECK-NEXT: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 4 -// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} -// CHECK: CXXOperatorCallExpr {{.*}} 'void':'void' '()' -// CHECK-NOT: IntelReqdSubGroupSizeAttr {{.*}} -// CHECK: CXXConstructorDecl diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index 082fe8bf64594..c1fa227b9f73a 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -53,10 +53,6 @@ int main() { #ifndef TRIGGER_ERROR // CHECK-LABEL: FunctionDecl {{.*}} main 'int ()' // CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()' - // CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} // CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 4 @@ -67,6 +63,10 @@ int main() { // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 4 // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} + // CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} // CHECK: `-ReqdWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 2 diff --git a/clang/test/SemaSYCL/stall_enable.cpp b/clang/test/SemaSYCL/stall_enable.cpp index ec496f441469f..45ba33a58a906 100644 --- a/clang/test/SemaSYCL/stall_enable.cpp +++ b/clang/test/SemaSYCL/stall_enable.cpp @@ -6,7 +6,7 @@ using namespace cl::sycl; queue q; -[[intel::use_stall_enable_clusters]] void test() {} // expected-warning{{'use_stall_enable_clusters' attribute allowed only on a function directly called from a SYCL kernel}} +[[intel::use_stall_enable_clusters]] void test() {} //expected-warning{{'use_stall_enable_clusters' attribute ignored}} #ifdef TRIGGER_ERROR [[intel::use_stall_enable_clusters(1)]] void bar1() {} // expected-error{{'use_stall_enable_clusters' attribute takes no arguments}} diff --git a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp index 0d44415bf2394..59d212f64ded7 100644 --- a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp +++ b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp @@ -25,38 +25,38 @@ queue q; int main() { #ifdef CHECK_ERROR q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:233 {{'nullptr_t' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:220 {{'nullptr_t' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:233 {{'std::T' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'std::T' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:220 {{'std::T' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{type 'std::T' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:220 {{'Templated_kernel_name' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task>([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'std::U' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:220 {{'Templated_kernel_name' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{type 'std::U' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task>([=] {}); }); q.submit([&](handler &cgh) { - // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name2>' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233{{type 'std::Foo' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:220 {{'Templated_kernel_name2>' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220{{type 'std::Foo' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task>>([]() {}); }); q.submit([&](handler &cgh) { - // expected-error@Inputs/sycl.hpp:233 {{'TemplParamPack' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:220 {{'TemplParamPack' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task>([]() {}); }); diff --git a/clang/test/SemaSYCL/unnamed-kernel.cpp b/clang/test/SemaSYCL/unnamed-kernel.cpp index 97edfca682745..c4e3746f78679 100644 --- a/clang/test/SemaSYCL/unnamed-kernel.cpp +++ b/clang/test/SemaSYCL/unnamed-kernel.cpp @@ -33,8 +33,8 @@ struct MyWrapper { void test() { cl::sycl::queue q; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:220 {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+4{{in instantiation of function template specialization}} #endif class InvalidKernelName1 {}; @@ -43,8 +43,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName2' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:220 {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName2' should be globally-visible}} // expected-note@+4{{in instantiation of function template specialization}} #endif class InvalidKernelName2 {}; @@ -53,8 +53,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName0' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName0' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -62,8 +62,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName3' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:220 {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName3' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -76,8 +76,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'std::max_align_t' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'std::max_align_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:220 {{'std::max_align_t' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{type 'std::max_align_t' cannot be in the "std" namespace}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -86,8 +86,8 @@ struct MyWrapper { using InvalidAlias = InvalidKernelName4; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName4' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName4' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -96,16 +96,16 @@ struct MyWrapper { using InvalidAlias1 = InvalidKernelName5; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName5' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:220 {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName5' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task>([] {}); }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name2>' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:220 {{'Templated_kernel_name2>' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -117,8 +117,8 @@ struct MyWrapper { int main() { cl::sycl::queue q; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error-re@Inputs/sycl.hpp:233 {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{unnamed type used in a SYCL kernel name}} + // expected-error-re@Inputs/sycl.hpp:220 {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{unnamed type used in a SYCL kernel name}} // expected-note@+2{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); }); diff --git a/sycl/test/kernel_param/attr.cpp b/sycl/test/kernel_param/attr.cpp deleted file mode 100755 index 1037d23f54686..0000000000000 --- a/sycl/test/kernel_param/attr.cpp +++ /dev/null @@ -1,22 +0,0 @@ -// RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -// RUN: FileCheck %s --input-file %t.ll - -// Check copying of parallel_for kernel attributes to wrapper kernel. - -#include -using namespace cl::sycl; - -int main() { - range<1> Size{10}; - { - queue myQueue; - myQueue.submit([&](handler &cgh) { - cgh.parallel_for(Size, [=](item<1> ITEM) - [[intel::reqd_work_group_size(4)]]{}); - }); - } - - return 0; -} - -// CHECK: define {{.*}}spir_kernel void @{{.*}}__pf_kernel_wrapper{{.*}}reqd_work_group_size