From 694bbcbdb7e2a484f98b1f2bdb75dfa1f31c085d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 4 Aug 2020 16:33:28 -0700 Subject: [PATCH 01/17] immutable functions redo. WIP. Signed-off-by: Chris Perkins --- clang/include/clang/Basic/DiagnosticGroups.td | 4 +- .../clang/Basic/DiagnosticSemaKinds.td | 9 + clang/lib/Frontend/CompilerInvocation.cpp | 1 + clang/lib/Sema/SemaSYCL.cpp | 35 +++- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 12 +- .../CodeGenSYCL/address-space-cond-op.cpp | 2 +- clang/test/CodeGenSYCL/address-space-new.cpp | 4 +- .../CodeGenSYCL/address-space-of-returns.cpp | 2 +- .../address-space-parameter-conversions.cpp | 2 +- .../test/CodeGenSYCL/basic-kernel-wrapper.cpp | 2 +- clang/test/CodeGenSYCL/const-wg-init.cpp | 2 +- .../CodeGenSYCL/debug-info-srcpos-kernel.cpp | 2 +- clang/test/CodeGenSYCL/device-functions.cpp | 2 +- clang/test/CodeGenSYCL/device-variables.cpp | 2 +- .../emit-kernel-in-virtual-func.cpp | 2 +- .../emit-kernel-in-virtual-func2.cpp | 4 +- clang/test/CodeGenSYCL/esimd_metadata1.cpp | 2 +- clang/test/CodeGenSYCL/fpga_pipes.cpp | 2 +- clang/test/CodeGenSYCL/inheritance.cpp | 2 +- clang/test/CodeGenSYCL/inline_asm.cpp | 2 +- clang/test/CodeGenSYCL/inlining.cpp | 2 +- clang/test/CodeGenSYCL/int_header1.cpp | 2 +- .../test/CodeGenSYCL/int_header_inline_ns.cpp | 2 +- clang/test/CodeGenSYCL/integration_header.cpp | 2 +- .../CodeGenSYCL/intel-fpga-ivdep-array.cpp | 2 +- .../intel-fpga-ivdep-embedded-loops.cpp | 2 +- .../CodeGenSYCL/intel-fpga-ivdep-global.cpp | 2 +- clang/test/CodeGenSYCL/intel-fpga-local.cpp | 2 +- clang/test/CodeGenSYCL/intel-fpga-loops.cpp | 2 +- .../CodeGenSYCL/intel-fpga-mem-builtin.cpp | 2 +- .../intel-fpga-no-global-work-offset.cpp | 4 +- clang/test/CodeGenSYCL/intel-fpga-reg.cpp | 2 +- .../CodeGenSYCL/intel-max-global-work-dim.cpp | 4 +- .../CodeGenSYCL/intel-max-work-group-size.cpp | 4 +- clang/test/CodeGenSYCL/intel-restrict.cpp | 2 +- .../test/CodeGenSYCL/kernel-by-reference.cpp | 42 +++++ clang/test/CodeGenSYCL/kernel-name.cpp | 2 +- .../CodeGenSYCL/kernel-param-acc-array-ih.cpp | 2 +- .../CodeGenSYCL/kernel-param-acc-array.cpp | 2 +- .../kernel-param-member-acc-array-ih.cpp | 2 +- .../kernel-param-member-acc-array.cpp | 2 +- .../CodeGenSYCL/kernel-param-pod-array-ih.cpp | 2 +- .../CodeGenSYCL/kernel-param-pod-array.cpp | 2 +- clang/test/CodeGenSYCL/kernel_functor.cpp | 63 +------ .../CodeGenSYCL/kernel_name_with_typedefs.cpp | 4 +- clang/test/CodeGenSYCL/kernelname-enum.cpp | 16 +- clang/test/CodeGenSYCL/loop_unroll.cpp | 2 +- clang/test/CodeGenSYCL/module-id.cpp | 2 +- clang/test/CodeGenSYCL/noexcept.cpp | 2 +- .../test/CodeGenSYCL/num-simd-work-items.cpp | 4 +- clang/test/CodeGenSYCL/remove-ur-inst.cpp | 2 +- .../test/CodeGenSYCL/reqd-sub-group-size.cpp | 8 +- .../test/CodeGenSYCL/reqd-work-group-size.cpp | 6 +- clang/test/CodeGenSYCL/sampler.cpp | 2 +- clang/test/CodeGenSYCL/spir-calling-conv.cpp | 2 +- clang/test/CodeGenSYCL/spir-enum.cpp | 2 +- .../test/CodeGenSYCL/spir-opencl-version.cpp | 2 +- .../CodeGenSYCL/sycl-device-static-init.cpp | 2 +- .../CodeGenSYCL/sycl-multi-kernel-attr.cpp | 5 +- clang/test/CodeGenSYCL/unique-stable-name.cpp | 2 +- clang/test/CodeGenSYCL/usm-int-header.cpp | 2 +- clang/test/CodeGenSYCL/virtual-types.cpp | 2 +- clang/test/CodeGenSYCL/wrapped-accessor.cpp | 2 +- clang/test/SemaSYCL/Inputs/sycl.hpp | 4 +- .../SemaSYCL/accessor-type-diagnostics.cpp | 4 +- .../test/SemaSYCL/accessors-targets-image.cpp | 2 +- clang/test/SemaSYCL/accessors-targets.cpp | 2 +- .../SemaSYCL/allow-constexpr-recursion.cpp | 4 +- .../test/SemaSYCL/array-kernel-param-neg.cpp | 6 +- clang/test/SemaSYCL/array-kernel-param.cpp | 2 +- clang/test/SemaSYCL/basic-kernel-wrapper.cpp | 2 +- .../SemaSYCL/built-in-type-kernel-arg.cpp | 2 +- .../SemaSYCL/call-to-undefined-function.cpp | 4 +- .../deferred-diagnostics-aux-builtin.cpp | 4 +- .../SemaSYCL/deferred-diagnostics-emit.cpp | 4 +- clang/test/SemaSYCL/fake-accessors.cpp | 2 +- clang/test/SemaSYCL/float128.cpp | 6 +- clang/test/SemaSYCL/forward-decl.cpp | 4 +- clang/test/SemaSYCL/fpga_pipes.cpp | 2 +- clang/test/SemaSYCL/implicit_kernel_type.cpp | 50 ++++-- clang/test/SemaSYCL/inheritance.cpp | 2 +- clang/test/SemaSYCL/inline-asm.cpp | 8 +- clang/test/SemaSYCL/intel-fpga-local.cpp | 4 +- clang/test/SemaSYCL/intel-fpga-loops.cpp | 4 +- .../test/SemaSYCL/intel-fpga-mem-builtin.cpp | 2 +- clang/test/SemaSYCL/intel-fpga-reg.cpp | 2 +- .../SemaSYCL/intel-max-global-work-dim.cpp | 27 +-- .../SemaSYCL/intel-max-work-group-size.cpp | 19 ++- .../SemaSYCL/intel-reqd-work-group-size.cpp | 28 ++-- clang/test/SemaSYCL/intel-restrict.cpp | 9 +- clang/test/SemaSYCL/kernel-function-type.cpp | 4 +- clang/test/SemaSYCL/kernelname-enum.cpp | 10 +- .../SemaSYCL/lambda_implicit_capture_this.cpp | 4 +- clang/test/SemaSYCL/loop_unroll.cpp | 4 +- clang/test/SemaSYCL/mangle-kernel.cpp | 2 +- .../SemaSYCL/markfunction-astconsumer.cpp | 4 +- clang/test/SemaSYCL/no-vtables.cpp | 5 +- clang/test/SemaSYCL/no-vtables2.cpp | 5 +- clang/test/SemaSYCL/non-std-layout-param.cpp | 9 +- .../non-trivially-copyable-kernel-param.cpp | 4 +- clang/test/SemaSYCL/num_simd_work_items.cpp | 16 +- clang/test/SemaSYCL/prohibit-thread-local.cpp | 4 +- .../test/SemaSYCL/reference-kernel-param.cpp | 4 +- .../SemaSYCL/reqd-sub-group-size-device.cpp | 14 +- .../SemaSYCL/reqd-work-group-size-device.cpp | 18 +- clang/test/SemaSYCL/restrict-recursion.cpp | 10 +- clang/test/SemaSYCL/restrict-recursion2.cpp | 4 +- clang/test/SemaSYCL/restrict-recursion3.cpp | 4 +- clang/test/SemaSYCL/restrict-recursion4.cpp | 4 +- clang/test/SemaSYCL/sampler.cpp | 2 +- .../SemaSYCL/spec_const_and_accesor_crash.cpp | 2 +- clang/test/SemaSYCL/sycl-callstack.cpp | 7 +- clang/test/SemaSYCL/sycl-cconv.cpp | 4 +- .../SemaSYCL/sycl-device-const-static.cpp | 4 +- .../SemaSYCL/sycl-device-static-restrict.cpp | 4 +- .../SemaSYCL/sycl-dllimport-dllexport.cpp | 10 +- clang/test/SemaSYCL/sycl-esimd.cpp | 12 +- clang/test/SemaSYCL/sycl-fptr-lambda.cpp | 4 +- clang/test/SemaSYCL/sycl-pseudo-dtor.cpp | 4 +- clang/test/SemaSYCL/sycl-restrict.cpp | 10 +- clang/test/SemaSYCL/sycl-varargs-cconv.cpp | 2 +- clang/test/SemaSYCL/tls_error.cpp | 4 +- clang/test/SemaSYCL/unevaluated-function.cpp | 4 +- clang/test/SemaSYCL/unnamed-kernel.cpp | 4 +- clang/test/SemaSYCL/unsupported_math.cpp | 4 +- clang/test/SemaSYCL/variadic-func-call.cpp | 4 +- clang/test/SemaSYCL/wrapped-accessor.cpp | 2 +- sycl/include/CL/sycl/handler.hpp | 115 ++++++++++++- sycl/include/CL/sycl/queue.hpp | 154 +++++++++++++++--- sycl/test/functor/functor_inheritance.cpp | 2 +- sycl/test/functor/kernel_functor.cpp | 8 +- sycl/test/hier_par/hier_par_basic.cpp | 8 +- sycl/test/separate-compile/same-kernel.cpp | 4 +- sycl/test/sub_group/attributes.cpp | 2 +- 134 files changed, 641 insertions(+), 396 deletions(-) create mode 100644 clang/test/CodeGenSYCL/kernel-by-reference.cpp 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..77017bdcfac05 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10995,6 +10995,7 @@ def warn_boolean_attribute_argument_is_not_valid: Warning< def err_sycl_attibute_cannot_be_applied_here : Error<"%0 attribute cannot be applied to a " "static function or function in an anonymous namespace">; +<<<<<<< HEAD def err_sycl_compiletime_property_duplication : Error< "Can't apply %0 property twice to the same accessor">; def err_sycl_invalid_property_list_param_number : Error< @@ -11005,6 +11006,14 @@ 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; +>>>>>>> immutable functions redo. WIP. 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/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f339c4a69b3be..969ad1c680e59 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -707,7 +707,39 @@ 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->param_begin())->getType(); + // In SYCL 2020 kernels are now passed by reference. + if (KernelParamTy->isReferenceType()) + return KernelParamTy->getPointeeCXXRecordDecl(); + + // SYCL 1.2.1 + return KernelParamTy->getAsCXXRecordDecl(); +} + +static void checkKernelAndCaller(Sema &SemaRef, FunctionDecl *Caller, + const CXXRecordDecl *KernelObj) { + // check captures + if (KernelObj->isLambda()) { + for (const LambdaCapture &LC : KernelObj->captures()) + if (LC.capturesThis() && LC.isImplicit()) + SemaRef.Diag(LC.getLocation(), diag::err_implicit_this_capture); + } + + // check that calling kernel conforms to spec + assert(Caller->param_size() >= 1 && "missing kernel function argument."); + QualType KernelParamTy = (*Caller->param_begin())->getType(); + if (KernelParamTy->isReferenceType()) { + // passing by reference, so emit warning if not using SYCL 2020 + if (SemaRef.LangOpts.SYCLVersion < 2020) + SemaRef.Diag(Caller->getLocation(), + diag::warn_sycl_pass_by_reference_future); + } else { + // passing by value. emit warning if using SYCL 2020 or greater + if (SemaRef.LangOpts.SYCLVersion > 2017) + SemaRef.Diag(Caller->getLocation(), + diag::warn_sycl_pass_by_value_deprecated); + } } /// Creates a kernel parameter descriptor @@ -2284,6 +2316,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, // The first argument to the KernelCallerFunc is the lambda object. const CXXRecordDecl *KernelObj = getKernelObjectType(KernelCallerFunc); assert(KernelObj && "invalid kernel caller"); + checkKernelAndCaller(*this, KernelCallerFunc, KernelObj); // Calculate both names, since Integration headers need both. std::string CalculatedName, StableName; 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/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..e6b958d130caa 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. @@ -7,7 +7,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-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/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..022d34ad59227 100644 --- a/clang/test/SemaSYCL/implicit_kernel_type.cpp +++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp @@ -1,13 +1,43 @@ // 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; +//using namespace cl::sycl; void function() { } @@ -20,11 +50,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 +66,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 +74,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..cfde3fd4d7c3b 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,22 +24,21 @@ 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 @@ -46,12 +46,13 @@ 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()() {} + 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..737afb7d0327d 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -33,6 +33,10 @@ #include #include +#if !CL_SYCL_LANGUAGE_VERSION || CL_SYCL_LANGUAGE_VERSION < 2020 +#define __SYCL_NONCONST_FUNCTOR__ +#endif + template @@ -719,14 +723,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 +749,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))); } @@ -839,7 +857,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 +878,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)); } @@ -929,7 +963,11 @@ class __SYCL_EXPORT handler { template void parallel_for(range NumWorkItems, id WorkItemOffset, +#ifdef __SYCL_NONCONST_FUNCTOR__ KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -961,7 +999,12 @@ 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, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -988,7 +1031,12 @@ 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, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif intel::detail::reduCGFunc(*this, KernelFunc, Range, Redu, Redu.getUserAccessor()); } @@ -1001,7 +1049,12 @@ 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, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif intel::detail::reduCGFunc(*this, KernelFunc, Range, Redu, Redu.getUSMPointer()); } @@ -1020,7 +1073,12 @@ 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, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif shared_ptr_class QueueCopy = MQueue; auto RWAcc = Redu.getReadWriteScalarAcc(*this); intel::detail::reduCGFunc(*this, KernelFunc, Range, Redu, @@ -1056,7 +1114,12 @@ 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, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif // 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 +1194,11 @@ class __SYCL_EXPORT handler { template void parallel_for_work_group(range NumWorkGroups, +#ifdef __SYCL_NONCONST_FUNCTOR__ KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1164,7 +1231,11 @@ class __SYCL_EXPORT handler { int Dims> void parallel_for_work_group(range NumWorkGroups, range WorkGroupSize, +#ifdef __SYCL_NONCONST_FUNCTOR__ KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1266,7 +1337,12 @@ 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, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1306,7 +1382,11 @@ class __SYCL_EXPORT handler { template void parallel_for(kernel Kernel, range NumWorkItems, +#ifdef __SYCL_NONCONST_FUNCTOR__ KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1341,7 +1421,12 @@ class __SYCL_EXPORT handler { template void parallel_for(kernel Kernel, range NumWorkItems, - id WorkItemOffset, KernelType KernelFunc) { + id WorkItemOffset, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1378,7 +1463,11 @@ class __SYCL_EXPORT handler { template void parallel_for(kernel Kernel, nd_range NDRange, +#ifdef __SYCL_NONCONST_FUNCTOR__ KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1420,7 +1509,11 @@ class __SYCL_EXPORT handler { template void parallel_for_work_group(kernel Kernel, range NumWorkGroups, +#ifdef __SYCL_NONCONST_FUNCTOR__ KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1458,7 +1551,11 @@ class __SYCL_EXPORT handler { int Dims> void parallel_for_work_group(kernel Kernel, range NumWorkGroups, range WorkGroupSize, +#ifdef __SYCL_NONCONST_FUNCTOR__ KernelType KernelFunc) { +#else + const KernelType &KernelFunc) { +#endif throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index 5338567f5eda3..09c7da71a3dcb 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -361,21 +361,47 @@ 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(A,B,C,D,E) and +// CODELOCARG(B,C,D) Similarly, the KernelFunc param is simplified to be +// KERNELFUNCPARAM(B,F,G) Once the queue kernel functions are defined, these +// macros are #undef immediately. Compare this first definition of single_task +// here, with the second one below. replace CODELOCPARAM(A,B,C,D,E) with nothing +// or : , const detail::code_location &CodeLoc = +// detail::code_location::current() replace CODELOCARG(B,C,D) with nothign or : +// const detail::code_location &CodeLoc = {}; +#define A , +#define B const +#define C detail::code_location +#define D &CodeLoc +#define E detail::code_location::current() +#ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA +#define CODELOCPARAM(a, b, c, d, e) a b c d = e + +#define CODELOCARG(b, c, d) +#else +#define CODELOCPARAM(a, b, c, d, e) + +#define CODELOCARG(b, c, d) b c d = {} +#endif +// replace KERNELFUNCPARAM(B,F,G) with KernelType KernelFunc +// or const KernelType &KernelFunc +#define F KernelType +#define G KernelFunc +#ifdef __SYCL_NONCONST_FUNCTOR__ +#define KERNELFUNCPARAM(B, F, G) F G +#else +#define KERNELFUNCPARAM(B, F, G) B F &G +#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(B, F, G) CODELOCPARAM(A, B, C, D, E)) { + CODELOCARG(B, C, D); + return submit( [&](handler &CGH) { CGH.template single_task(KernelFunc); @@ -383,6 +409,18 @@ class __SYCL_EXPORT queue { CodeLoc); } +// Clean up CODELOC and KERNELFUNC macros. +#undef A +#undef B +#undef C +#undef D +#undef E +#undef F +#undef G +#undef CODELOCPARAM +#undef CODELOCARG +#undef KERNELFUNCPARAM + /// single_task version with a kernel represented as a lambda. /// /// \param DepEvent is an event that specifies the kernel dependencies @@ -390,7 +428,13 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event single_task( - event DepEvent, KernelType KernelFunc + event DepEvent, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc +#else + const KernelType &KernelFunc +#endif + #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -415,7 +459,13 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event single_task( - const vector_class &DepEvents, KernelType KernelFunc + const vector_class &DepEvents, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc +#else + const KernelType &KernelFunc +#endif + #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -440,7 +490,13 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for( - range<1> NumWorkItems, KernelType KernelFunc + range<1> NumWorkItems, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc +#else + const KernelType &KernelFunc +#endif + #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -460,7 +516,13 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for( - range<2> NumWorkItems, KernelType KernelFunc + range<2> NumWorkItems, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc +#else + const KernelType &KernelFunc +#endif + #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -480,7 +542,13 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for( - range<3> NumWorkItems, KernelType KernelFunc + range<3> NumWorkItems, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc +#else + const KernelType &KernelFunc +#endif + #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -501,7 +569,13 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for( - range<1> NumWorkItems, event DepEvent, KernelType KernelFunc + range<1> NumWorkItems, event DepEvent, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc +#else + const KernelType &KernelFunc +#endif + #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -523,7 +597,13 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for( - range<2> NumWorkItems, event DepEvent, KernelType KernelFunc + range<2> NumWorkItems, event DepEvent, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc +#else + const KernelType &KernelFunc +#endif + #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -545,7 +625,13 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for( - range<3> NumWorkItems, event DepEvent, KernelType KernelFunc + range<3> NumWorkItems, event DepEvent, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc +#else + const KernelType &KernelFunc +#endif + #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -569,7 +655,7 @@ class __SYCL_EXPORT queue { template event parallel_for( range<1> NumWorkItems, const vector_class &DepEvents, - KernelType KernelFunc + const KernelType &KernelFunc #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -593,7 +679,7 @@ class __SYCL_EXPORT queue { template event parallel_for( range<2> NumWorkItems, const vector_class &DepEvents, - KernelType KernelFunc + const KernelType &KernelFunc #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -617,7 +703,7 @@ class __SYCL_EXPORT queue { template event parallel_for( range<3> NumWorkItems, const vector_class &DepEvents, - KernelType KernelFunc + const KernelType &KernelFunc #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -640,7 +726,8 @@ class __SYCL_EXPORT queue { template event parallel_for( - range NumWorkItems, id WorkItemOffset, KernelType KernelFunc + range NumWorkItems, id WorkItemOffset, + const KernelType &KernelFunc #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -669,7 +756,7 @@ class __SYCL_EXPORT queue { int Dims> event parallel_for( range NumWorkItems, id WorkItemOffset, event DepEvent, - KernelType KernelFunc + const KernelType &KernelFunc #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -700,7 +787,13 @@ class __SYCL_EXPORT queue { int Dims> event parallel_for( range NumWorkItems, id WorkItemOffset, - const vector_class &DepEvents, KernelType KernelFunc + const vector_class &DepEvents, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc +#else + const KernelType &KernelFunc +#endif + #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -728,7 +821,13 @@ class __SYCL_EXPORT queue { template event parallel_for( - nd_range ExecutionRange, KernelType KernelFunc + nd_range ExecutionRange, +#ifdef __SYCL_NONCONST_FUNCTOR__ + KernelType KernelFunc +#else + const KernelType &KernelFunc +#endif + #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -756,7 +855,8 @@ class __SYCL_EXPORT queue { template event parallel_for( - nd_range ExecutionRange, event DepEvent, KernelType KernelFunc + nd_range ExecutionRange, event DepEvent, + const KernelType &KernelFunc #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() @@ -787,7 +887,7 @@ class __SYCL_EXPORT queue { int Dims> event parallel_for( nd_range ExecutionRange, const vector_class &DepEvents, - KernelType KernelFunc + const KernelType &KernelFunc #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA , const detail::code_location &CodeLoc = detail::code_location::current() 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(); \ } \ }; From d68ba917b35e7124272ab47c256a640bd26ada35 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 4 Aug 2020 18:24:41 -0700 Subject: [PATCH 02/17] hooked up -sycl-std flag to set SYCL version constant. Needs test. Signed-off-by: Chris Perkins --- clang/lib/Driver/ToolChains/Clang.cpp | 11 +++++++++++ sycl/include/CL/sycl/queue.hpp | 9 ++++++--- 2 files changed, 17 insertions(+), 3 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 61ee698ffbd9e..283a458db2fc8 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -4160,9 +4160,20 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, if (SYCLStdArg) { SYCLStdArg->render(Args, CmdArgs); CmdArgs.push_back("-fsycl-std-layout-kernel-params"); + + unsigned version = llvm::StringSwitch(SYCLStdArg->getValue()) + .Cases("2017", "1.2.1", "121", "sycl-1.2.1", 2017) + .Case("2020", 2020) + .Default(0); + if (version == 2017) + CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=121"); + else if (version == 2020) + CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=2020"); + } else { // Ensure the default version in SYCL mode is 1.2.1 (aka 2017) CmdArgs.push_back("-sycl-std=2017"); + CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=121"); } } diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index 09c7da71a3dcb..bdd0823cd3528 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -366,10 +366,13 @@ class __SYCL_EXPORT queue { // CODELOCARG(B,C,D) Similarly, the KernelFunc param is simplified to be // KERNELFUNCPARAM(B,F,G) Once the queue kernel functions are defined, these // macros are #undef immediately. Compare this first definition of single_task -// here, with the second one below. replace CODELOCPARAM(A,B,C,D,E) with nothing +// here, with the second one below. + +// replace CODELOCPARAM(A,B,C,D,E) with nothing // or : , const detail::code_location &CodeLoc = -// detail::code_location::current() replace CODELOCARG(B,C,D) with nothign or : -// const detail::code_location &CodeLoc = {}; +// detail::code_location::current() +// replace CODELOCARG(B,C,D) with nothing +// or : const detail::code_location &CodeLoc = {} #define A , #define B const #define C detail::code_location From 8390646aaee35f4eca8300a606a83ab11df03d29 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 4 Aug 2020 20:23:17 -0700 Subject: [PATCH 03/17] simplified macros - tricky re-use not worth it. Signed-off-by: Chris Perkins --- sycl/include/CL/sycl/queue.hpp | 60 ++++++++++++++-------------------- 1 file changed, 24 insertions(+), 36 deletions(-) diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index bdd0823cd3528..880e6f5bf652b 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -361,40 +361,35 @@ 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(A,B,C,D,E) and -// CODELOCARG(B,C,D) Similarly, the KernelFunc param is simplified to be -// KERNELFUNCPARAM(B,F,G) Once the queue kernel functions are defined, these -// macros are #undef immediately. Compare this first definition of single_task -// here, with the second one below. - -// replace CODELOCPARAM(A,B,C,D,E) with nothing -// or : , const detail::code_location &CodeLoc = -// detail::code_location::current() -// replace CODELOCARG(B,C,D) with nothing -// or : const detail::code_location &CodeLoc = {} -#define A , -#define B const -#define C detail::code_location -#define D &CodeLoc -#define E detail::code_location::current() + // 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. Compare this first definition of + // single_task here, with the second one below. + + // 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, b, c, d, e) a b c d = e +#define CODELOCPARAM(a) \ + , const detail::code_location a = detail::code_location::current() -#define CODELOCARG(b, c, d) +#define CODELOCARG(a) #else -#define CODELOCPARAM(a, b, c, d, e) +#define CODELOCPARAM(a) -#define CODELOCARG(b, c, d) b c d = {} +#define CODELOCARG(a) const detail::code_location a = {} #endif -// replace KERNELFUNCPARAM(B,F,G) with KernelType KernelFunc -// or const KernelType &KernelFunc -#define F KernelType -#define G KernelFunc +// replace KERNELFUNCPARAM(KernelFunc) with KernelType KernelFunc +// or const KernelType &KernelFunc #ifdef __SYCL_NONCONST_FUNCTOR__ -#define KERNELFUNCPARAM(B, F, G) F G +#define KERNELFUNCPARAM(a) KernelType a #else -#define KERNELFUNCPARAM(B, F, G) B F &G +#define KERNELFUNCPARAM(a) const KernelType &a #endif /// single_task version with a kernel represented as a lambda. @@ -402,8 +397,8 @@ 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(KERNELFUNCPARAM(B, F, G) CODELOCPARAM(A, B, C, D, E)) { - CODELOCARG(B, C, D); + event single_task(KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { + CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { @@ -413,13 +408,6 @@ class __SYCL_EXPORT queue { } // Clean up CODELOC and KERNELFUNC macros. -#undef A -#undef B -#undef C -#undef D -#undef E -#undef F -#undef G #undef CODELOCPARAM #undef CODELOCARG #undef KERNELFUNCPARAM From f5df1d581d9ea9c2d6a00f7e71939aba188086d5 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 4 Aug 2020 20:53:33 -0700 Subject: [PATCH 04/17] test added Signed-off-by: Chris Perkins --- clang/test/CodeGenSYCL/sycl-std-flag.cpp | 26 ++++++++++++++++++++++++ 1 file changed, 26 insertions(+) create mode 100644 clang/test/CodeGenSYCL/sycl-std-flag.cpp diff --git a/clang/test/CodeGenSYCL/sycl-std-flag.cpp b/clang/test/CodeGenSYCL/sycl-std-flag.cpp new file mode 100644 index 0000000000000..0411c19227853 --- /dev/null +++ b/clang/test/CodeGenSYCL/sycl-std-flag.cpp @@ -0,0 +1,26 @@ +// RUN: %clang++ -fsycl -fsyntax-only -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang++ -fsycl -fsyntax-only -sycl-std=121 -DSYCL2017 %s +// RUN: %clang++ -fsycl -fsyntax-only -sycl-std=1.2.1 -DSYCL2017 %s +// RUN: %clang++ -fsycl -fsyntax-only -sycl-std=sycl-1.2.1 -DSYCL2017 %s +// RUN: %clang++ -fsycl -fsyntax-only -sycl-std=2020 -DSYCL2020 %s + + +#if defined(SYCL2017) + #if !CL_SYCL_LANGUAGE_VERSION + static_assert(false , "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used"); + #endif + #if CL_SYCL_LANGUAGE_VERSION > 121 + static_assert(false, "-sycl-std flag should have set CL_SYCL_LANGUAGE_VERSION to 121"); + #endif +#endif + +#if defined(SYCL2020) + #if !CL_SYCL_LANGUAGE_VERSION + static_assert(false , "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used"); + #endif + #if CL_SYCL_LANGUAGE_VERSION < 2020 + static_assert(false , "-sycl-std=2020 flag should set CL_SYCL_LANGUAGE_VERSION to 2020"); + #endif +#endif + +// expected-no-diagnostics \ No newline at end of file From 02dcc87d7454b990ff8693ed5de6f577e0cea932 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 5 Aug 2020 10:08:54 -0700 Subject: [PATCH 05/17] last changes, clang-formattery. Signed-off-by: Chris Perkins --- clang/test/CodeGenSYCL/sycl-std-flag.cpp | 35 +- .../SemaSYCL/intel-max-global-work-dim.cpp | 2 +- sycl/include/CL/sycl/queue.hpp | 307 ++++-------------- 3 files changed, 82 insertions(+), 262 deletions(-) diff --git a/clang/test/CodeGenSYCL/sycl-std-flag.cpp b/clang/test/CodeGenSYCL/sycl-std-flag.cpp index 0411c19227853..e2b4c352127e0 100644 --- a/clang/test/CodeGenSYCL/sycl-std-flag.cpp +++ b/clang/test/CodeGenSYCL/sycl-std-flag.cpp @@ -1,26 +1,27 @@ -// RUN: %clang++ -fsycl -fsyntax-only -sycl-std=2017 -DSYCL2017 %s -// RUN: %clang++ -fsycl -fsyntax-only -sycl-std=121 -DSYCL2017 %s -// RUN: %clang++ -fsycl -fsyntax-only -sycl-std=1.2.1 -DSYCL2017 %s -// RUN: %clang++ -fsycl -fsyntax-only -sycl-std=sycl-1.2.1 -DSYCL2017 %s -// RUN: %clang++ -fsycl -fsyntax-only -sycl-std=2020 -DSYCL2020 %s +// REQUIRES: clang-driver +// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=2017 -DSYCL2017 %s +// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=121 -DSYCL2017 %s +// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=1.2.1 -DSYCL2017 %s +// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=sycl-1.2.1 -DSYCL2017 %s +// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=2020 -DSYCL2020 %s #if defined(SYCL2017) - #if !CL_SYCL_LANGUAGE_VERSION - static_assert(false , "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used"); - #endif - #if CL_SYCL_LANGUAGE_VERSION > 121 - static_assert(false, "-sycl-std flag should have set CL_SYCL_LANGUAGE_VERSION to 121"); - #endif +#if !CL_SYCL_LANGUAGE_VERSION +static_assert(false, "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used"); +#endif +#if CL_SYCL_LANGUAGE_VERSION > 121 +static_assert(false, "-sycl-std flag should have set CL_SYCL_LANGUAGE_VERSION to 121"); +#endif #endif #if defined(SYCL2020) - #if !CL_SYCL_LANGUAGE_VERSION - static_assert(false , "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used"); - #endif - #if CL_SYCL_LANGUAGE_VERSION < 2020 - static_assert(false , "-sycl-std=2020 flag should set CL_SYCL_LANGUAGE_VERSION to 2020"); - #endif +#if !CL_SYCL_LANGUAGE_VERSION +static_assert(false, "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used"); +#endif +#if CL_SYCL_LANGUAGE_VERSION < 2020 +static_assert(false, "-sycl-std=2020 flag should set CL_SYCL_LANGUAGE_VERSION to 2020"); +#endif #endif // expected-no-diagnostics \ No newline at end of file diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim.cpp index cfde3fd4d7c3b..c039787b6d8cf 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim.cpp @@ -45,7 +45,7 @@ struct TRIFuncObjGood2 { 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}} + [[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 {} }; diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index 880e6f5bf652b..8d6a01d07af0a 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -365,8 +365,7 @@ class __SYCL_EXPORT queue { // 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. Compare this first definition of - // single_task here, with the second one below. + // these macros are #undef immediately. // replace CODELOCPARAM(&CodeLoc) with nothing // or : , const detail::code_location &CodeLoc = @@ -407,10 +406,6 @@ class __SYCL_EXPORT queue { CodeLoc); } -// Clean up CODELOC and KERNELFUNC macros. -#undef CODELOCPARAM -#undef CODELOCARG -#undef KERNELFUNCPARAM /// single_task version with a kernel represented as a lambda. /// @@ -418,22 +413,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, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc -#else - const KernelType &KernelFunc -#endif - -#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); @@ -449,22 +431,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, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc -#else - const KernelType &KernelFunc -#endif - -#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); @@ -480,22 +449,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, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc -#else - const KernelType &KernelFunc -#endif - -#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); } @@ -506,22 +462,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, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc -#else - const KernelType &KernelFunc -#endif - -#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); } @@ -532,22 +475,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, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc -#else - const KernelType &KernelFunc -#endif - -#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); } @@ -559,22 +489,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, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc -#else - const KernelType &KernelFunc -#endif - -#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); } @@ -587,22 +504,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, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc -#else - const KernelType &KernelFunc -#endif - -#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); } @@ -615,22 +519,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, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc -#else - const KernelType &KernelFunc -#endif - -#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); } @@ -644,17 +535,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, - const 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); } @@ -668,17 +552,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, - const 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); } @@ -692,17 +569,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, - const 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); } @@ -716,17 +586,9 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template - event parallel_for( - range NumWorkItems, id WorkItemOffset, - const 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( @@ -745,17 +607,10 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template - event parallel_for( - range NumWorkItems, id WorkItemOffset, event DepEvent, - const 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); @@ -776,23 +631,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, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc -#else - const KernelType &KernelFunc -#endif - -#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); @@ -811,22 +653,9 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template - event parallel_for( - nd_range ExecutionRange, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc -#else - const KernelType &KernelFunc -#endif - -#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, @@ -845,17 +674,9 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template - event parallel_for( - nd_range ExecutionRange, event DepEvent, - const 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); @@ -876,17 +697,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, - const 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); @@ -896,6 +710,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() From b0cffd8199090d8c9c0e83768b6548dbf0488a3d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 5 Aug 2020 10:18:10 -0700 Subject: [PATCH 06/17] clang-formatota Signed-off-by: Chris Perkins --- sycl/include/CL/sycl/queue.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index 8d6a01d07af0a..adbc33533514b 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -406,7 +406,6 @@ class __SYCL_EXPORT queue { CodeLoc); } - /// single_task version with a kernel represented as a lambda. /// /// \param DepEvent is an event that specifies the kernel dependencies From 7d43180024a2346113bc4bbe0b819c0c8d8b68b4 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 5 Aug 2020 13:08:20 -0700 Subject: [PATCH 07/17] rebased to remove conflicts --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 3 --- 1 file changed, 3 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 77017bdcfac05..f3b49747761f2 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10995,7 +10995,6 @@ def warn_boolean_attribute_argument_is_not_valid: Warning< def err_sycl_attibute_cannot_be_applied_here : Error<"%0 attribute cannot be applied to a " "static function or function in an anonymous namespace">; -<<<<<<< HEAD def err_sycl_compiletime_property_duplication : Error< "Can't apply %0 property twice to the same accessor">; def err_sycl_invalid_property_list_param_number : Error< @@ -11006,14 +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; ->>>>>>> immutable functions redo. WIP. 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 " From daf304c69ed402738619ea729b1ffb23190c3ab3 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 5 Aug 2020 14:59:22 -0700 Subject: [PATCH 08/17] new test needed flag Signed-off-by: Chris Perkins --- clang/test/SemaSYCL/buffer_location.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) 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" From 388a6db8a3a3ae889266c1ae6a4f4eeef6c396c4 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 12 Aug 2020 14:39:46 -0700 Subject: [PATCH 09/17] rebased, incorporated latest PR feedback Signed-off-by: Chris Perkins --- clang/lib/Driver/ToolChains/Clang.cpp | 6 -- clang/lib/Frontend/InitPreprocessor.cpp | 3 + clang/test/CodeGenSYCL/sycl-std-flag.cpp | 27 ----- clang/test/Preprocessor/sycl-macro.cpp | 5 + clang/test/SemaSYCL/implicit_kernel_type.cpp | 1 - sycl/include/CL/sycl/handler.hpp | 97 +++++------------- sycl/include/CL/sycl/queue.hpp | 102 +++++++++---------- 7 files changed, 84 insertions(+), 157 deletions(-) delete mode 100644 clang/test/CodeGenSYCL/sycl-std-flag.cpp diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 283a458db2fc8..4f100cb68eb38 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -4165,15 +4165,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, .Cases("2017", "1.2.1", "121", "sycl-1.2.1", 2017) .Case("2020", 2020) .Default(0); - if (version == 2017) - CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=121"); - else if (version == 2020) - CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=2020"); - } else { // Ensure the default version in SYCL mode is 1.2.1 (aka 2017) CmdArgs.push_back("-sycl-std=2017"); - CmdArgs.push_back("-DCL_SYCL_LANGUAGE_VERSION=121"); } } diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index e0342209d7814..a8c78a5846c58 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -465,6 +465,9 @@ static void InitializeStandardPredefinedMacros(const TargetInfo &TI, // SYCL Version is set to a value when building SYCL applications if (LangOpts.SYCLVersion == 2017) Builder.defineMacro("CL_SYCL_LANGUAGE_VERSION", "121"); + else if (LangOpts.SYCLVersion == 2020) + Builder.defineMacro("CL_SYCL_LANGUAGE_VERSION", "202001"); + if (LangOpts.SYCLValueFitInMaxInt) Builder.defineMacro("__SYCL_ID_QUERIES_FIT_IN_INT__", "1"); } diff --git a/clang/test/CodeGenSYCL/sycl-std-flag.cpp b/clang/test/CodeGenSYCL/sycl-std-flag.cpp deleted file mode 100644 index e2b4c352127e0..0000000000000 --- a/clang/test/CodeGenSYCL/sycl-std-flag.cpp +++ /dev/null @@ -1,27 +0,0 @@ -// REQUIRES: clang-driver - -// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=2017 -DSYCL2017 %s -// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=121 -DSYCL2017 %s -// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=1.2.1 -DSYCL2017 %s -// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=sycl-1.2.1 -DSYCL2017 %s -// RUN: %clangxx -fsycl -fsyntax-only -sycl-std=2020 -DSYCL2020 %s - -#if defined(SYCL2017) -#if !CL_SYCL_LANGUAGE_VERSION -static_assert(false, "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used"); -#endif -#if CL_SYCL_LANGUAGE_VERSION > 121 -static_assert(false, "-sycl-std flag should have set CL_SYCL_LANGUAGE_VERSION to 121"); -#endif -#endif - -#if defined(SYCL2020) -#if !CL_SYCL_LANGUAGE_VERSION -static_assert(false, "CL_SYCL_LANGUAGE_VERSION should be defined when -sycl-std flag is used"); -#endif -#if CL_SYCL_LANGUAGE_VERSION < 2020 -static_assert(false, "-sycl-std=2020 flag should set CL_SYCL_LANGUAGE_VERSION to 2020"); -#endif -#endif - -// expected-no-diagnostics \ No newline at end of file diff --git a/clang/test/Preprocessor/sycl-macro.cpp b/clang/test/Preprocessor/sycl-macro.cpp index 21f94e9fac2df..24145987637b2 100644 --- a/clang/test/Preprocessor/sycl-macro.cpp +++ b/clang/test/Preprocessor/sycl-macro.cpp @@ -1,11 +1,14 @@ // 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 \ // RUN: --check-prefix=CHECK-NO-SYCL_FIT_IN_INT %s + + // CHECK-NOT:#define __SYCL_DEVICE_ONLY__ 1 // CHECK-NOT:#define SYCL_EXTERNAL // CHECK-NOT:#define CL_SYCL_LANGUAGE_VERSION 121 @@ -14,6 +17,8 @@ // CHECK-SYCL-STD:#define CL_SYCL_LANGUAGE_VERSION 121 // CHECK-SYCL-STD:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1 +// CHECK-SYCL-STD-2020:#define CL_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/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp index 022d34ad59227..b80dbc14ce380 100644 --- a/clang/test/SemaSYCL/implicit_kernel_type.cpp +++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp @@ -37,7 +37,6 @@ class handler { // expected-no-diagnostics #endif -//using namespace cl::sycl; void function() { } diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 737afb7d0327d..2c9feb2b5eba5 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -947,6 +947,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. /// @@ -963,11 +971,7 @@ class __SYCL_EXPORT handler { template void parallel_for(range NumWorkItems, id WorkItemOffset, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1000,11 +1004,7 @@ class __SYCL_EXPORT handler { template void parallel_for(nd_range ExecutionRange, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1032,11 +1032,7 @@ class __SYCL_EXPORT handler { detail::enable_if_t parallel_for(nd_range Range, Reduction Redu, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { intel::detail::reduCGFunc(*this, KernelFunc, Range, Redu, Redu.getUserAccessor()); } @@ -1050,11 +1046,7 @@ class __SYCL_EXPORT handler { detail::enable_if_t parallel_for(nd_range Range, Reduction Redu, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { intel::detail::reduCGFunc(*this, KernelFunc, Range, Redu, Redu.getUSMPointer()); } @@ -1074,11 +1066,7 @@ class __SYCL_EXPORT handler { detail::enable_if_t parallel_for(nd_range Range, Reduction Redu, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { shared_ptr_class QueueCopy = MQueue; auto RWAcc = Redu.getReadWriteScalarAcc(*this); intel::detail::reduCGFunc(*this, KernelFunc, Range, Redu, @@ -1115,11 +1103,7 @@ class __SYCL_EXPORT handler { int Dims, typename Reduction> detail::enable_if_t parallel_for(nd_range Range, Reduction Redu, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _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 @@ -1194,11 +1178,7 @@ class __SYCL_EXPORT handler { template void parallel_for_work_group(range NumWorkGroups, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1231,11 +1211,7 @@ class __SYCL_EXPORT handler { int Dims> void parallel_for_work_group(range NumWorkGroups, range WorkGroupSize, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1337,12 +1313,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, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + void single_task(kernel Kernel, _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1382,11 +1353,7 @@ class __SYCL_EXPORT handler { template void parallel_for(kernel Kernel, range NumWorkItems, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1421,12 +1388,7 @@ class __SYCL_EXPORT handler { template void parallel_for(kernel Kernel, range NumWorkItems, - id WorkItemOffset, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + id WorkItemOffset, _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1463,11 +1425,7 @@ class __SYCL_EXPORT handler { template void parallel_for(kernel Kernel, nd_range NDRange, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1509,11 +1467,7 @@ class __SYCL_EXPORT handler { template void parallel_for_work_group(kernel Kernel, range NumWorkGroups, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1551,11 +1505,7 @@ class __SYCL_EXPORT handler { int Dims> void parallel_for_work_group(kernel Kernel, range NumWorkGroups, range WorkGroupSize, -#ifdef __SYCL_NONCONST_FUNCTOR__ - KernelType KernelFunc) { -#else - const KernelType &KernelFunc) { -#endif + _KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; @@ -1579,6 +1529,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 adbc33533514b..2c6ffdfdfed74 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -362,33 +362,33 @@ class __SYCL_EXPORT queue { } // 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, + // 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 + // replace _CODELOCPARAM(&CodeLoc) with nothing // or : , const detail::code_location &CodeLoc = // detail::code_location::current() - // replace CODELOCARG(&CodeLoc) with nothing + // replace _CODELOCARG(&CodeLoc) with nothing // or : const detail::code_location &CodeLoc = {} #ifndef DISABLE_SYCL_INSTRUMENTATION_METADATA -#define CODELOCPARAM(a) \ +#define _CODELOCPARAM(a) \ , const detail::code_location a = detail::code_location::current() -#define CODELOCARG(a) +#define _CODELOCARG(a) #else -#define CODELOCPARAM(a) +#define _CODELOCPARAM(a) -#define CODELOCARG(a) const detail::code_location a = {} +#define _CODELOCARG(a) const detail::code_location a = {} #endif -// replace KERNELFUNCPARAM(KernelFunc) with KernelType KernelFunc +// replace _KERNELFUNCPARAM(KernelFunc) with KernelType KernelFunc // or const KernelType &KernelFunc #ifdef __SYCL_NONCONST_FUNCTOR__ -#define KERNELFUNCPARAM(a) KernelType a +#define _KERNELFUNCPARAM(a) KernelType a #else -#define KERNELFUNCPARAM(a) const KernelType &a +#define _KERNELFUNCPARAM(a) const KernelType &a #endif /// single_task version with a kernel represented as a lambda. @@ -396,8 +396,8 @@ 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(KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + event single_task(_KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { @@ -413,8 +413,8 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event single_task(event DepEvent, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvent); @@ -431,8 +431,8 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event single_task(const vector_class &DepEvents, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvents); @@ -449,8 +449,8 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for(range<1> NumWorkItems, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, KernelFunc, CodeLoc); } @@ -462,8 +462,8 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for(range<2> NumWorkItems, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, KernelFunc, CodeLoc); } @@ -475,8 +475,8 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for(range<3> NumWorkItems, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, KernelFunc, CodeLoc); } @@ -489,8 +489,8 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for(range<1> NumWorkItems, event DepEvent, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvent, KernelFunc, CodeLoc); } @@ -504,8 +504,8 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for(range<2> NumWorkItems, event DepEvent, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvent, KernelFunc, CodeLoc); } @@ -519,8 +519,8 @@ class __SYCL_EXPORT queue { /// \param CodeLoc contains the code location of user code template event parallel_for(range<3> NumWorkItems, event DepEvent, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvent, KernelFunc, CodeLoc); } @@ -536,8 +536,8 @@ class __SYCL_EXPORT queue { template event parallel_for(range<1> NumWorkItems, const vector_class &DepEvents, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvents, KernelFunc, CodeLoc); } @@ -553,8 +553,8 @@ class __SYCL_EXPORT queue { template event parallel_for(range<2> NumWorkItems, const vector_class &DepEvents, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvents, KernelFunc, CodeLoc); } @@ -570,8 +570,8 @@ class __SYCL_EXPORT queue { template event parallel_for(range<3> NumWorkItems, const vector_class &DepEvents, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return parallel_for_impl(NumWorkItems, DepEvents, KernelFunc, CodeLoc); } @@ -586,8 +586,8 @@ class __SYCL_EXPORT queue { template event parallel_for(range NumWorkItems, id WorkItemOffset, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.template parallel_for( @@ -608,8 +608,8 @@ class __SYCL_EXPORT queue { int Dims> event parallel_for(range NumWorkItems, id WorkItemOffset, event DepEvent, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvent); @@ -632,8 +632,8 @@ class __SYCL_EXPORT queue { int Dims> event parallel_for(range NumWorkItems, id WorkItemOffset, const vector_class &DepEvents, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvents); @@ -653,8 +653,8 @@ class __SYCL_EXPORT queue { template event parallel_for(nd_range ExecutionRange, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.template parallel_for(ExecutionRange, @@ -674,8 +674,8 @@ class __SYCL_EXPORT queue { template event parallel_for(nd_range ExecutionRange, event DepEvent, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvent); @@ -698,8 +698,8 @@ class __SYCL_EXPORT queue { int Dims> event parallel_for(nd_range ExecutionRange, const vector_class &DepEvents, - KERNELFUNCPARAM(KernelFunc) CODELOCPARAM(&CodeLoc)) { - CODELOCARG(&CodeLoc); + _KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) { + _CODELOCARG(&CodeLoc); return submit( [&](handler &CGH) { CGH.depends_on(DepEvents); @@ -710,9 +710,9 @@ class __SYCL_EXPORT queue { } // Clean up CODELOC and KERNELFUNC macros. -#undef CODELOCPARAM -#undef CODELOCARG -#undef KERNELFUNCPARAM +#undef _CODELOCPARAM +#undef _CODELOCARG +#undef _KERNELFUNCPARAM /// Returns whether the queue is in order or OoO /// From 6bace1ef56d5a84504614025233bf3bf65f219d5 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 12 Aug 2020 14:50:05 -0700 Subject: [PATCH 10/17] overlooked test. Signed-off-by: Chris Perkins --- clang/test/SemaSYCL/accessor-type-diagnostics.cpp | 7 +------ 1 file changed, 1 insertion(+), 6 deletions(-) diff --git a/clang/test/SemaSYCL/accessor-type-diagnostics.cpp b/clang/test/SemaSYCL/accessor-type-diagnostics.cpp index e6b958d130caa..a409e35527e4e 100644 --- a/clang/test/SemaSYCL/accessor-type-diagnostics.cpp +++ b/clang/test/SemaSYCL/accessor-type-diagnostics.cpp @@ -6,11 +6,6 @@ using namespace cl::sycl; -template -__attribute__((sycl_kernel)) void kernel(const 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(); From dc5863d60866bf5622f7ead63722c3a2d730d0e0 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 12 Aug 2020 14:55:29 -0700 Subject: [PATCH 11/17] clang-format is grumpy Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 969ad1c680e59..f200b9cc456c8 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -707,7 +707,7 @@ getKernelInvocationKind(FunctionDecl *KernelCallerFunc) { static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) { assert(Caller->getNumParams() > 0 && "Insufficient kernel parameters"); - + QualType KernelParamTy = (*Caller->param_begin())->getType(); // In SYCL 2020 kernels are now passed by reference. if (KernelParamTy->isReferenceType()) From bbf2f9528f25e3d17e5a170fb3199e9ee2937888 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 12 Aug 2020 16:59:34 -0700 Subject: [PATCH 12/17] move conformance check Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 39 +++++++++++++------------------------ 1 file changed, 13 insertions(+), 26 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f200b9cc456c8..ff7499ba64667 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -708,7 +708,7 @@ getKernelInvocationKind(FunctionDecl *KernelCallerFunc) { static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) { assert(Caller->getNumParams() > 0 && "Insufficient kernel parameters"); - QualType KernelParamTy = (*Caller->param_begin())->getType(); + QualType KernelParamTy = Caller->getParamDecl(0)->getType(); // In SYCL 2020 kernels are now passed by reference. if (KernelParamTy->isReferenceType()) return KernelParamTy->getPointeeCXXRecordDecl(); @@ -717,30 +717,6 @@ static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) { return KernelParamTy->getAsCXXRecordDecl(); } -static void checkKernelAndCaller(Sema &SemaRef, FunctionDecl *Caller, - const CXXRecordDecl *KernelObj) { - // check captures - if (KernelObj->isLambda()) { - for (const LambdaCapture &LC : KernelObj->captures()) - if (LC.capturesThis() && LC.isImplicit()) - SemaRef.Diag(LC.getLocation(), diag::err_implicit_this_capture); - } - - // check that calling kernel conforms to spec - assert(Caller->param_size() >= 1 && "missing kernel function argument."); - QualType KernelParamTy = (*Caller->param_begin())->getType(); - if (KernelParamTy->isReferenceType()) { - // passing by reference, so emit warning if not using SYCL 2020 - if (SemaRef.LangOpts.SYCLVersion < 2020) - SemaRef.Diag(Caller->getLocation(), - diag::warn_sycl_pass_by_reference_future); - } else { - // passing by value. emit warning if using SYCL 2020 or greater - if (SemaRef.LangOpts.SYCLVersion > 2017) - SemaRef.Diag(Caller->getLocation(), - diag::warn_sycl_pass_by_value_deprecated); - } -} /// Creates a kernel parameter descriptor /// \param Src field declaration to construct name from @@ -2280,6 +2256,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); @@ -2316,7 +2304,6 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, // The first argument to the KernelCallerFunc is the lambda object. const CXXRecordDecl *KernelObj = getKernelObjectType(KernelCallerFunc); assert(KernelObj && "invalid kernel caller"); - checkKernelAndCaller(*this, KernelCallerFunc, KernelObj); // Calculate both names, since Integration headers need both. std::string CalculatedName, StableName; From d84c1e9c0300fef58059fd6e64b49d15964f27d0 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 12 Aug 2020 17:04:06 -0700 Subject: [PATCH 13/17] hmm - git clang-format is skipping certain whitespace changes. Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ff7499ba64667..ac0ea7b125185 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -717,7 +717,6 @@ static const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) { return KernelParamTy->getAsCXXRecordDecl(); } - /// Creates a kernel parameter descriptor /// \param Src field declaration to construct name from /// \param Ty the desired parameter type From 0f9fe01403c850ebe318849d9632f77ef84c4513 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 13 Aug 2020 08:52:44 -0700 Subject: [PATCH 14/17] removed unused var. stray from earlier change Signed-off-by: Chris Perkins --- clang/lib/Driver/ToolChains/Clang.cpp | 5 ----- 1 file changed, 5 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 4f100cb68eb38..61ee698ffbd9e 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -4160,11 +4160,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, if (SYCLStdArg) { SYCLStdArg->render(Args, CmdArgs); CmdArgs.push_back("-fsycl-std-layout-kernel-params"); - - unsigned version = llvm::StringSwitch(SYCLStdArg->getValue()) - .Cases("2017", "1.2.1", "121", "sycl-1.2.1", 2017) - .Case("2020", 2020) - .Default(0); } else { // Ensure the default version in SYCL mode is 1.2.1 (aka 2017) CmdArgs.push_back("-sycl-std=2017"); From b87d642072b5f9af22f3a2ec34205678a6bddb1e Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 13 Aug 2020 09:32:16 -0700 Subject: [PATCH 15/17] latest review requests Signed-off-by: Chris Perkins --- sycl/include/CL/sycl/handler.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 2c9feb2b5eba5..084f1a4b2641d 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -33,7 +33,8 @@ #include #include -#if !CL_SYCL_LANGUAGE_VERSION || CL_SYCL_LANGUAGE_VERSION < 2020 +// CL_SYCL_LANGUAGE_VERSION is 4 digit year followed by 2 digit revision +#if !CL_SYCL_LANGUAGE_VERSION || CL_SYCL_LANGUAGE_VERSION < 202001 #define __SYCL_NONCONST_FUNCTOR__ #endif From 728dd4558c9921f20b514542b0ec6505afc427eb Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 13 Aug 2020 11:00:57 -0700 Subject: [PATCH 16/17] clang-format, of course Signed-off-by: Chris Perkins --- clang/test/Preprocessor/sycl-macro.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/clang/test/Preprocessor/sycl-macro.cpp b/clang/test/Preprocessor/sycl-macro.cpp index 24145987637b2..e80c1fe1d3bfa 100644 --- a/clang/test/Preprocessor/sycl-macro.cpp +++ b/clang/test/Preprocessor/sycl-macro.cpp @@ -7,8 +7,6 @@ // RUN: %clang_cc1 -fno-sycl-id-queries-fit-in-int %s -E -dM | FileCheck \ // RUN: --check-prefix=CHECK-NO-SYCL_FIT_IN_INT %s - - // CHECK-NOT:#define __SYCL_DEVICE_ONLY__ 1 // CHECK-NOT:#define SYCL_EXTERNAL // CHECK-NOT:#define CL_SYCL_LANGUAGE_VERSION 121 From 391ba54e16dfa2afe589b2ab2043edcb0f53778d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 14 Aug 2020 09:35:18 -0700 Subject: [PATCH 17/17] new SYCL_LANGUAGE_VERSION macro and etc. Signed-off-by: Chris Perkins --- clang/lib/Frontend/InitPreprocessor.cpp | 8 +++++--- clang/test/Preprocessor/sycl-macro.cpp | 3 ++- sycl/include/CL/sycl/handler.hpp | 6 +++--- sycl/test/basic_tests/macros.cpp | 3 +-- sycl/test/basic_tests/set_arg_error.cpp | 2 +- 5 files changed, 12 insertions(+), 10 deletions(-) diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index a8c78a5846c58..b48917f7b1eb2 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -463,10 +463,12 @@ 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"); - else if (LangOpts.SYCLVersion == 2020) - Builder.defineMacro("CL_SYCL_LANGUAGE_VERSION", "202001"); + 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/test/Preprocessor/sycl-macro.cpp b/clang/test/Preprocessor/sycl-macro.cpp index e80c1fe1d3bfa..8a0be0841da47 100644 --- a/clang/test/Preprocessor/sycl-macro.cpp +++ b/clang/test/Preprocessor/sycl-macro.cpp @@ -13,9 +13,10 @@ // 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 CL_SYCL_LANGUAGE_VERSION 202001 +// 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/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index 084f1a4b2641d..03b3d5eb9d74f 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -33,8 +33,8 @@ #include #include -// CL_SYCL_LANGUAGE_VERSION is 4 digit year followed by 2 digit revision -#if !CL_SYCL_LANGUAGE_VERSION || CL_SYCL_LANGUAGE_VERSION < 202001 +// 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 @@ -812,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 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);