diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td index a55227673b656..117590d6c6adf 100644 --- a/clang/include/clang/Basic/DiagnosticGroups.td +++ b/clang/include/clang/Basic/DiagnosticGroups.td @@ -1128,7 +1128,9 @@ def OpenMP : DiagGroup<"openmp", [ ]>; // SYCL warnings -def SyclStrict : DiagGroup<"sycl-strict">; +def Sycl2017Compat : DiagGroup<"sycl-2017-compat">; +def Sycl2020Compat : DiagGroup<"sycl-2020-compat">; +def SyclStrict : DiagGroup<"sycl-strict", [ Sycl2017Compat, Sycl2020Compat]>; def SyclTarget : DiagGroup<"sycl-target">; // Backend warnings. diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 62a46831da8fb..f3b49747761f2 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11005,6 +11005,12 @@ def err_sycl_invalid_property_list_template_param : Error< "%select{property_list|property_list pack argument|buffer_location}0 " "template parameter must be a " "%select{parameter pack|type|non-negative integer}1">; +def warn_sycl_pass_by_value_deprecated + : Warning<"Passing kernel functions by value is deprecated in SYCL 2020">, + InGroup; +def warn_sycl_pass_by_reference_future + : Warning<"Passing of kernel functions by reference is a SYCL 2020 extension">, + InGroup; def warn_sycl_attibute_function_raw_ptr : Warning<"SYCL 1.2.1 specification does not allow %0 attribute applied " "to a function with a raw pointer " diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index 010383c1d5f52..023c31c533398 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -2586,6 +2586,7 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, if (const Arg *A = Args.getLastArg(OPT_sycl_std_EQ)) { Opts.SYCLVersion = llvm::StringSwitch(A->getValue()) .Cases("2017", "1.2.1", "121", "sycl-1.2.1", 2017) + .Case("2020", 2020) .Default(0U); if (Opts.SYCLVersion == 0U) { diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index e0342209d7814..b48917f7b1eb2 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -463,8 +463,13 @@ static void InitializeStandardPredefinedMacros(const TargetInfo &TI, if (LangOpts.SYCL) { // SYCL Version is set to a value when building SYCL applications - if (LangOpts.SYCLVersion == 2017) + if (LangOpts.SYCLVersion == 2017) { Builder.defineMacro("CL_SYCL_LANGUAGE_VERSION", "121"); + Builder.defineMacro("SYCL_LANGUAGE_VERSION", "201707"); + } else if (LangOpts.SYCLVersion == 2020) { + Builder.defineMacro("SYCL_LANGUAGE_VERSION", "202001"); + } + if (LangOpts.SYCLValueFitInMaxInt) Builder.defineMacro("__SYCL_ID_QUERIES_FIT_IN_INT__", "1"); } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f339c4a69b3be..ac0ea7b125185 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -707,7 +707,14 @@ getKernelInvocationKind(FunctionDecl *KernelCallerFunc) { static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) { assert(Caller->getNumParams() > 0 && "Insufficient kernel parameters"); - return Caller->getParamDecl(0)->getType()->getAsCXXRecordDecl(); + + QualType KernelParamTy = Caller->getParamDecl(0)->getType(); + // In SYCL 2020 kernels are now passed by reference. + if (KernelParamTy->isReferenceType()) + return KernelParamTy->getPointeeCXXRecordDecl(); + + // SYCL 1.2.1 + return KernelParamTy->getAsCXXRecordDecl(); } /// Creates a kernel parameter descriptor @@ -2248,6 +2255,18 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, SyclKernelFieldChecker FieldChecker(*this); SyclKernelUnionChecker UnionChecker(*this); + // check that calling kernel conforms to spec + QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType(); + if (KernelParamTy->isReferenceType()) { + // passing by reference, so emit warning if not using SYCL 2020 + if (LangOpts.SYCLVersion < 2020) + Diag(KernelFunc->getLocation(), diag::warn_sycl_pass_by_reference_future); + } else { + // passing by value. emit warning if using SYCL 2020 or greater + if (LangOpts.SYCLVersion > 2017) + Diag(KernelFunc->getLocation(), diag::warn_sycl_pass_by_value_deprecated); + } + KernelObjVisitor Visitor{*this}; DiagnosingSYCLKernel = true; Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker); diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index de3e8999961b0..60d5d2bc2307c 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -271,26 +271,26 @@ class spec_constant { #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) template -ATTR_SYCL_KERNEL void kernel_single_task(KernelType kernelFunc) { +ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } template ATTR_SYCL_KERNEL void -kernel_parallel_for(KernelType KernelFunc) { +kernel_parallel_for(const KernelType &KernelFunc) { KernelFunc(id()); } template ATTR_SYCL_KERNEL void -kernel_parallel_for_work_group(KernelType KernelFunc) { +kernel_parallel_for_work_group(const KernelType &KernelFunc) { KernelFunc(group()); } class handler { public: template - void parallel_for(range numWorkItems, KernelType kernelFunc) { + void parallel_for(range numWorkItems, const KernelType &kernelFunc) { using NameT = typename get_kernel_name_t::name; #ifdef __SYCL_DEVICE_ONLY__ kernel_parallel_for(kernelFunc); @@ -300,7 +300,7 @@ class handler { } template - void parallel_for_work_group(range numWorkGroups, range WorkGroupSize, KernelType kernelFunc) { + void parallel_for_work_group(range numWorkGroups, range WorkGroupSize, const KernelType &kernelFunc) { using NameT = typename get_kernel_name_t::name; #ifdef __SYCL_DEVICE_ONLY__ kernel_parallel_for_work_group(kernelFunc); @@ -311,7 +311,7 @@ class handler { } template - void single_task(KernelType kernelFunc) { + void single_task(const KernelType &kernelFunc) { using NameT = typename get_kernel_name_t::name; #ifdef __SYCL_DEVICE_ONLY__ kernel_single_task(kernelFunc); diff --git a/clang/test/CodeGenSYCL/address-space-cond-op.cpp b/clang/test/CodeGenSYCL/address-space-cond-op.cpp index b43ed4e380f49..b2428dca72dfd 100644 --- a/clang/test/CodeGenSYCL/address-space-cond-op.cpp +++ b/clang/test/CodeGenSYCL/address-space-cond-op.cpp @@ -24,7 +24,7 @@ S foo(bool cond, S &lhs, S rhs) { } template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/address-space-new.cpp b/clang/test/CodeGenSYCL/address-space-new.cpp index 827215cec8ef4..1caf5d49dd206 100644 --- a/clang/test/CodeGenSYCL/address-space-new.cpp +++ b/clang/test/CodeGenSYCL/address-space-new.cpp @@ -110,13 +110,11 @@ void test() { // CHECK: call spir_func void @{{.*}}bar{{.*}}(%struct.{{.*}}.HasX addrspace(4)* align 4 dereferenceable(4) %[[SECOND]]) } - template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } - int main() { kernel_single_task([]() { test(); }); return 0; diff --git a/clang/test/CodeGenSYCL/address-space-of-returns.cpp b/clang/test/CodeGenSYCL/address-space-of-returns.cpp index 69104c261ddf7..24bd762bb28d8 100644 --- a/clang/test/CodeGenSYCL/address-space-of-returns.cpp +++ b/clang/test/CodeGenSYCL/address-space-of-returns.cpp @@ -29,7 +29,7 @@ A ret_agg() { // CHECK: define spir_func void @{{.*}}ret_agg{{.*}}(%struct.{{.*}}.A addrspace(4)* {{.*}} %agg.result) template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp index 632453486c2f9..cbb645009bb34 100644 --- a/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp +++ b/clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp @@ -195,7 +195,7 @@ void usages2() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } int main() { diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp index 31795fc73b776..7c89a0154ec27 100644 --- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp @@ -6,7 +6,7 @@ #include "sycl.hpp" template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/const-wg-init.cpp b/clang/test/CodeGenSYCL/const-wg-init.cpp index 07a2a746dffbd..a31a008b17520 100644 --- a/clang/test/CodeGenSYCL/const-wg-init.cpp +++ b/clang/test/CodeGenSYCL/const-wg-init.cpp @@ -4,7 +4,7 @@ template __attribute__((sycl_kernel)) void -kernel_parallel_for_work_group(KernelType KernelFunc) { +kernel_parallel_for_work_group(const KernelType &KernelFunc) { cl::sycl::group<1> G; KernelFunc(G); } diff --git a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp index fcff6aae4f763..5cfd9a939913d 100644 --- a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp +++ b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp @@ -11,7 +11,7 @@ #include template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp index d52ba4c13a7f7..f2c180a80ccf3 100644 --- a/clang/test/CodeGenSYCL/device-functions.cpp +++ b/clang/test/CodeGenSYCL/device-functions.cpp @@ -13,7 +13,7 @@ T bar(T arg) { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/device-variables.cpp b/clang/test/CodeGenSYCL/device-variables.cpp index 83d6c4258995c..e02ac0f6ae42b 100644 --- a/clang/test/CodeGenSYCL/device-variables.cpp +++ b/clang/test/CodeGenSYCL/device-variables.cpp @@ -11,7 +11,7 @@ static constexpr int my_array[1] = {42}; void foo(const test_type &) {} template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp b/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp index 081d6b5d93c70..2c4734b509c60 100644 --- a/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp +++ b/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp b/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp index 2120176c5cd3a..15d11b379fa2d 100644 --- a/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp +++ b/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp @@ -1,8 +1,8 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { - kernelFunc(); +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { + kernelFunc(); } template diff --git a/clang/test/CodeGenSYCL/esimd_metadata1.cpp b/clang/test/CodeGenSYCL/esimd_metadata1.cpp index 554f508ef4f32..0f776e9b43ddf 100644 --- a/clang/test/CodeGenSYCL/esimd_metadata1.cpp +++ b/clang/test/CodeGenSYCL/esimd_metadata1.cpp @@ -9,7 +9,7 @@ // 3. Proper module !spirv.Source metadata is generated template -void kernel(Func f) __attribute__((sycl_kernel)) { +void kernel(const Func &f) __attribute__((sycl_kernel)) { f(); } diff --git a/clang/test/CodeGenSYCL/fpga_pipes.cpp b/clang/test/CodeGenSYCL/fpga_pipes.cpp index bf1ed38afe24f..237f9988f5893 100644 --- a/clang/test/CodeGenSYCL/fpga_pipes.cpp +++ b/clang/test/CodeGenSYCL/fpga_pipes.cpp @@ -45,7 +45,7 @@ class pipe { }; template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp index 4ac785336fb39..f648b244e552d 100644 --- a/clang/test/CodeGenSYCL/inheritance.cpp +++ b/clang/test/CodeGenSYCL/inheritance.cpp @@ -24,7 +24,7 @@ struct base { struct derived : base, second_base { int a; - void operator()() { + void operator()() const { } }; diff --git a/clang/test/CodeGenSYCL/inline_asm.cpp b/clang/test/CodeGenSYCL/inline_asm.cpp index db1956673a15a..be63c64f07e7a 100644 --- a/clang/test/CodeGenSYCL/inline_asm.cpp +++ b/clang/test/CodeGenSYCL/inline_asm.cpp @@ -3,7 +3,7 @@ class kernel; template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [100 x i32], align 4 // CHECK: %[[IDX:.*]] = getelementptr inbounds [100 x i32], [100 x i32]* %[[ARRAY_A]], i64 0, i64 0 int a[100], i = 0; diff --git a/clang/test/CodeGenSYCL/inlining.cpp b/clang/test/CodeGenSYCL/inlining.cpp index da01f2b70a8cb..a816d16f88a01 100644 --- a/clang/test/CodeGenSYCL/inlining.cpp +++ b/clang/test/CodeGenSYCL/inlining.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice %s -S -emit-llvm -o - | FileCheck %s template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp index 65b1b6cdab056..76e37c4a45b5e 100644 --- a/clang/test/CodeGenSYCL/int_header1.cpp +++ b/clang/test/CodeGenSYCL/int_header1.cpp @@ -22,7 +22,7 @@ #include "sycl.hpp" template -__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/int_header_inline_ns.cpp b/clang/test/CodeGenSYCL/int_header_inline_ns.cpp index dedeb0345f0c5..efe4e12e98530 100644 --- a/clang/test/CodeGenSYCL/int_header_inline_ns.cpp +++ b/clang/test/CodeGenSYCL/int_header_inline_ns.cpp @@ -8,7 +8,7 @@ #include "sycl.hpp" template -__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index e285bfce8f536..65c1e22be1889 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -68,7 +68,7 @@ #include "sycl.hpp" template -__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } struct x {}; diff --git a/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp b/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp index 707eeff1ea392..a59d4733e59c3 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp @@ -178,7 +178,7 @@ void ivdep_struct() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp index c90671c8f04a0..72434fe3480af 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp @@ -148,7 +148,7 @@ void ivdep_embedded_multiple_dimensions() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp b/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp index df745e91d63d8..c6a0298336136 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp @@ -80,7 +80,7 @@ void ivdep_conflicting_safelen() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-local.cpp b/clang/test/CodeGenSYCL/intel-fpga-local.cpp index 41a483f33babc..ad02d32e2aca1 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-local.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-local.cpp @@ -256,7 +256,7 @@ void field_addrspace_cast() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp index 0556feadcf7a2..ce7fb46b9c66c 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-loops.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-loops.cpp @@ -95,7 +95,7 @@ void speculated_iterations() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp b/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp index e330fffb8eb65..abc75cf10dd85 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp @@ -71,7 +71,7 @@ void foo(float *A, int *B, State *C, State &D) { // CHECK-DAG: attributes [[ATT]] = { readnone } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp index d1352b190fa94..3b2ffd1cae5e4 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp @@ -2,11 +2,11 @@ class Foo { public: - [[intelfpga::no_global_work_offset(1)]] void operator()() {} + [[intelfpga::no_global_work_offset(1)]] void operator()() const {} }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/intel-fpga-reg.cpp b/clang/test/CodeGenSYCL/intel-fpga-reg.cpp index 9428243813f40..066ccc4064933 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-reg.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-reg.cpp @@ -143,7 +143,7 @@ void foo() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp index 5208db6ec3908..0fcd1ab014864 100644 --- a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp +++ b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp @@ -2,11 +2,11 @@ class Foo { public: - [[intelfpga::max_global_work_dim(1)]] void operator()() {} + [[intelfpga::max_global_work_dim(1)]] void operator()() const {} }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp index 13bbb54f34198..e9b4360701e09 100644 --- a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp @@ -2,11 +2,11 @@ class Foo { public: - [[intelfpga::max_work_group_size(1, 1, 1)]] void operator()() {} + [[intelfpga::max_work_group_size(1, 1, 1)]] void operator()() const {} }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/intel-restrict.cpp b/clang/test/CodeGenSYCL/intel-restrict.cpp index ec4d9f9685449..12b870ba625f1 100644 --- a/clang/test/CodeGenSYCL/intel-restrict.cpp +++ b/clang/test/CodeGenSYCL/intel-restrict.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device %s -emit-llvm -triple spir64-unknown-unknown-sycldevice -o - | FileCheck %s template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/kernel-by-reference.cpp b/clang/test/CodeGenSYCL/kernel-by-reference.cpp new file mode 100644 index 0000000000000..575ed9af6dac5 --- /dev/null +++ b/clang/test/CodeGenSYCL/kernel-by-reference.cpp @@ -0,0 +1,42 @@ +// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -sycl-std=2020 -DSYCL2020 %s +// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -Wno-sycl-strict -DNODIAG %s +// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only -sycl-std=2020 -Wno-sycl-strict -DNODIAG %s + +// SYCL 1.2/2017 - kernel functions passed directly. (Also no const requirement, though mutable lambdas never supported) +template +#if defined(SYCL2020) +// expected-warning@+2 {{Passing kernel functions by value is deprecated in SYCL 2020}} +#endif +__attribute__((sycl_kernel)) void sycl_2017_single_task(Func kernelFunc) { + kernelFunc(); +} + +// SYCL 2020 - kernel functions are passed by reference. +template +#if defined(SYCL2017) +// expected-warning@+2 {{Passing of kernel functions by reference is a SYCL 2020 extension}} +#endif +__attribute__((sycl_kernel)) void sycl_2020_single_task(const Func &kernelFunc) { + kernelFunc(); +} + +int do_nothing(int i) { + return i + 1; +} + +// ensure both compile. +int main() { + sycl_2017_single_task([]() { + do_nothing(10); + }); + + sycl_2020_single_task([]() { + do_nothing(11); + }); + + return 0; +} +#if defined(NODIAG) +// expected-no-diagnostics +#endif \ No newline at end of file diff --git a/clang/test/CodeGenSYCL/kernel-name.cpp b/clang/test/CodeGenSYCL/kernel-name.cpp index 69298f9f8b164..167203fd266a7 100644 --- a/clang/test/CodeGenSYCL/kernel-name.cpp +++ b/clang/test/CodeGenSYCL/kernel-name.cpp @@ -9,7 +9,7 @@ inline namespace cl { using namespace cl::sycl; template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp index 902ad7ddc339c..7bf277437a032 100644 --- a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp @@ -37,7 +37,7 @@ using namespace cl::sycl; template -__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp index ec8ac8bc01f5f..cf17a9d7e3e83 100644 --- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp @@ -7,7 +7,7 @@ using namespace cl::sycl; template -__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp index ddfaca9664851..242d3c2f50697 100644 --- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp @@ -37,7 +37,7 @@ using namespace cl::sycl; template -__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp index 1b1b25dcd3ff4..b77eecfc85d68 100644 --- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp @@ -7,7 +7,7 @@ using namespace cl::sycl; template -__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp index d4a5c8d5995a0..b0dc2132fb06a 100755 --- a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp @@ -39,7 +39,7 @@ using namespace cl::sycl; template -__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp index 3ff1669dc4d17..78fb63c714768 100644 --- a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp @@ -7,7 +7,7 @@ using namespace cl::sycl; template -__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/kernel_functor.cpp b/clang/test/CodeGenSYCL/kernel_functor.cpp index 8712895c20656..b75e4d465f888 100644 --- a/clang/test/CodeGenSYCL/kernel_functor.cpp +++ b/clang/test/CodeGenSYCL/kernel_functor.cpp @@ -9,32 +9,9 @@ constexpr auto sycl_read_write = cl::sycl::access::mode::read_write; constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer; // Case 1: -// - functor class is defined in an anonymous namespace -// - the '()' operator: -// * does not have parameters (to be used in 'single_task'). -// * has no 'const' qualifier -namespace { - class Functor1 { - public: - Functor1(int X_, cl::sycl::accessor &Acc_) : - X(X_), Acc(Acc_) - {} - - void operator()() { - Acc.use(X); - } - - private: - int X; - cl::sycl::accessor Acc; - }; -} - -// Case 2: // - functor class is defined in a namespace // - the '()' operator: // * does not have parameters (to be used in 'single_task'). -// * has the 'const' qualifier namespace ns { class Functor2 { public: @@ -52,31 +29,10 @@ namespace ns { }; } -// Case 3: -// - functor class is templated and defined in the translation unit scope -// - the '()' operator: -// * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for'). -// * has no 'const' qualifier -template class TmplFunctor { -public: - TmplFunctor(T X_, cl::sycl::accessor &Acc_) : - X(X_), Acc(Acc_) - {} - - void operator()(cl::sycl::id<1> id) { - Acc.use(id, X); - } - -private: - T X; - cl::sycl::accessor Acc; -}; - -// Case 4: +// Case 2: // - functor class is templated and defined in the translation unit scope // - the '()' operator: // * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for'). -// * has the 'const' qualifier template class TmplConstFunctor { public: TmplConstFunctor(T X_, cl::sycl::accessor &Acc_) : @@ -99,12 +55,6 @@ int foo(int X) { cl::sycl::queue Q; cl::sycl::buffer Buf(A, 1); - Q.submit([&](cl::sycl::handler& cgh) { - auto Acc = Buf.get_access(cgh); - Functor1 F(X, Acc); - - cgh.single_task(F); - }); Q.submit([&](cl::sycl::handler& cgh) { auto Acc = Buf.get_access(cgh); ns::Functor2 F(X, Acc); @@ -129,13 +79,6 @@ template T bar(T X) { { cl::sycl::queue Q; cl::sycl::buffer Buf(A, ARR_LEN(A)); - - Q.submit([&](cl::sycl::handler& cgh) { - auto Acc = Buf.template get_access(cgh); - TmplFunctor F(X, Acc); - - cgh.parallel_for(cl::sycl::range<1>(ARR_LEN(A)), F); - }); // Spice with lambdas to make sure functors and lambdas work together. Q.submit([&](cl::sycl::handler& cgh) { auto Acc = Buf.template get_access(cgh); @@ -165,12 +108,8 @@ int main() { const int Gold2 = 80; #ifndef __SYCL_DEVICE_ONLY__ - cl::sycl::detail::KernelInfo::getName(); - // CHECK: Functor1 cl::sycl::detail::KernelInfo::getName(); // CHECK: ns::Functor2 - cl::sycl::detail::KernelInfo>::getName(); - // CHECK: TmplFunctor cl::sycl::detail::KernelInfo>::getName(); // CHECK: TmplConstFunctor #endif // __SYCL_DEVICE_ONLY__ diff --git a/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp b/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp index 7a1bb7bb447d4..543b3ddfc3f5b 100644 --- a/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp +++ b/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp @@ -4,12 +4,12 @@ #include "sycl.hpp" template -__attribute__((sycl_kernel)) void single_task(KernelType kernelFunc) { +__attribute__((sycl_kernel)) void single_task(const KernelType &kernelFunc) { kernelFunc(); } struct dummy_functor { - void operator()() {} + void operator()() const {} }; typedef int int_t; diff --git a/clang/test/CodeGenSYCL/kernelname-enum.cpp b/clang/test/CodeGenSYCL/kernelname-enum.cpp index 8580f1c532794..3529fa1534b00 100644 --- a/clang/test/CodeGenSYCL/kernelname-enum.cpp +++ b/clang/test/CodeGenSYCL/kernelname-enum.cpp @@ -40,43 +40,43 @@ enum class no_type_set { template class dummy_functor_1 { public: - void operator()() {} + void operator()() const {} }; template class dummy_functor_2 { public: - void operator()() {} + void operator()() const {} }; template class dummy_functor_3 { public: - void operator()() {} + void operator()() const {} }; template class dummy_functor_4 { public: - void operator()() {} + void operator()() const {} }; template class dummy_functor_5 { public: - void operator()() {} + void operator()() const {} }; template class dummy_functor_6 { public: - void operator()() {} + void operator()() const {} }; template class dummy_functor_7 { public: - void operator()() {} + void operator()() const {} }; namespace type_argument_template_enum { @@ -105,7 +105,7 @@ class Baz; template class T> class dummy_functor_8 { public: - void operator()() {} + void operator()() const {} }; int main() { diff --git a/clang/test/CodeGenSYCL/loop_unroll.cpp b/clang/test/CodeGenSYCL/loop_unroll.cpp index 7e80c92db5231..3b9928873a6e3 100644 --- a/clang/test/CodeGenSYCL/loop_unroll.cpp +++ b/clang/test/CodeGenSYCL/loop_unroll.cpp @@ -30,7 +30,7 @@ void disable() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/module-id.cpp b/clang/test/CodeGenSYCL/module-id.cpp index d120ee295c288..08ad6178ddbf9 100644 --- a/clang/test/CodeGenSYCL/module-id.cpp +++ b/clang/test/CodeGenSYCL/module-id.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/noexcept.cpp b/clang/test/CodeGenSYCL/noexcept.cpp index 816f2c43ebfe2..c0f4e55400a90 100644 --- a/clang/test/CodeGenSYCL/noexcept.cpp +++ b/clang/test/CodeGenSYCL/noexcept.cpp @@ -48,7 +48,7 @@ void foo_cleanup() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/num-simd-work-items.cpp b/clang/test/CodeGenSYCL/num-simd-work-items.cpp index 8b8b8ba22d0da..5613969b88a91 100644 --- a/clang/test/CodeGenSYCL/num-simd-work-items.cpp +++ b/clang/test/CodeGenSYCL/num-simd-work-items.cpp @@ -2,11 +2,11 @@ class Foo { public: - [[intelfpga::num_simd_work_items(1)]] void operator()() {} + [[intelfpga::num_simd_work_items(1)]] void operator()() const {} }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/remove-ur-inst.cpp b/clang/test/CodeGenSYCL/remove-ur-inst.cpp index 77792438d4811..4b799f0996ab1 100644 --- a/clang/test/CodeGenSYCL/remove-ur-inst.cpp +++ b/clang/test/CodeGenSYCL/remove-ur-inst.cpp @@ -4,7 +4,7 @@ SYCL_EXTERNAL void doesNotReturn() throw() __attribute__((__noreturn__)); template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp index d95b97180b3f6..1b62cfa810525 100644 --- a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp @@ -2,14 +2,14 @@ class Functor16 { public: - [[cl::intel_reqd_sub_group_size(16)]] void operator()() {} + [[cl::intel_reqd_sub_group_size(16)]] void operator()() const {} }; [[cl::intel_reqd_sub_group_size(8)]] void foo() {} class Functor { public: - void operator()() { + void operator()() const { foo(); } }; @@ -17,11 +17,11 @@ class Functor { template class Functor5 { public: - [[cl::intel_reqd_sub_group_size(SIZE)]] void operator()() {} + [[cl::intel_reqd_sub_group_size(SIZE)]] void operator()() const {} }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp index bfb08c7ce6c2d..72ba3492d2ed0 100644 --- a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp @@ -2,20 +2,20 @@ class Functor32x16x16 { public: - [[cl::reqd_work_group_size(32, 16, 16)]] void operator()() {} + [[cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {} }; [[cl::reqd_work_group_size(8, 1, 1)]] void f8x1x1() {} class Functor { public: - void operator()() { + void operator()() const { f8x1x1(); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp index 39755aa6b9aec..68c856f4532a3 100644 --- a/clang/test/CodeGenSYCL/sampler.cpp +++ b/clang/test/CodeGenSYCL/sampler.cpp @@ -43,7 +43,7 @@ struct sampler_wrapper { }; template -__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp index bed31dcb96e48..35f1b280881e6 100644 --- a/clang/test/CodeGenSYCL/spir-calling-conv.cpp +++ b/clang/test/CodeGenSYCL/spir-calling-conv.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/spir-enum.cpp b/clang/test/CodeGenSYCL/spir-enum.cpp index 738d1337e02ab..185ef144c7dd3 100644 --- a/clang/test/CodeGenSYCL/spir-enum.cpp +++ b/clang/test/CodeGenSYCL/spir-enum.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/spir-opencl-version.cpp b/clang/test/CodeGenSYCL/spir-opencl-version.cpp index 43e187b463156..5729020883565 100644 --- a/clang/test/CodeGenSYCL/spir-opencl-version.cpp +++ b/clang/test/CodeGenSYCL/spir-opencl-version.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/sycl-device-static-init.cpp b/clang/test/CodeGenSYCL/sycl-device-static-init.cpp index b811b1061bfd6..95b220388d894 100644 --- a/clang/test/CodeGenSYCL/sycl-device-static-init.cpp +++ b/clang/test/CodeGenSYCL/sycl-device-static-init.cpp @@ -28,7 +28,7 @@ template const int BaseInit::var = 9; template struct BaseInit; template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } int main() { diff --git a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp index 3df365b3fef10..457766d4c5c11 100644 --- a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp +++ b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp @@ -2,12 +2,11 @@ class Functor { public: - [[cl::intel_reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() {} + [[cl::intel_reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {} }; - template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/unique-stable-name.cpp b/clang/test/CodeGenSYCL/unique-stable-name.cpp index f24bef81bb8c8..66ca499e6cdae 100644 --- a/clang/test/CodeGenSYCL/unique-stable-name.cpp +++ b/clang/test/CodeGenSYCL/unique-stable-name.cpp @@ -33,7 +33,7 @@ void lambda_in_dependent_function() { {DEF_IN_MACRO();}{DEF_IN_MACRO();} template -[[clang::sycl_kernel]] void kernel_single_task(KernelType kernelFunc) { +[[clang::sycl_kernel]] void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/usm-int-header.cpp b/clang/test/CodeGenSYCL/usm-int-header.cpp index 5b3ffa00050fa..298a7464ac588 100644 --- a/clang/test/CodeGenSYCL/usm-int-header.cpp +++ b/clang/test/CodeGenSYCL/usm-int-header.cpp @@ -17,7 +17,7 @@ #include template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/virtual-types.cpp b/clang/test/CodeGenSYCL/virtual-types.cpp index b78f93712f9a0..48bb1b17456f5 100644 --- a/clang/test/CodeGenSYCL/virtual-types.cpp +++ b/clang/test/CodeGenSYCL/virtual-types.cpp @@ -1,6 +1,6 @@ // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-linux-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/wrapped-accessor.cpp b/clang/test/CodeGenSYCL/wrapped-accessor.cpp index 0355a2a17b15b..bf511f6ae9fa7 100644 --- a/clang/test/CodeGenSYCL/wrapped-accessor.cpp +++ b/clang/test/CodeGenSYCL/wrapped-accessor.cpp @@ -34,7 +34,7 @@ template struct AccWrapper { Acc accessor; }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/Preprocessor/sycl-macro.cpp b/clang/test/Preprocessor/sycl-macro.cpp index 21f94e9fac2df..8a0be0841da47 100644 --- a/clang/test/Preprocessor/sycl-macro.cpp +++ b/clang/test/Preprocessor/sycl-macro.cpp @@ -1,6 +1,7 @@ // RUN: %clang_cc1 %s -E -dM | FileCheck %s // RUN: %clang_cc1 %s -fsycl -fsycl-id-queries-fit-in-int -fsycl-is-host -sycl-std=2017 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD %s // RUN: %clang_cc1 %s -fsycl -fsycl-id-queries-fit-in-int -fsycl-is-device -sycl-std=2017 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD %s +// RUN: %clang_cc1 %s -fsycl -fsycl-id-queries-fit-in-int -fsycl-is-device -sycl-std=2020 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-2020 %s // RUN: %clang_cc1 %s -fsycl -fsycl-id-queries-fit-in-int -fsycl-is-device -sycl-std=1.2.1 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-DEVICE %s // RUNx: %clang_cc1 %s -fsycl -fsycl-id-queries-fit-in-int -fsycl-is-device -E -dM -fms-compatibility | FileCheck --check-prefix=CHECK-MSVC %s // RUN: %clang_cc1 -fno-sycl-id-queries-fit-in-int %s -E -dM | FileCheck \ @@ -12,8 +13,11 @@ // CHECK-NOT:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1 // CHECK-SYCL-STD:#define CL_SYCL_LANGUAGE_VERSION 121 +// CHECK-SYCL-STD:#define SYCL_LANGUAGE_VERSION 201707 // CHECK-SYCL-STD:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1 +// CHECK-SYCL-STD-2020:#define SYCL_LANGUAGE_VERSION 202001 + // CHECK-SYCL-STD-DEVICE:#define __SYCL_DEVICE_ONLY__ 1 // CHECK-SYCL-STD-DEVICE:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1 diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 1b3ab8b20763f..5e3478abf5176 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -202,13 +202,13 @@ struct get_kernel_name_t { }; #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) template -ATTR_SYCL_KERNEL void kernel_single_task(KernelType kernelFunc) { +ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } class handler { public: template - void single_task(KernelType kernelFunc) { + void single_task(const KernelType &kernelFunc) { using NameT = typename get_kernel_name_t::name; #ifdef __SYCL_DEVICE_ONLY__ kernel_single_task(kernelFunc); diff --git a/clang/test/SemaSYCL/accessor-type-diagnostics.cpp b/clang/test/SemaSYCL/accessor-type-diagnostics.cpp index 3a814fbf78011..a409e35527e4e 100644 --- a/clang/test/SemaSYCL/accessor-type-diagnostics.cpp +++ b/clang/test/SemaSYCL/accessor-type-diagnostics.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -I %S/Inputs -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only -Wno-sycl-2017-compat %s // // Ensure SYCL type restrictions are applied to accessors as well. @@ -6,11 +6,6 @@ using namespace cl::sycl; -template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { - kernelFunc(); -} - //alias template template using int128alias_t = __uint128_t; @@ -42,7 +37,7 @@ int main() { // -- Accessor of struct that contains a prohibited type. accessor struct_acc; - kernel( + kernel_single_task( [=]() { ok_acc.use(); diff --git a/clang/test/SemaSYCL/accessors-targets-image.cpp b/clang/test/SemaSYCL/accessors-targets-image.cpp index 18fac9940cb1f..385b1d11bf219 100644 --- a/clang/test/SemaSYCL/accessors-targets-image.cpp +++ b/clang/test/SemaSYCL/accessors-targets-image.cpp @@ -8,7 +8,7 @@ using namespace cl::sycl; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp index dbaab2664e95c..b589bb462c69c 100644 --- a/clang/test/SemaSYCL/accessors-targets.cpp +++ b/clang/test/SemaSYCL/accessors-targets.cpp @@ -8,7 +8,7 @@ using namespace cl::sycl; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/allow-constexpr-recursion.cpp b/clang/test/SemaSYCL/allow-constexpr-recursion.cpp index fed70a02f7021..6b17af6351e22 100644 --- a/clang/test/SemaSYCL/allow-constexpr-recursion.cpp +++ b/clang/test/SemaSYCL/allow-constexpr-recursion.cpp @@ -1,7 +1,7 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -verify -fsyntax-only -std=c++20 -Werror=vla %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-sycl-2017-compat -verify -fsyntax-only -std=c++20 -Werror=vla %s template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/array-kernel-param-neg.cpp b/clang/test/SemaSYCL/array-kernel-param-neg.cpp index 0618014c9fb10..5ad8569feebe8 100755 --- a/clang/test/SemaSYCL/array-kernel-param-neg.cpp +++ b/clang/test/SemaSYCL/array-kernel-param-neg.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-sycl-2017-compat -verify -fsyntax-only %s // This test checks if compiler reports compilation error on an attempt to pass // an array of non-trivially copyable structs as SYCL kernel parameter or @@ -20,11 +20,11 @@ class E { int i[]; public: - int operator()() { return i[0]; } + int operator()() const { return i[0]; } }; template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index 548077be4f82b..2308d6a3360dc 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -8,7 +8,7 @@ using namespace cl::sycl; template -__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void a_kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp index 54ae68ae5b8ce..9d32dd1445453 100644 --- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp @@ -9,7 +9,7 @@ template struct AccWrapper { Acc accessor; }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index bd2daaa01ba41..65742af0cbe1c 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s -// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -verify -pedantic -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -Wno-sycl-2017-compat -verify -pedantic -DTRIGGER_ERROR %s #include "sycl.hpp" diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp index ef3751cd76557..f9a219f407529 100644 --- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp +++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp @@ -6,7 +6,7 @@ #include template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/call-to-undefined-function.cpp b/clang/test/SemaSYCL/call-to-undefined-function.cpp index cc30af261936e..8fd8a6fd7ca73 100644 --- a/clang/test/SemaSYCL/call-to-undefined-function.cpp +++ b/clang/test/SemaSYCL/call-to-undefined-function.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -verify -fsyntax-only %s void defined() { } @@ -9,7 +9,7 @@ void undefined(); SYCL_EXTERNAL void undefinedExternal(); template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp index 8da7b60bcdf67..b4df077725393 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp @@ -1,10 +1,10 @@ -// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl -fsycl-is-device -aux-triple x86_64-unknown-linux-gnu -verify -fsyntax-only %s +// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl -fsycl-is-device -aux-triple x86_64-unknown-linux-gnu -Wno-sycl-2017-compat -verify -fsyntax-only %s inline namespace cl { namespace sycl { template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // expected-note@+1 {{called by 'kernel_single_task -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // expected-note@+1 3{{called by 'kernel_single_task; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp index 0904e740d5083..b61787a153fd0 100644 --- a/clang/test/SemaSYCL/float128.cpp +++ b/clang/test/SemaSYCL/float128.cpp @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only %s -// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsycl -fsycl-is-device -fsyntax-only %s +// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -verify -fsyntax-only %s +// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsycl -fsycl-is-device -Wno-sycl-2017-compat -fsyntax-only %s typedef __float128 BIGTY; @@ -63,7 +63,7 @@ void foo2(){}; __float128 foo(__float128 P) { return P; } template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { // expected-note@+1 6{{called by 'kernel}} kernelFunc(); } diff --git a/clang/test/SemaSYCL/forward-decl.cpp b/clang/test/SemaSYCL/forward-decl.cpp index 4a7a741991edb..7d252af5abc89 100644 --- a/clang/test/SemaSYCL/forward-decl.cpp +++ b/clang/test/SemaSYCL/forward-decl.cpp @@ -1,8 +1,8 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -pedantic %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -pedantic %s // expected-no-diagnostics template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/fpga_pipes.cpp b/clang/test/SemaSYCL/fpga_pipes.cpp index 7f988c0ff49e0..ef25438a08162 100644 --- a/clang/test/SemaSYCL/fpga_pipes.cpp +++ b/clang/test/SemaSYCL/fpga_pipes.cpp @@ -34,7 +34,7 @@ template pipe_storage Storage5 __attribute__((io_pipe_id(N))); template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp index 97a0ca90a2bc9..b80dbc14ce380 100644 --- a/clang/test/SemaSYCL/implicit_kernel_type.cpp +++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp @@ -1,13 +1,42 @@ // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -Werror=sycl-strict -DERROR -// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -Wsycl-strict -DWARN -// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s -Werror=sycl-strict -#include +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -Wsycl-strict -DWARN +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s -Werror=sycl-strict + +// SYCL 1.2 Definitions +template +__attribute__((sycl_kernel)) void sycl_121_single_task(Func kernelFunc) { + kernelFunc(); +} + +class event {}; +class queue { +public: + template + event submit(T cgf) { return event{}; } +}; +class auto_name {}; +template +struct get_kernel_name_t { + using name = Name; +}; +class handler { +public: + template + void single_task(KernelType kernelFunc) { + using NameT = typename get_kernel_name_t::name; +#ifdef __SYCL_DEVICE_ONLY__ + sycl_121_single_task(kernelFunc); +#else + kernelFunc(); +#endif + } +}; +// -- /Definitions #ifdef __SYCL_UNNAMED_LAMBDA__ // expected-no-diagnostics #endif -using namespace cl::sycl; void function() { } @@ -20,11 +49,11 @@ struct myWrapper { class myWrapper2; int main() { - cl::sycl::queue q; + queue q; #ifndef __SYCL_UNNAMED_LAMBDA__ // expected-note@+1 {{InvalidKernelName1 declared here}} class InvalidKernelName1 {}; - q.submit([&](cl::sycl::handler &h) { + q.submit([&](handler &h) { // expected-error@+1 {{kernel needs to have a globally-visible name}} h.single_task([]() {}); }); @@ -36,7 +65,7 @@ int main() { // expected-error@+3 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} // expected-note@+2 {{fake_kernel declared here}} #endif - cl::sycl::kernel_single_task([]() { function(); }); + sycl_121_single_task([]() { function(); }); #if defined(WARN) // expected-warning@+6 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} // expected-note@+5 {{fake_kernel2 declared here}} @@ -44,10 +73,10 @@ int main() { // expected-error@+3 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} // expected-note@+2 {{fake_kernel2 declared here}} #endif - cl::sycl::kernel_single_task([]() { + sycl_121_single_task([]() { auto l = [](auto f) { f(); }; }); - cl::sycl::kernel_single_task([]() { function(); }); - cl::sycl::kernel_single_task([]() { function(); }); + sycl_121_single_task([]() { function(); }); + sycl_121_single_task([]() { function(); }); return 0; } diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp index ff0b263449a35..b801c416cada2 100644 --- a/clang/test/SemaSYCL/inheritance.cpp +++ b/clang/test/SemaSYCL/inheritance.cpp @@ -24,7 +24,7 @@ struct base { struct derived : base, second_base { int a; - void operator()() { + void operator()() const { } }; diff --git a/clang/test/SemaSYCL/inline-asm.cpp b/clang/test/SemaSYCL/inline-asm.cpp index 36237ee3fd38a..bb8a03b7a4090 100644 --- a/clang/test/SemaSYCL/inline-asm.cpp +++ b/clang/test/SemaSYCL/inline-asm.cpp @@ -1,6 +1,6 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify %s -DLINUX_ASM -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify %s -DLINUX_ASM -DSPIR_CHECK -triple spir64-unknown-unknown-sycldevice -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -triple x86_64-windows -fasm-blocks %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s -DLINUX_ASM +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s -DLINUX_ASM -DSPIR_CHECK -triple spir64-unknown-unknown-sycldevice +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -triple x86_64-windows -fasm-blocks %s #ifndef SPIR_CHECK //expected-no-diagnostics @@ -50,7 +50,7 @@ void bar() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); #ifdef LINUX_ASM __asm__("int3"); diff --git a/clang/test/SemaSYCL/intel-fpga-local.cpp b/clang/test/SemaSYCL/intel-fpga-local.cpp index cb406b84f6e38..4d637a4a9213f 100644 --- a/clang/test/SemaSYCL/intel-fpga-local.cpp +++ b/clang/test/SemaSYCL/intel-fpga-local.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -fcxx-exceptions -fsyntax-only -ast-dump -Wno-sycl-2017-compat -verify -pedantic %s | FileCheck %s //CHECK: FunctionDecl{{.*}}check_ast void check_ast() @@ -781,7 +781,7 @@ struct templ_st { }; template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/intel-fpga-loops.cpp b/clang/test/SemaSYCL/intel-fpga-loops.cpp index d32c3cfa4a0c8..164d720eb1560 100644 --- a/clang/test/SemaSYCL/intel-fpga-loops.cpp +++ b/clang/test/SemaSYCL/intel-fpga-loops.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -pedantic %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -pedantic %s // Test for Intel FPGA loop attributes applied not to a loop void foo() { @@ -371,7 +371,7 @@ void max_concurrency_dependent() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/intel-fpga-mem-builtin.cpp b/clang/test/SemaSYCL/intel-fpga-mem-builtin.cpp index fa3709b4eccb5..bfc3d803ec44f 100644 --- a/clang/test/SemaSYCL/intel-fpga-mem-builtin.cpp +++ b/clang/test/SemaSYCL/intel-fpga-mem-builtin.cpp @@ -56,7 +56,7 @@ void foo(float *A, int *B, State *C) { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } int main() { diff --git a/clang/test/SemaSYCL/intel-fpga-reg.cpp b/clang/test/SemaSYCL/intel-fpga-reg.cpp index 5d09258d39170..213038b44a81d 100644 --- a/clang/test/SemaSYCL/intel-fpga-reg.cpp +++ b/clang/test/SemaSYCL/intel-fpga-reg.cpp @@ -57,7 +57,7 @@ void foo() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } int main() { diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim.cpp index d391924773f1e..c039787b6d8cf 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim.cpp @@ -1,15 +1,16 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s -// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -verify %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-host -Wno-sycl-2017-compat -fsyntax-only -verify %s #ifndef __SYCL_DEVICE_ONLY__ struct FuncObj { [[intelfpga::max_global_work_dim(1)]] // expected-no-diagnostics - void operator()() {} + void + operator()() const {} }; template -void kernel(Func kernelFunc) { +void kernel(const Func &kernelFunc) { kernelFunc(); } @@ -23,35 +24,35 @@ void foo() { [[intelfpga::max_global_work_dim(2)]] void func_do_not_ignore() {} struct FuncObj { - [[intelfpga::max_global_work_dim(1)]] - void operator()() {} + [[intelfpga::max_global_work_dim(1)]] void operator()() const {} }; struct TRIFuncObjGood1 { [[intelfpga::max_global_work_dim(0)]] [[intelfpga::max_work_group_size(1, 1, 1)]] - [[cl::reqd_work_group_size(1, 1, 1)]] - void operator()() {} + [[cl::reqd_work_group_size(1, 1, 1)]] void + operator()() const {} }; struct TRIFuncObjGood2 { [[intelfpga::max_global_work_dim(3)]] [[intelfpga::max_work_group_size(8, 1, 1)]] - [[cl::reqd_work_group_size(4, 1, 1)]] - void operator()() {} + [[cl::reqd_work_group_size(4, 1, 1)]] void + operator()() const {} }; #ifdef TRIGGER_ERROR struct TRIFuncObjBad { [[intelfpga::max_global_work_dim(0)]] [[intelfpga::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}} - [[cl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}} - void operator()() {} + [[cl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}} + void + operator()() const {} }; #endif // TRIGGER_ERROR template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp index f5b2a67599e8b..210944f6b2fdd 100644 --- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp @@ -1,15 +1,16 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s -// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -verify %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -Wno-sycl-2017-compat -verify %s #ifndef __SYCL_DEVICE_ONLY__ struct FuncObj { [[intelfpga::max_work_group_size(1, 1, 1)]] // expected-no-diagnostics - void operator()() {} + void + operator()() const {} }; template -void kernel(Func kernelFunc) { +void kernel(const Func &kernelFunc) { kernelFunc(); } @@ -23,20 +24,20 @@ void foo() { [[intelfpga::max_work_group_size(2, 2, 2)]] void func_do_not_ignore() {} struct FuncObj { - [[intelfpga::max_work_group_size(4, 4, 4)]] - void operator()() {} + [[intelfpga::max_work_group_size(4, 4, 4)]] void operator()() const {} }; #ifdef TRIGGER_ERROR struct DAFuncObj { [[intelfpga::max_work_group_size(4, 4, 4)]] [[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} - void operator()() {} + void + operator()() const {} }; #endif // TRIGGER_ERROR template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp index 361b3d083c0ef..09fe0934029b2 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size.cpp @@ -1,16 +1,16 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s -// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -fsyntax-only -verify -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-host -Wno-sycl-2017-compat -fsyntax-only -verify %s #ifndef __SYCL_DEVICE_ONLY__ // expected-no-diagnostics class Functor { public: - [[intel::reqd_work_group_size(4)]] void operator()() {} + [[intel::reqd_work_group_size(4)]] void operator()() const {} }; template -void kernel(Func kernelFunc) { +void kernel(const Func &kernelFunc) { kernelFunc(); } @@ -32,50 +32,50 @@ void bar() { #ifdef TRIGGER_ERROR class Functor32 { public: - [[cl::reqd_work_group_size(32)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} + [[cl::reqd_work_group_size(32)]] void operator()() const {} // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} }; class Functor33 { public: - [[intel::reqd_work_group_size(32, -4)]] void operator()() {} // expected-error {{'reqd_work_group_size' attribute requires a non-negative integral compile time constant expression}} + [[intel::reqd_work_group_size(32, -4)]] void operator()() const {} // expected-error {{'reqd_work_group_size' attribute requires a non-negative integral compile time constant expression}} }; #endif // TRIGGER_ERROR class Functor16 { public: - [[intel::reqd_work_group_size(16)]] void operator()() {} + [[intel::reqd_work_group_size(16)]] void operator()() const {} }; class Functor64 { public: - [[intel::reqd_work_group_size(64, 64)]] void operator()() {} + [[intel::reqd_work_group_size(64, 64)]] void operator()() const {} }; class Functor16x16x16 { public: - [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() {} + [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {} }; class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} public: - [[intel::reqd_work_group_size(8)]] void operator()() { // expected-note {{conflicting attribute is here}} + [[intel::reqd_work_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} f4x1x1(); } }; class Functor { public: - void operator()() { + void operator()() const { f4x1x1(); } }; class FunctorAttr { public: - __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() {} + __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() const {} }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp index b998b261f4320..c878cec15be3b 100644 --- a/clang/test/SemaSYCL/intel-restrict.cpp +++ b/clang/test/SemaSYCL/intel-restrict.cpp @@ -1,15 +1,14 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DCHECKDIAG -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 -DCHECKDIAG -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 | FileCheck %s [[intel::kernel_args_restrict]] void func_do_not_ignore() {} struct FuncObj { - [[intel::kernel_args_restrict]] - void operator()() {} + [[intel::kernel_args_restrict]] void operator()() const {} }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); #ifdef CHECKDIAG [[intel::kernel_args_restrict]] int invalid = 42; // expected-error{{'kernel_args_restrict' attribute only applies to functions}} diff --git a/clang/test/SemaSYCL/kernel-function-type.cpp b/clang/test/SemaSYCL/kernel-function-type.cpp index 7b33ec5ffd07a..86b0399d79ed9 100644 --- a/clang/test/SemaSYCL/kernel-function-type.cpp +++ b/clang/test/SemaSYCL/kernel-function-type.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s // expected-no-diagnostics // The kernel_single_task call is emitted as an OpenCL kernel function. The call @@ -29,7 +29,7 @@ void foo(j e, d k) { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp index 22a9f96acc50a..8ed02e06f7473 100644 --- a/clang/test/SemaSYCL/kernelname-enum.cpp +++ b/clang/test/SemaSYCL/kernelname-enum.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -Wno-sycl-2017-compat -verify %s #include "sycl.hpp" @@ -26,26 +26,26 @@ enum class scoped_enum_no_type_set { template class dummy_functor_1 { public: - void operator()() {} + void operator()() const {} }; // expected-error@+2 {{kernel name is invalid. Unscoped enum requires fixed underlying type}} template class dummy_functor_2 { public: - void operator()() {} + void operator()() const {} }; template class dummy_functor_3 { public: - void operator()() {} + void operator()() const {} }; template class dummy_functor_4 { public: - void operator()() {} + void operator()() const {} }; int main() { diff --git a/clang/test/SemaSYCL/lambda_implicit_capture_this.cpp b/clang/test/SemaSYCL/lambda_implicit_capture_this.cpp index 98212d3ca6448..3868601f353f8 100644 --- a/clang/test/SemaSYCL/lambda_implicit_capture_this.cpp +++ b/clang/test/SemaSYCL/lambda_implicit_capture_this.cpp @@ -1,7 +1,7 @@ -// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/loop_unroll.cpp b/clang/test/SemaSYCL/loop_unroll.cpp index 2e7e205532fd1..13ecae08d7430 100644 --- a/clang/test/SemaSYCL/loop_unroll.cpp +++ b/clang/test/SemaSYCL/loop_unroll.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -pedantic %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -pedantic %s template void bar() { @@ -60,7 +60,7 @@ void foo() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/mangle-kernel.cpp b/clang/test/SemaSYCL/mangle-kernel.cpp index 8cc2950019b45..7b515dd508309 100644 --- a/clang/test/SemaSYCL/mangle-kernel.cpp +++ b/clang/test/SemaSYCL/mangle-kernel.cpp @@ -5,7 +5,7 @@ #include template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/markfunction-astconsumer.cpp b/clang/test/SemaSYCL/markfunction-astconsumer.cpp index c4883f64a8c36..8315c890fc35e 100644 --- a/clang/test/SemaSYCL/markfunction-astconsumer.cpp +++ b/clang/test/SemaSYCL/markfunction-astconsumer.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -verify -fsyntax-only -std=c++17 %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s void bar(); template @@ -7,7 +7,7 @@ void usage(T func) { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/no-vtables.cpp b/clang/test/SemaSYCL/no-vtables.cpp index 1891ed278d444..68f9953bf6784 100644 --- a/clang/test/SemaSYCL/no-vtables.cpp +++ b/clang/test/SemaSYCL/no-vtables.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -verify -fsyntax-only -emit-llvm-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -verify -Wno-sycl-2017-compat -fsyntax-only -emit-llvm-only %s // expected-no-diagnostics // Should never fail, since the type is never used in kernel code. @@ -17,9 +17,8 @@ void always_uses() { void usage() { } - template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } int main() { diff --git a/clang/test/SemaSYCL/no-vtables2.cpp b/clang/test/SemaSYCL/no-vtables2.cpp index e2468121291f4..92bde4db57b33 100644 --- a/clang/test/SemaSYCL/no-vtables2.cpp +++ b/clang/test/SemaSYCL/no-vtables2.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only %s struct Base { virtual void f() const {} @@ -28,9 +28,8 @@ Inherit usage() { usage_child(); } - template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } int main() { diff --git a/clang/test/SemaSYCL/non-std-layout-param.cpp b/clang/test/SemaSYCL/non-std-layout-param.cpp index 0a8ffbf798f94..0698b3b7573a3 100644 --- a/clang/test/SemaSYCL/non-std-layout-param.cpp +++ b/clang/test/SemaSYCL/non-std-layout-param.cpp @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-std-layout-kernel-params -verify -fsyntax-only %s -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-std-layout-kernel-params -verify -Wno-sycl-2017-compat -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -fsyntax-only %s // This test checks if compiler reports compilation error on an attempt to pass // non-standard layout struct object as SYCL kernel parameter. @@ -15,11 +15,10 @@ struct C : public Base { }; template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } - void test() { C C0; C0.Y=0; @@ -30,7 +29,7 @@ void test() { } struct Kernel { - void operator()() { + void operator()() const { (void) c1; (void) c2; (void) p; diff --git a/clang/test/SemaSYCL/non-trivially-copyable-kernel-param.cpp b/clang/test/SemaSYCL/non-trivially-copyable-kernel-param.cpp index 4453e5273f9a7..00d1ec0bc0418 100644 --- a/clang/test/SemaSYCL/non-trivially-copyable-kernel-param.cpp +++ b/clang/test/SemaSYCL/non-trivially-copyable-kernel-param.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -Wno-sycl-2017-compat -fsyntax-only %s // This test checks if compiler reports compilation error on an attempt to pass // a struct with non-trivially copyable type as SYCL kernel parameter. @@ -22,7 +22,7 @@ struct D { }; template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/num_simd_work_items.cpp b/clang/test/SemaSYCL/num_simd_work_items.cpp index 638539917a283..d6528ec4e665d 100644 --- a/clang/test/SemaSYCL/num_simd_work_items.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items.cpp @@ -1,15 +1,16 @@ -// RUN: %clang_cc1 %s -fsycl -fsycl-is-device -triple spir64 -fsyntax-only -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsycl -fsycl-is-device -triple spir64 -fsyntax-only -ast-dump | FileCheck %s -// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -verify %s +// RUN: %clang_cc1 %s -fsycl -fsycl-is-device -triple spir64 -fsyntax-only -Wno-sycl-2017-compat -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsycl -fsycl-is-device -triple spir64 -fsyntax-only -Wno-sycl-2017-compat -ast-dump | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-host -fsyntax-only -Wno-sycl-2017-compat -verify %s #ifndef __SYCL_DEVICE_ONLY__ struct FuncObj { [[intelfpga::num_simd_work_items(42)]] // expected-no-diagnostics - void operator()() {} + void + operator()() const {} }; template -void kernel(Func kernelFunc) { +void kernel(const Func &kernelFunc) { kernelFunc(); } @@ -23,12 +24,11 @@ void foo() { [[intelfpga::num_simd_work_items(2)]] void func_do_not_ignore() {} struct FuncObj { - [[intelfpga::num_simd_work_items(42)]] - void operator()() {} + [[intelfpga::num_simd_work_items(42)]] void operator()() const {} }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/prohibit-thread-local.cpp b/clang/test/SemaSYCL/prohibit-thread-local.cpp index e07353e862479..1cdfb4c5ea3b6 100644 --- a/clang/test/SemaSYCL/prohibit-thread-local.cpp +++ b/clang/test/SemaSYCL/prohibit-thread-local.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify -Wno-sycl-2017-compat -fsyntax-only %s thread_local const int prohobit_ns_scope = 0; thread_local int prohobit_ns_scope2 = 0; @@ -41,7 +41,7 @@ template __attribute__((sycl_kernel)) // expected-note@+2 2{{called by}} void -kernel_single_task(Func kernelFunc) { kernelFunc(); } +kernel_single_task(const Func &kernelFunc) { kernelFunc(); } int main() { // expected-note@+1 2{{called by}} diff --git a/clang/test/SemaSYCL/reference-kernel-param.cpp b/clang/test/SemaSYCL/reference-kernel-param.cpp index 3bd548e23a6fb..06809a8fcd1c3 100644 --- a/clang/test/SemaSYCL/reference-kernel-param.cpp +++ b/clang/test/SemaSYCL/reference-kernel-param.cpp @@ -1,10 +1,10 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -Wno-sycl-2017-compat -fsyntax-only %s // This test checks if compiler reports compilation error on an attempt to pass // a reference as SYCL kernel parameter. template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp index ffa4b176207a0..56fca9002d436 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s [[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}} // expected-note@-1 {{conflicting attribute is here}} @@ -9,30 +9,30 @@ class Functor16 { public: // expected-warning@+2 {{attribute 'intel_reqd_sub_group_size' is deprecated}} // expected-note@+1 {{did you mean to use 'intel::reqd_sub_group_size' instead?}} - [[cl::intel_reqd_sub_group_size(16)]] void operator()() {} + [[cl::intel_reqd_sub_group_size(16)]] void operator()() const {} }; class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} public: - [[intel::reqd_sub_group_size(8)]] void operator()() { // expected-note {{conflicting attribute is here}} + [[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} foo(); } }; class Functor4 { public: - [[intel::reqd_sub_group_size(12)]] void operator()() {} + [[intel::reqd_sub_group_size(12)]] void operator()() const {} }; class Functor { public: - void operator()() { + void operator()() const { foo(); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp index 7175322fb43c4..65bb0b13474d7 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s [[cl::reqd_work_group_size(4, 1, 1)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} // expected-note@-1 {{conflicting attribute is here}} @@ -13,7 +13,7 @@ class Functor16 { public: - [[cl::reqd_work_group_size(16, 1, 1)]] [[cl::reqd_work_group_size(16, 1, 1)]] void operator()() {} + [[cl::reqd_work_group_size(16, 1, 1)]] [[cl::reqd_work_group_size(16, 1, 1)]] void operator()() const {} }; #ifdef TRIGGER_ERROR @@ -21,35 +21,35 @@ class Functor32 { public: //expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different parameters}} // expected-error@+1{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} - [[cl::reqd_work_group_size(32, 1, 1)]] [[cl::reqd_work_group_size(1, 1, 32)]] void operator()() {} + [[cl::reqd_work_group_size(32, 1, 1)]] [[cl::reqd_work_group_size(1, 1, 32)]] void operator()() const {} }; #endif class Functor16x16x16 { public: - [[cl::reqd_work_group_size(16, 16, 16)]] void operator()() {} + [[cl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} }; class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} public: - [[cl::reqd_work_group_size(1, 1, 8)]] void operator()() { // expected-note {{conflicting attribute is here}} + [[cl::reqd_work_group_size(1, 1, 8)]] void operator()() const { // expected-note {{conflicting attribute is here}} f4x1x1(); } }; class Functor { public: - void operator()() { + void operator()() const { f4x1x1(); } }; class FunctorAttr { public: - __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() {} + __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() const {} }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/restrict-recursion.cpp b/clang/test/SemaSYCL/restrict-recursion.cpp index 3547e2f793e19..3e29271561fe2 100644 --- a/clang/test/SemaSYCL/restrict-recursion.cpp +++ b/clang/test/SemaSYCL/restrict-recursion.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -verify -fsyntax-only -std=c++17 %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s // This recursive function is not called from sycl kernel, // so it should not be diagnosed. @@ -57,7 +57,7 @@ bool isa_B(void) { } template -__attribute__((sycl_kernel)) void kernel(L l) { +__attribute__((sycl_kernel)) void kernel(const L &l) { l(); } @@ -85,13 +85,13 @@ int addInt(int n, int m) { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } template - // expected-note@+1 2{{function implemented using recursion declared here}} -__attribute__((sycl_kernel)) void kernel_single_task2(Func kernelFunc) { +// expected-note@+1 2{{function implemented using recursion declared here}} +__attribute__((sycl_kernel)) void kernel_single_task2(const Func &kernelFunc) { kernelFunc(); // expected-error@+1 2{{SYCL kernel cannot call a recursive function}} kernel_single_task2(kernelFunc); diff --git a/clang/test/SemaSYCL/restrict-recursion2.cpp b/clang/test/SemaSYCL/restrict-recursion2.cpp index 51688fea7b1b6..b9eaffa4149b1 100644 --- a/clang/test/SemaSYCL/restrict-recursion2.cpp +++ b/clang/test/SemaSYCL/restrict-recursion2.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -verify -fsyntax-only -std=c++17 %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s // This recursive function is not called from sycl kernel, // so it should not be diagnosed. @@ -71,7 +71,7 @@ int addInt(int n, int m) { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/restrict-recursion3.cpp b/clang/test/SemaSYCL/restrict-recursion3.cpp index 1a4a671904169..f5289b3142e2b 100644 --- a/clang/test/SemaSYCL/restrict-recursion3.cpp +++ b/clang/test/SemaSYCL/restrict-recursion3.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-error=sycl-strict -verify -fsyntax-only -std=c++17 %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-sycl-2017-compat -Wno-error=sycl-strict -verify -fsyntax-only -std=c++17 %s // This recursive function is not called from sycl kernel, // so it should not be diagnosed. @@ -32,7 +32,7 @@ int addInt(int n, int m) { template // expected-note@+1 2{{function implemented using recursion declared here}} -__attribute__((sycl_kernel)) void kernel_single_task2(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task2(const Func &kernelFunc) { // expected-note@+1 {{called by 'kernel_single_task2}} kernelFunc(); // expected-warning@+1 2{{SYCL kernel cannot call a recursive function}} diff --git a/clang/test/SemaSYCL/restrict-recursion4.cpp b/clang/test/SemaSYCL/restrict-recursion4.cpp index c755f0eaec63a..dde093792cf26 100644 --- a/clang/test/SemaSYCL/restrict-recursion4.cpp +++ b/clang/test/SemaSYCL/restrict-recursion4.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-error=sycl-strict -verify -fsyntax-only -std=c++17 %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-sycl-2017-compat -Wno-error=sycl-strict -verify -fsyntax-only -std=c++17 %s // This recursive function is not called from sycl kernel, // so it should not be diagnosed. @@ -34,7 +34,7 @@ int addInt(int n, int m) { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // expected-note@+1 {{called by 'kernel_single_task}} kernelFunc(); } diff --git a/clang/test/SemaSYCL/sampler.cpp b/clang/test/SemaSYCL/sampler.cpp index 2a3171ce06f08..a1fb423b109ca 100644 --- a/clang/test/SemaSYCL/sampler.cpp +++ b/clang/test/SemaSYCL/sampler.cpp @@ -3,7 +3,7 @@ #include template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp b/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp index cbfba6e4be32d..54a6f46290268 100644 --- a/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp +++ b/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp @@ -5,7 +5,7 @@ #include template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/sycl-callstack.cpp b/clang/test/SemaSYCL/sycl-callstack.cpp index 15da9d68255f5..8c8a53753743d 100644 --- a/clang/test/SemaSYCL/sycl-callstack.cpp +++ b/clang/test/SemaSYCL/sycl-callstack.cpp @@ -1,10 +1,9 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -verify -fsyntax-only -std=c++17 %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s template -__attribute__((sycl_kernel)) -void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // expected-note@+1 {{called by 'kernel_single_task}} - kernelFunc(); + kernelFunc(); } void foo() { diff --git a/clang/test/SemaSYCL/sycl-cconv.cpp b/clang/test/SemaSYCL/sycl-cconv.cpp index f905417027c4a..a42d4e7622aaa 100644 --- a/clang/test/SemaSYCL/sycl-cconv.cpp +++ b/clang/test/SemaSYCL/sycl-cconv.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-windows-sycldevice -aux-triple x86_64-pc-windows-msvc -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-windows-sycldevice -aux-triple x86_64-pc-windows-msvc -fsyntax-only -Wno-sycl-2017-compat -verify %s // expected-no-warning@+1 __inline __cdecl int printf(char const* const _Format, ...) { return 0; } @@ -13,7 +13,7 @@ void bar() { template // expected-no-warning@+1 -__cdecl __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__cdecl __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // expected-error@+1{{SYCL kernel cannot call a variadic function}} printf("cannot call from here\n"); // expected-no-error@+1 diff --git a/clang/test/SemaSYCL/sycl-device-const-static.cpp b/clang/test/SemaSYCL/sycl-device-const-static.cpp index 0e90057c73d60..02a97810b7e5e 100644 --- a/clang/test/SemaSYCL/sycl-device-const-static.cpp +++ b/clang/test/SemaSYCL/sycl-device-const-static.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -Wno-sycl-2017-compat -fsyntax-only %s struct Base {}; struct S { @@ -37,7 +37,7 @@ void usage() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // expected-error@+1{{SYCL kernel cannot use a non-const static data variable}} static int z; // expected-note-re@+3{{called by 'kernel_single_task}} diff --git a/clang/test/SemaSYCL/sycl-device-static-restrict.cpp b/clang/test/SemaSYCL/sycl-device-static-restrict.cpp index 03132b4530b88..b3517770e5290 100644 --- a/clang/test/SemaSYCL/sycl-device-static-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-device-static-restrict.cpp @@ -1,8 +1,8 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -Wno-sycl-2017-compat -fsyntax-only %s const int glob1 = 1; int glob2 = 2; template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // expected-note-re@+1{{called by 'kernel_single_task}} kernelFunc(); } diff --git a/clang/test/SemaSYCL/sycl-dllimport-dllexport.cpp b/clang/test/SemaSYCL/sycl-dllimport-dllexport.cpp index b2c3fbf3aa1ff..e03d749a09db6 100644 --- a/clang/test/SemaSYCL/sycl-dllimport-dllexport.cpp +++ b/clang/test/SemaSYCL/sycl-dllimport-dllexport.cpp @@ -1,19 +1,19 @@ // RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fms-extensions \ // RUN: -aux-triple x86_64-unknown-linux-gnu -fsycl -fsycl-is-device \ -// RUN: -fsyntax-only -DWARNCHECK %s -o /dev/null 2>&1 | FileCheck %s +// RUN: -fsyntax-only -Wno-sycl-2017-compat -DWARNCHECK %s -o /dev/null 2>&1 | FileCheck %s // check random triple aux-triple with sycl-device -// RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -fsyntax-only \ +// RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -Wno-sycl-2017-compat -fsyntax-only \ // RUN: -fms-extensions -DWARNCHECK %s -o /dev/null 2>&1 | FileCheck %s // check without -aux-triple but sycl-device // RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -fsycl \ // RUN: -fsycl-is-device -aux-triple x86_64-pc-windows-msvc -fms-extensions \ -// RUN: -fsyntax-only -DWARNCHECK %s -o /dev/null 2>&1 | \ +// RUN: -fsyntax-only -Wno-sycl-2017-compat -DWARNCHECK %s -o /dev/null 2>&1 | \ // RUN: FileCheck %s --check-prefixes CHECKALL // check -aux-tripe without sycl-device -// RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -fsyntax-only \ +// RUN: %clang_cc1 -triple spir64-unknown-windows-sycldevice -Wno-sycl-2017-compat -fsyntax-only \ // RUN: -aux-triple x86_64-pc-windows-msvc -fsycl -fsycl-is-device \ // RUN: -fms-extensions -verify %s // check error message when dllimport function gets called in sycl-kernel code @@ -51,7 +51,7 @@ int foobar() // expected-warning {{'foobar' redeclared without 'dllimport' attr } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/sycl-esimd.cpp b/clang/test/SemaSYCL/sycl-esimd.cpp index f235394b9971c..68ef557818ddc 100644 --- a/clang/test/SemaSYCL/sycl-esimd.cpp +++ b/clang/test/SemaSYCL/sycl-esimd.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-explicit-simd -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-explicit-simd -fsyntax-only -Wno-sycl-2017-compat -verify %s // ----------- Negative tests @@ -12,7 +12,7 @@ bar() {} // -- ESIMD kernel can't call functions with required subgroup size != 1 template -void kernel0(F f) __attribute__((sycl_kernel)) { +void kernel0(const F &f) __attribute__((sycl_kernel)) { f(); } @@ -27,7 +27,7 @@ void test0() { // -- Usual kernel can't call ESIMD function template -void kernel1(F f) __attribute__((sycl_kernel)) { +void kernel1(const F &f) __attribute__((sycl_kernel)) { f(); } @@ -43,7 +43,7 @@ void test1() { // -- Kernel-function call, both have the attribute, lambda kernel. template -void kernel2(F f) __attribute__((sycl_kernel)) { +void kernel2(const F &f) __attribute__((sycl_kernel)) { f(); } @@ -64,12 +64,12 @@ class A { // -- Functor object kernel. template -void kernel3(F f) __attribute__((sycl_kernel)) { +void kernel3(const F &f) __attribute__((sycl_kernel)) { f(); } struct Kernel3 { - void operator()() __attribute__((sycl_explicit_simd)) {} + void operator()() const __attribute__((sycl_explicit_simd)) {} }; void bar3() { diff --git a/clang/test/SemaSYCL/sycl-fptr-lambda.cpp b/clang/test/SemaSYCL/sycl-fptr-lambda.cpp index ba44c7286ec2d..5ec68265a2ee9 100644 --- a/clang/test/SemaSYCL/sycl-fptr-lambda.cpp +++ b/clang/test/SemaSYCL/sycl-fptr-lambda.cpp @@ -1,8 +1,8 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -std=c++14 -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -std=c++14 -verify -Wno-sycl-2017-compat -fsyntax-only %s // expected-no-diagnostics template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/sycl-pseudo-dtor.cpp b/clang/test/SemaSYCL/sycl-pseudo-dtor.cpp index b4700cadd3f2f..cc839a18446ed 100644 --- a/clang/test/SemaSYCL/sycl-pseudo-dtor.cpp +++ b/clang/test/SemaSYCL/sycl-pseudo-dtor.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -Wno-sycl-2017-compat -fsyntax-only %s template struct functor_wrapper{ @@ -13,7 +13,7 @@ struct S { virtual void foo(); }; struct T { virtual ~T(); }; template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // expected-no-note@+1 using DATA_I = int; using DATA_S = S; diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index 00b215d7d459b..2c963f95ff76d 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -1,6 +1,6 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -Wno-return-type -verify -fsyntax-only -std=c++17 %s -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -fno-sycl-allow-func-ptr -Wno-return-type -verify -fsyntax-only -std=c++17 %s -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -DALLOW_FP=1 -fsycl-allow-func-ptr -Wno-return-type -verify -fsyntax-only -std=c++17 %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -fno-sycl-allow-func-ptr -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -triple spir64 -DALLOW_FP=1 -fsycl-allow-func-ptr -Wno-return-type -verify -Wno-sycl-2017-compat -fsyntax-only -std=c++17 %s namespace std { class type_info; @@ -86,7 +86,7 @@ bool isa_B(A *a) { } template -__attribute__((sycl_kernel)) void kernel1(L l) { +__attribute__((sycl_kernel)) void kernel1(const L &l) { l(); //#rtti_kernel // expected-note 2{{called by 'kernel1 -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); //#call_kernelFunc // expected-note 3{{called by 'kernel_single_task -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { kernelFunc(); //expected-note 2+ {{called by 'kernel_single_task}} } diff --git a/clang/test/SemaSYCL/tls_error.cpp b/clang/test/SemaSYCL/tls_error.cpp index eb89616105955..bacecf7827803 100644 --- a/clang/test/SemaSYCL/tls_error.cpp +++ b/clang/test/SemaSYCL/tls_error.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64 -verify -Wno-sycl-2017-compat -fsyntax-only %s extern __thread void* __once_callable; // expected-no-error extern __thread void (*__once_call)(); // expected-no-error @@ -14,7 +14,7 @@ void usage() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { //expected-note@+1{{called by}} kernelFunc(); } diff --git a/clang/test/SemaSYCL/unevaluated-function.cpp b/clang/test/SemaSYCL/unevaluated-function.cpp index 6e4f3b0fa3f57..49ffe954ab7a9 100644 --- a/clang/test/SemaSYCL/unevaluated-function.cpp +++ b/clang/test/SemaSYCL/unevaluated-function.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -verify -fsyntax-only %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -verify -Wno-sycl-2017-compat -fsyntax-only %s // Check that a function used in an unevaluated context is not subject // to delayed device diagnostics. @@ -24,7 +24,7 @@ bool foo3() { } template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { // expected-note@+1 1{{called by}} kernelFunc(); } diff --git a/clang/test/SemaSYCL/unnamed-kernel.cpp b/clang/test/SemaSYCL/unnamed-kernel.cpp index 33e5656ac8220..45954f325b4d6 100644 --- a/clang/test/SemaSYCL/unnamed-kernel.cpp +++ b/clang/test/SemaSYCL/unnamed-kernel.cpp @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -Wno-sycl-2017-compat -verify %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -Wno-sycl-2017-compat -verify %s #include #ifdef __SYCL_UNNAMED_LAMBDA__ diff --git a/clang/test/SemaSYCL/unsupported_math.cpp b/clang/test/SemaSYCL/unsupported_math.cpp index e93bf61283588..3c0de837dcd77 100644 --- a/clang/test/SemaSYCL/unsupported_math.cpp +++ b/clang/test/SemaSYCL/unsupported_math.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s extern "C" float sinf(float); extern "C" float cosf(float); extern "C" float logf(float); @@ -6,7 +6,7 @@ extern "C" double sin(double); extern "C" double cos(double); extern "C" double log(double); template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/SemaSYCL/variadic-func-call.cpp b/clang/test/SemaSYCL/variadic-func-call.cpp index 896ba16e46b9e..bdb7ace1690b6 100644 --- a/clang/test/SemaSYCL/variadic-func-call.cpp +++ b/clang/test/SemaSYCL/variadic-func-call.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown -fsyntax-only -Wno-sycl-2017-compat -verify %s void variadic(int, ...) {} namespace NS { @@ -18,7 +18,7 @@ void foo() { void overloaded(int, int) {} void overloaded(int, ...) {} template -__attribute__((sycl_kernel)) void task(Func KF) { +__attribute__((sycl_kernel)) void task(const Func &KF) { KF(); // expected-note 2 {{called by 'task}} } diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp index 9f9a3196024c1..c6ebc3c48730e 100644 --- a/clang/test/SemaSYCL/wrapped-accessor.cpp +++ b/clang/test/SemaSYCL/wrapped-accessor.cpp @@ -9,7 +9,7 @@ template struct AccWrapper { Acc accessor; }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 83cc08f88cbbf..03b3d5eb9d74f 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -33,6 +33,11 @@ #include #include +// SYCL_LANGUAGE_VERSION is 4 digit year followed by 2 digit revision +#if !SYCL_LANGUAGE_VERSION || SYCL_LANGUAGE_VERSION < 202001 +#define __SYCL_NONCONST_FUNCTOR__ +#endif + template @@ -719,14 +724,24 @@ class __SYCL_EXPORT handler { // NOTE: the name of this function - "kernel_single_task" - is used by the // Front End to determine kernel invocation kind. template - __attribute__((sycl_kernel)) void kernel_single_task(KernelType KernelFunc) { + __attribute__((sycl_kernel)) void +#ifdef __SYCL_NONCONST_FUNCTOR__ + kernel_single_task(KernelType KernelFunc) { +#else + kernel_single_task(const KernelType &KernelFunc) { +#endif KernelFunc(); } // NOTE: the name of these functions - "kernel_parallel_for" - are used by the // Front End to determine kernel invocation kind. template - __attribute__((sycl_kernel)) void kernel_parallel_for(KernelType KernelFunc) { + __attribute__((sycl_kernel)) void +#ifdef __SYCL_NONCONST_FUNCTOR__ + kernel_parallel_for(KernelType KernelFunc) { +#else + kernel_parallel_for(const KernelType &KernelFunc) { +#endif KernelFunc( detail::Builder::getElement(static_cast(nullptr))); } @@ -735,7 +750,11 @@ class __SYCL_EXPORT handler { // used by the Front End to determine kernel invocation kind. template __attribute__((sycl_kernel)) void +#ifdef __SYCL_NONCONST_FUNCTOR__ kernel_parallel_for_work_group(KernelType KernelFunc) { +#else + kernel_parallel_for_work_group(const KernelType &KernelFunc) { +#endif KernelFunc( detail::Builder::getElement(static_cast(nullptr))); } @@ -793,7 +812,7 @@ class __SYCL_EXPORT handler { template struct ShouldEnableSetArg { static constexpr bool value = std::is_trivially_copyable>::value -#if CL_SYCL_LANGUAGE_VERSION && CL_SYCL_LANGUAGE_VERSION <= 121 +#if SYCL_LANGUAGE_VERSION && SYCL_LANGUAGE_VERSION <= 201707 && std::is_standard_layout>::value #endif || is_same_type::value // Sampler @@ -839,7 +858,11 @@ class __SYCL_EXPORT handler { /// /// \param KernelFunc is a SYCL kernel function. template +#ifdef __SYCL_NONCONST_FUNCTOR__ void single_task(KernelType KernelFunc) { +#else + void single_task(const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -856,17 +879,29 @@ class __SYCL_EXPORT handler { } template +#ifdef __SYCL_NONCONST_FUNCTOR__ void parallel_for(range<1> NumWorkItems, KernelType KernelFunc) { +#else + void parallel_for(range<1> NumWorkItems, const KernelType &KernelFunc) { +#endif parallel_for_lambda_impl(NumWorkItems, std::move(KernelFunc)); } template +#ifdef __SYCL_NONCONST_FUNCTOR__ void parallel_for(range<2> NumWorkItems, KernelType KernelFunc) { +#else + void parallel_for(range<2> NumWorkItems, const KernelType &KernelFunc) { +#endif parallel_for_lambda_impl(NumWorkItems, std::move(KernelFunc)); } template +#ifdef __SYCL_NONCONST_FUNCTOR__ void parallel_for(range<3> NumWorkItems, KernelType KernelFunc) { +#else + void parallel_for(range<3> NumWorkItems, const KernelType &KernelFunc) { +#endif parallel_for_lambda_impl(NumWorkItems, std::move(KernelFunc)); } @@ -913,6 +948,14 @@ class __SYCL_EXPORT handler { MCGType = detail::CG::CODEPLAY_HOST_TASK; } +// replace _KERNELFUNCPARAM(KernelFunc) with KernelType KernelFunc +// or const KernelType &KernelFunc +#ifdef __SYCL_NONCONST_FUNCTOR__ +#define _KERNELFUNCPARAM(a) KernelType a +#else +#define _KERNELFUNCPARAM(a) const KernelType &a +#endif + /// Defines and invokes a SYCL kernel function for the specified range and /// offset. /// @@ -929,7 +972,7 @@ class __SYCL_EXPORT handler { template void parallel_for(range NumWorkItems, id WorkItemOffset, - KernelType KernelFunc) { + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -961,7 +1004,8 @@ class __SYCL_EXPORT handler { /// \param KernelFunc is a SYCL kernel function. template - void parallel_for(nd_range ExecutionRange, KernelType KernelFunc) { + void parallel_for(nd_range ExecutionRange, + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -988,7 +1032,8 @@ class __SYCL_EXPORT handler { int Dims, typename Reduction> detail::enable_if_t - parallel_for(nd_range Range, Reduction Redu, KernelType KernelFunc) { + parallel_for(nd_range Range, Reduction Redu, + _KERNELFUNCPARAM(KernelFunc)) { intel::detail::reduCGFunc(*this, KernelFunc, Range, Redu, Redu.getUserAccessor()); } @@ -1001,7 +1046,8 @@ class __SYCL_EXPORT handler { int Dims, typename Reduction> detail::enable_if_t - parallel_for(nd_range Range, Reduction Redu, KernelType KernelFunc) { + parallel_for(nd_range Range, Reduction Redu, + _KERNELFUNCPARAM(KernelFunc)) { intel::detail::reduCGFunc(*this, KernelFunc, Range, Redu, Redu.getUSMPointer()); } @@ -1020,7 +1066,8 @@ class __SYCL_EXPORT handler { int Dims, typename Reduction> detail::enable_if_t - parallel_for(nd_range Range, Reduction Redu, KernelType KernelFunc) { + parallel_for(nd_range Range, Reduction Redu, + _KERNELFUNCPARAM(KernelFunc)) { shared_ptr_class QueueCopy = MQueue; auto RWAcc = Redu.getReadWriteScalarAcc(*this); intel::detail::reduCGFunc(*this, KernelFunc, Range, Redu, @@ -1056,7 +1103,8 @@ class __SYCL_EXPORT handler { template detail::enable_if_t - parallel_for(nd_range Range, Reduction Redu, KernelType KernelFunc) { + parallel_for(nd_range Range, Reduction Redu, + _KERNELFUNCPARAM(KernelFunc)) { // This parallel_for() is lowered to the following sequence: // 1) Call a kernel that a) call user's lambda function and b) performs // one iteration of reduction, storing the partial reductions/sums @@ -1131,7 +1179,7 @@ class __SYCL_EXPORT handler { template void parallel_for_work_group(range NumWorkGroups, - KernelType KernelFunc) { + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1164,7 +1212,7 @@ class __SYCL_EXPORT handler { int Dims> void parallel_for_work_group(range NumWorkGroups, range WorkGroupSize, - KernelType KernelFunc) { + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1266,7 +1314,7 @@ class __SYCL_EXPORT handler { /// \param KernelFunc is a lambda that is used if device, queue is bound to, /// is a host device. template - void single_task(kernel Kernel, KernelType KernelFunc) { + void single_task(kernel Kernel, _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1306,7 +1354,7 @@ class __SYCL_EXPORT handler { template void parallel_for(kernel Kernel, range NumWorkItems, - KernelType KernelFunc) { + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1341,7 +1389,7 @@ class __SYCL_EXPORT handler { template void parallel_for(kernel Kernel, range NumWorkItems, - id WorkItemOffset, KernelType KernelFunc) { + id WorkItemOffset, _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1378,7 +1426,7 @@ class __SYCL_EXPORT handler { template void parallel_for(kernel Kernel, nd_range NDRange, - KernelType KernelFunc) { + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1420,7 +1468,7 @@ class __SYCL_EXPORT handler { template void parallel_for_work_group(kernel Kernel, range NumWorkGroups, - KernelType KernelFunc) { + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1458,7 +1506,7 @@ class __SYCL_EXPORT handler { int Dims> void parallel_for_work_group(kernel Kernel, range NumWorkGroups, range WorkGroupSize, - KernelType KernelFunc) { + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1482,6 +1530,9 @@ class __SYCL_EXPORT handler { #endif // __SYCL_DEVICE_ONLY__ } + // Clean up KERNELFUNC macro. +#undef _KERNELFUNCPARAM + // Explicit copy operations API /// Copies the content of memory object accessed by Src into the memory diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index 5338567f5eda3..2c6ffdfdfed74 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -361,21 +361,44 @@ class __SYCL_EXPORT queue { return submit([=](handler &CGH) { CGH.prefetch(Ptr, Count); }); } + // having _TWO_ mid-param #ifdefs makes the functions very difficult to read. + // Here we simplify the &CodeLoc declaration to be _CODELOCPARAM(&CodeLoc) and + // _CODELOCARG(&CodeLoc) Similarly, the KernelFunc param is simplified to be + // _KERNELFUNCPARAM(KernelFunc) Once the queue kernel functions are defined, + // these macros are #undef immediately. + + // replace _CODELOCPARAM(&CodeLoc) with nothing + // or : , const detail::code_location &CodeLoc = + // detail::code_location::current() + // replace _CODELOCARG(&CodeLoc) with nothing + // or : const detail::code_location &CodeLoc = {} + +#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA +#define _CODELOCPARAM(a) \ + , const detail::code_location a = detail::code_location::current() + +#define _CODELOCARG(a) +#else +#define _CODELOCPARAM(a) + +#define _CODELOCARG(a) const detail::code_location a = {} +#endif +// replace _KERNELFUNCPARAM(KernelFunc) with KernelType KernelFunc +// or const KernelType &KernelFunc +#ifdef __SYCL_NONCONST_FUNCTOR__ +#define _KERNELFUNCPARAM(a) KernelType a +#else +#define _KERNELFUNCPARAM(a) const KernelType &a +#endif + /// single_task version with a kernel represented as a lambda. /// /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event single_task( - KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event single_task(_KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); + return submit( [&](handler &CGH) { CGH.template single_task(KernelFunc); @@ -389,16 +412,9 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event single_task( - event DepEvent, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event single_task(event DepEvent, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvent); @@ -414,16 +430,9 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event single_task( - const vector_class &DepEvents, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event single_task(const vector_class &DepEvents, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvents); @@ -439,16 +448,9 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event parallel_for( - range<1> NumWorkItems, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range<1> NumWorkItems, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, KernelFunc, CodeLoc); } @@ -459,16 +461,9 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event parallel_for( - range<2> NumWorkItems, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range<2> NumWorkItems, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, KernelFunc, CodeLoc); } @@ -479,16 +474,9 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event parallel_for( - range<3> NumWorkItems, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range<3> NumWorkItems, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, KernelFunc, CodeLoc); } @@ -500,16 +488,9 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event parallel_for( - range<1> NumWorkItems, event DepEvent, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range<1> NumWorkItems, event DepEvent, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvent, KernelFunc, CodeLoc); } @@ -522,16 +503,9 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event parallel_for( - range<2> NumWorkItems, event DepEvent, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range<2> NumWorkItems, event DepEvent, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvent, KernelFunc, CodeLoc); } @@ -544,16 +518,9 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event parallel_for( - range<3> NumWorkItems, event DepEvent, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range<3> NumWorkItems, event DepEvent, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvent, KernelFunc, CodeLoc); } @@ -567,17 +534,10 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event parallel_for( - range<1> NumWorkItems, const vector_class &DepEvents, - KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range<1> NumWorkItems, + const vector_class &DepEvents, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvents, KernelFunc, CodeLoc); } @@ -591,17 +551,10 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event parallel_for( - range<2> NumWorkItems, const vector_class &DepEvents, - KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range<2> NumWorkItems, + const vector_class &DepEvents, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvents, KernelFunc, CodeLoc); } @@ -615,17 +568,10 @@ class __SYCL_EXPORT queue { /// \param KernelFunc is the Kernel functor or lambda /// \param CodeLoc contains the code location of user code template - event parallel_for( - range<3> NumWorkItems, const vector_class &DepEvents, - KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range<3> NumWorkItems, + const vector_class &DepEvents, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvents, KernelFunc, CodeLoc); } @@ -639,16 +585,9 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template - event parallel_for( - range NumWorkItems, id WorkItemOffset, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range NumWorkItems, id WorkItemOffset, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.template parallel_for( @@ -667,17 +606,10 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template - event parallel_for( - range NumWorkItems, id WorkItemOffset, event DepEvent, - KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range NumWorkItems, id WorkItemOffset, + event DepEvent, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvent); @@ -698,17 +630,10 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template - event parallel_for( - range NumWorkItems, id WorkItemOffset, - const vector_class &DepEvents, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(range NumWorkItems, id WorkItemOffset, + const vector_class &DepEvents, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvents); @@ -727,16 +652,9 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template - event parallel_for( - nd_range ExecutionRange, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(nd_range ExecutionRange, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.template parallel_for(ExecutionRange, @@ -755,16 +673,9 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template - event parallel_for( - nd_range ExecutionRange, event DepEvent, KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(nd_range ExecutionRange, event DepEvent, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvent); @@ -785,17 +696,10 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template - event parallel_for( - nd_range ExecutionRange, const vector_class &DepEvents, - KernelType KernelFunc -#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA - , - const detail::code_location &CodeLoc = detail::code_location::current() -#endif - ) { -#ifdef DISABLE_SYCL_INSTRUMENTATION_METADATA - const detail::code_location &CodeLoc = {}; -#endif + event parallel_for(nd_range ExecutionRange, + const vector_class &DepEvents, + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvents); @@ -805,6 +709,11 @@ class __SYCL_EXPORT queue { CodeLoc); } +// Clean up CODELOC and KERNELFUNC macros. +#undef _CODELOCPARAM +#undef _CODELOCARG +#undef _KERNELFUNCPARAM + /// Returns whether the queue is in order or OoO /// /// Equivalent to has_property() diff --git a/sycl/test/basic_tests/macros.cpp b/sycl/test/basic_tests/macros.cpp index 2be95b94964a6..726f070406e2f 100644 --- a/sycl/test/basic_tests/macros.cpp +++ b/sycl/test/basic_tests/macros.cpp @@ -11,8 +11,7 @@ #include int main() { - std::cout << "SYCL language version: " << CL_SYCL_LANGUAGE_VERSION - << std::endl; + std::cout << "SYCL language version: " << SYCL_LANGUAGE_VERSION << std::endl; std::cout << "SYCL compiler version: " << __SYCL_COMPILER_VERSION << std::endl; return 0; diff --git a/sycl/test/basic_tests/set_arg_error.cpp b/sycl/test/basic_tests/set_arg_error.cpp index 52df555902b44..d4215766c7a56 100644 --- a/sycl/test/basic_tests/set_arg_error.cpp +++ b/sycl/test/basic_tests/set_arg_error.cpp @@ -44,7 +44,7 @@ int main() { 5, ntc); h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} 4, NonTriviallyCopyable{}); -#if CL_SYCL_LANGUAGE_VERSION && CL_SYCL_LANGUAGE_VERSION <= 121 +#if SYCL_LANGUAGE_VERSION && SYCL_LANGUAGE_VERSION <= 201707 NonStdLayout nstd; h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} 6, nstd); diff --git a/sycl/test/functor/functor_inheritance.cpp b/sycl/test/functor/functor_inheritance.cpp index 3a3d5218ef6ec..436e8ce74a7d0 100644 --- a/sycl/test/functor/functor_inheritance.cpp +++ b/sycl/test/functor/functor_inheritance.cpp @@ -36,7 +36,7 @@ struct Derived : public Base, public SecondBase { int _A, int _B, int _C, int _D, int _E, cl::sycl::accessor &_Acc) : A(_A), Acc(_Acc), /*Out(_Out),*/ Base(_B, _C, _D), SecondBase(_E) {} - void operator()() { + void operator()() const { Acc[0] = this->A + this->B + this->InnerObj.C + this->InnerObj.D + this->E; } diff --git a/sycl/test/functor/kernel_functor.cpp b/sycl/test/functor/kernel_functor.cpp index 9dd5e0f2fecdf..9bed0b451ca25 100644 --- a/sycl/test/functor/kernel_functor.cpp +++ b/sycl/test/functor/kernel_functor.cpp @@ -23,7 +23,7 @@ constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer; // - functor class is defined in an anonymous namespace // - the '()' operator: // * does not have parameters (to be used in 'single_task'). -// * has no 'const' qualifier +// * has the 'const' qualifier namespace { class Functor1 { public: @@ -32,7 +32,7 @@ class Functor1 { cl::sycl::accessor &Acc_) : X(X_), Acc(Acc_) {} - void operator()() { Acc[0] += X; } + void operator()() const { Acc[0] += X; } private: int X; @@ -66,14 +66,14 @@ class Functor2 { // - functor class is templated and defined in the translation unit scope // - the '()' operator: // * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for'). -// * has no 'const' qualifier +// * has the 'const' qualifier template class TmplFunctor { public: TmplFunctor( T X_, cl::sycl::accessor &Acc_) : X(X_), Acc(Acc_) {} - void operator()(cl::sycl::id<1> id) { Acc[id] += X; } + void operator()(cl::sycl::id<1> id) const { Acc[id] += X; } private: T X; diff --git a/sycl/test/hier_par/hier_par_basic.cpp b/sycl/test/hier_par/hier_par_basic.cpp index d1a94ea1a7112..5670663e2e149 100644 --- a/sycl/test/hier_par/hier_par_basic.cpp +++ b/sycl/test/hier_par/hier_par_basic.cpp @@ -55,7 +55,7 @@ struct PFWIFunctor { : wg_chunk(wg_chunk), wg_size(wg_size), wg_offset(wg_offset), range_length(range_length), v(v), dev_ptr(dev_ptr) {} - void operator()(h_item<1> i) { + void operator()(h_item<1> i) const { // number of buf elements per work item: size_t wi_chunk = (wg_chunk + wg_size - 1) / wg_size; auto id = i.get_physical_local_id().get(0); @@ -82,7 +82,7 @@ struct PFWGFunctor { : wg_chunk(wg_chunk), range_length(range_length), dev_ptr(dev_ptr), addend(addend), n_iter(n_iter) {} - void operator()(group<1> g) { + void operator()(group<1> g) const { int v = addend; // to check constant initializer works too size_t wg_offset = wg_chunk * g.get_id(0); size_t wg_size = g.get_local_range(0); @@ -95,13 +95,13 @@ struct PFWGFunctor { } // Dummy operator '()' to make sure compiler can handle multiple '()' // operators/ and pick the right one for PFWG kernel code generation. - void operator()(int ind, int val) { dev_ptr[ind] += val; } + void operator()(int ind, int val) const { dev_ptr[ind] += val; } const size_t wg_chunk; const size_t range_length; const int n_iter; const int addend; - AccTy dev_ptr; + mutable AccTy dev_ptr; }; int main() { diff --git a/sycl/test/separate-compile/same-kernel.cpp b/sycl/test/separate-compile/same-kernel.cpp index 66ca32780f3cd..ad72decfa3604 100644 --- a/sycl/test/separate-compile/same-kernel.cpp +++ b/sycl/test/separate-compile/same-kernel.cpp @@ -26,9 +26,7 @@ class TestFnObj { TestFnObj(buffer &buf, handler &cgh) : data(buf.get_access(cgh)) {} accessor data; - void operator()(id<1> item) { - data[item] = item[0]; - } + void operator()(id<1> item) const { data[item] = item[0]; } }; void kernel2(); diff --git a/sycl/test/sub_group/attributes.cpp b/sycl/test/sub_group/attributes.cpp index d8173d2d1cf72..650ba23dd8c76 100644 --- a/sycl/test/sub_group/attributes.cpp +++ b/sycl/test/sub_group/attributes.cpp @@ -22,7 +22,7 @@ class KernelFunctor##SIZE { \ public: \ [[cl::intel_reqd_sub_group_size(SIZE)]] void \ - operator()(cl::sycl::nd_item<1> Item) { \ + operator()(cl::sycl::nd_item<1> Item) const { \ const auto GID = Item.get_global_id(); \ } \ };