diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index ca24e1a81839d..afcd9ce41dc31 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -13172,9 +13172,8 @@ def err_sycl_num_kernel_wrong_reqd_wg_size : Error< def err_sycl_invalid_aspect_argument : Error< "%0 attribute argument is invalid; argument must be device aspect of type sycl::aspect">; -def warn_sycl_pass_by_value_deprecated - : Warning<"passing kernel functions by value is deprecated in SYCL 2020">, - InGroup, ShowInSystemHeader; +def err_sycl_kernel_pass_by_value + : Error<"SYCL kernel function must be passed by reference">; def warn_sycl_potentially_invalid_as_cast : Warning< "explicit cast from %0 to %1 potentially leads to an invalid address space" " cast in the resulting code">, InGroup, diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 07a6c25dced91..5a05e06ba48d0 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1153,17 +1153,15 @@ getKernelInvocationKind(FunctionDecl *KernelCallerFunc) { // The SYCL kernel's 'object type' used for diagnostics and naming/mangling is // the first parameter to a function template using the sycl_kernel -// attribute. In SYCL 1.2.1, this was passed by value, -// and in SYCL 2020, it is passed by reference. +// attribute. static QualType GetSYCLKernelObjectType(const FunctionDecl *KernelCaller) { assert(KernelCaller->getNumParams() > 0 && "Insufficient kernel parameters"); QualType KernelParamTy = KernelCaller->getParamDecl(0)->getType(); - // SYCL 2020 kernels are passed by reference. + // The kernel parameter is supposed to be a reference, but diagnostics are + // done at a later stage, so we accept a by-value kernel function type here. if (KernelParamTy->isReferenceType()) KernelParamTy = KernelParamTy->getPointeeType(); - - // SYCL 1.2.1 return KernelParamTy.getUnqualifiedType(); } @@ -5243,13 +5241,9 @@ void SemaSYCL::CheckSYCLKernelCall(FunctionDecl *KernelFunc, } } - // check that calling kernel conforms to spec - QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType(); - if (not KernelParamTy->isReferenceType()) { - // passing by value. emit warning if using SYCL 2020 or greater - if (SemaRef.LangOpts.getSYCLVersion() >= LangOptions::SYCL_2020) - Diag(KernelFunc->getLocation(), diag::warn_sycl_pass_by_value_deprecated); - } + // SYCL only supports passing kernel functions by reference. + if (!KernelFunc->getParamDecl(0)->getType()->isReferenceType()) + Diag(KernelFunc->getLocation(), diag::err_sycl_kernel_pass_by_value); // Do not visit invalid kernel object. if (KernelObj->isInvalidDecl()) @@ -5396,8 +5390,8 @@ void SemaSYCL::SetSYCLKernelNames() { // // Example of kernel caller function: // template -// __attribute__((sycl_kernel)) void kernel_caller_function(KernelType -// KernelFuncObj) { +// __attribute__((sycl_kernel)) +// void kernel_caller_function(const KernelType &KernelFuncObj) { // KernelFuncObj(); // } // diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 4245fb0f658c9..c6b8811eac7c1 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -558,11 +558,6 @@ ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_ha kernelFunc(kh); } -template -ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) { // #KernelSingleTask2017 - kernelFunc(); -} - template ATTR_SYCL_KERNEL void kernel_parallel_for(const KernelType &KernelFunc) { @@ -624,16 +619,6 @@ class handler { kernel_single_task(kernelFunc, kh); #else kernelFunc(kh); -#endif - } - - template - void single_task_2017(KernelType kernelFunc) { - using NameT = typename get_kernel_name_t::name; -#ifdef __SYCL_DEVICE_ONLY__ - kernel_single_task_2017(kernelFunc); -#else - kernelFunc(); #endif } }; diff --git a/clang/test/CodeGenSYCL/esimd-private-global.cpp b/clang/test/CodeGenSYCL/esimd-private-global.cpp index cd0c8b2e0526c..9e969a12dd187 100644 --- a/clang/test/CodeGenSYCL/esimd-private-global.cpp +++ b/clang/test/CodeGenSYCL/esimd-private-global.cpp @@ -7,7 +7,7 @@ __attribute__((opencl_private)) __attribute__((sycl_explicit_simd)) __attribute_ // CHECK: @vc = {{.+}} i32 0, align 4 #[[ATTR:[0-9]+]] template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/esimd_metadata2.cpp b/clang/test/CodeGenSYCL/esimd_metadata2.cpp index 30eb70de2eb20..6a39e18a0fe42 100644 --- a/clang/test/CodeGenSYCL/esimd_metadata2.cpp +++ b/clang/test/CodeGenSYCL/esimd_metadata2.cpp @@ -10,17 +10,17 @@ __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func // CHECK-ESIMD-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_cm() #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]] {{.*}}{ // CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}esimd_funcv() #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[SGSIZE1]] {{.*}}{ // CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}shared_funcv() #{{[0-9]+}}{{.*}} { -// CHECK-ESIMD-DAG: define linkonce_odr spir_func void @_ZN12ESIMDFunctorclEv({{.*}}) #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}} {{.*}}{ +// CHECK-ESIMD-DAG: define linkonce_odr spir_func void @_ZNK12ESIMDFunctorclEv({{.*}}) #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}} {{.*}}{ class ESIMDFunctor { public: - void operator()() __attribute__((sycl_explicit_simd)) { + void operator()() const __attribute__((sycl_explicit_simd)) { esimd_func(); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/union-kernel-param-ih.cpp b/clang/test/CodeGenSYCL/union-kernel-param-ih.cpp index 4f219c386b528..a689eba0a05e1 100644 --- a/clang/test/CodeGenSYCL/union-kernel-param-ih.cpp +++ b/clang/test/CodeGenSYCL/union-kernel-param-ih.cpp @@ -35,7 +35,7 @@ union MyUnion { }; 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/union-kernel-param.cpp b/clang/test/CodeGenSYCL/union-kernel-param.cpp index 0909e4720935d..c2176a5d9ea5a 100644 --- a/clang/test/CodeGenSYCL/union-kernel-param.cpp +++ b/clang/test/CodeGenSYCL/union-kernel-param.cpp @@ -9,7 +9,7 @@ union MyUnion { }; 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/unique_stable_name.cpp b/clang/test/CodeGenSYCL/unique_stable_name.cpp index 3ab7e3b8f2e7a..e3e5318bf3567 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name.cpp @@ -61,29 +61,29 @@ struct Derp { }; template -[[clang::sycl_kernel]] void kernel_single_task(KernelType kernelFunc) { +[[clang::sycl_kernel]] void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } template -void unnamed_kernel_single_task(KernelType kernelFunc) { +void unnamed_kernel_single_task(const KernelType &kernelFunc) { kernel_single_task(kernelFunc); } template -void not_kernel_single_task(KernelType kernelFunc) { +void not_kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } int main() { not_kernel_single_task(func); - // CHECK: call void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(ptr noundef @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv) + // CHECK: call void @_Z22not_kernel_single_taskIZ4mainE7kernel2FPKcvEEvRKT0_(ptr noundef nonnull @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv) auto l1 = []() { return 1; }; auto l2 = [](decltype(l1) *l = nullptr) { return 2; }; kernel_single_task(l2); puts(__builtin_sycl_unique_stable_name(decltype(l2))); - // CHECK: call void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_ + // CHECK: call void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvRKT0_ // CHECK: call void @puts(ptr noundef @[[LAMBDA_KERNEL3]]) constexpr const char str[] = "lalala"; @@ -98,10 +98,10 @@ int main() { puts(__builtin_sycl_unique_stable_name(int[++j])); // CHECK: call void @puts(ptr noundef @[[STRING]]) - // CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_ + // CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2FPKcvEEvRKT0_ // CHECK: declare noundef ptr @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv - // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_ - // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE0_S0_EvT0_ + // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvRKT0_ + // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE0_S0_EvRKT0_ unnamed_kernel_single_task( []() { diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index 2bf55b5dd1eb3..723e80fb6a2df 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -3,32 +3,32 @@ template -__attribute__((sycl_kernel)) void kernel(Func F){ +__attribute__((sycl_kernel)) void kernel(const Func &F){ F(); } template -void kernel_wrapper(Func F) { +void kernel_wrapper(const Func &F) { kernel(F); } template -__attribute__((sycl_kernel)) void kernel2(Func F){ +__attribute__((sycl_kernel)) void kernel2(const Func &F){ F(1); } template -void kernel2_wrapper(Func F) { +void kernel2_wrapper(const Func &F) { kernel2(F); } template -__attribute__((sycl_kernel)) void kernel3(Func F){ +__attribute__((sycl_kernel)) void kernel3(const Func &F){ F(1.1); } template -void kernel3_wrapper(Func F) { +void kernel3_wrapper(const Func &F) { kernel3(F); } diff --git a/clang/test/Frontend/Inputs/sycl.hpp b/clang/test/Frontend/Inputs/sycl.hpp index 1e5b4bd124e92..77d04e1bf0f85 100644 --- a/clang/test/Frontend/Inputs/sycl.hpp +++ b/clang/test/Frontend/Inputs/sycl.hpp @@ -400,11 +400,6 @@ ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_ha kernelFunc(kh); } -template -ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) { // #KernelSingleTask2017 - kernelFunc(); -} - template ATTR_SYCL_KERNEL void kernel_parallel_for(const KernelType &KernelFunc) { @@ -466,16 +461,6 @@ class handler { kernel_single_task(kernelFunc, kh); #else kernelFunc(kh); -#endif - } - - template - void single_task_2017(KernelType kernelFunc) { - using NameT = typename get_kernel_name_t::name; -#ifdef __SYCL_DEVICE_ONLY__ - kernel_single_task_2017(kernelFunc); -#else - kernelFunc(); #endif } }; diff --git a/clang/test/SemaSYCL/kernel-attribute-on-non-sycl.cpp b/clang/test/SemaSYCL/kernel-attribute-on-non-sycl.cpp index 2976283dbac1c..753fb528ddd74 100644 --- a/clang/test/SemaSYCL/kernel-attribute-on-non-sycl.cpp +++ b/clang/test/SemaSYCL/kernel-attribute-on-non-sycl.cpp @@ -9,6 +9,6 @@ #endif template -__attribute__((sycl_kernel)) void foo(T P); +__attribute__((sycl_kernel)) void foo(const T& P); template -[[clang::sycl_kernel]] void foo1(T P); +[[clang::sycl_kernel]] void foo1(const T& P); diff --git a/clang/test/SemaSYCL/kernel-attribute.cpp b/clang/test/SemaSYCL/kernel-attribute.cpp index bb0cb8f14b40e..7623b0c504dbd 100644 --- a/clang/test/SemaSYCL/kernel-attribute.cpp +++ b/clang/test/SemaSYCL/kernel-attribute.cpp @@ -11,43 +11,43 @@ __attribute__((sycl_kernel)) void foo(); // expected-warning {{'sycl_kernel' att // Attribute takes no arguments template -__attribute__((sycl_kernel(1))) void foo(T P); // expected-error {{'sycl_kernel' attribute takes no arguments}} +__attribute__((sycl_kernel(1))) void foo(const T &P); // expected-error {{'sycl_kernel' attribute takes no arguments}} template -[[clang::sycl_kernel(1)]] void foo1(T P);// expected-error {{'sycl_kernel' attribute takes no arguments}} +[[clang::sycl_kernel(1)]] void foo1(const T &P);// expected-error {{'sycl_kernel' attribute takes no arguments}} // At least two template parameters template -__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}} +__attribute__((sycl_kernel)) void foo(const T &P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}} template -[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}} +[[clang::sycl_kernel]] void foo1(const T &P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}} // First two template parameters cannot be non-type template parameters template -__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}} +__attribute__((sycl_kernel)) void foo(const T &P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}} template -[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}} +[[clang::sycl_kernel]] void foo1(const T &P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}} // Must return void template -__attribute__((sycl_kernel)) int foo(T P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}} +__attribute__((sycl_kernel)) int foo(const T &P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}} template -[[clang::sycl_kernel]] int foo1(T P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}} +[[clang::sycl_kernel]] int foo1(const T &P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}} // Must take at least one argument template __attribute__((sycl_kernel)) void foo(); // expected-warning {{function template with 'sycl_kernel' attribute must have at least one parameter}} template -[[clang::sycl_kernel]] void foo1(T t, A a); // no diagnostics +[[clang::sycl_kernel]] void foo1(const T &t, A a); // no diagnostics // No diagnostics template -__attribute__((sycl_kernel)) void foo(T P); +__attribute__((sycl_kernel)) void foo(const T &P); template -[[clang::sycl_kernel]] void foo1(T P); +[[clang::sycl_kernel]] void foo1(const T &P); #else // expected-no-diagnostics template -__attribute__((sycl_kernel)) void foo(T P); +__attribute__((sycl_kernel)) void foo(const T &P); #endif diff --git a/clang/test/SemaSYCL/kernel-by-val-error.cpp b/clang/test/SemaSYCL/kernel-by-val-error.cpp new file mode 100644 index 0000000000000..4e8dd8fd00401 --- /dev/null +++ b/clang/test/SemaSYCL/kernel-by-val-error.cpp @@ -0,0 +1,23 @@ + +// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fsycl-is-device -verify %s +// RUN: %clang_cc1 -fsycl-is-host -DHOST -fsyntax-only -verify %s + +// Kernel function argument must be passed by reference + +#ifndef HOST +template +__attribute__((sycl_kernel)) void foo(T P) {} // expected-error {{SYCL kernel function must be passed by reference}} +template +[[clang::sycl_kernel]] void bar(T P) {} // expected-error {{SYCL kernel function must be passed by reference}} +#else +// expected-no-diagnostics +template +__attribute__((sycl_kernel)) void foo(T P) {} +template +[[clang::sycl_kernel]] void bar(T P) {} +#endif + +void F() { + foo([](){}); + bar([](){}); +} \ No newline at end of file diff --git a/clang/test/SemaSYCL/kernel-not-functor.cpp b/clang/test/SemaSYCL/kernel-not-functor.cpp index 67f50d2679a89..4f704b286bae2 100644 --- a/clang/test/SemaSYCL/kernel-not-functor.cpp +++ b/clang/test/SemaSYCL/kernel-not-functor.cpp @@ -3,12 +3,12 @@ // RUNX: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s template -__attribute__((sycl_kernel)) void kernel(F kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const F &kernelFunc) { kernelFunc(); } template -void uses_kernel(F kernelFunc) { +void uses_kernel(const F &kernelFunc) { // expected-error@+1{{kernel parameter must be a lambda or function object}} kernel(kernelFunc); } diff --git a/clang/test/SemaSYCL/union-kernel-param.cpp b/clang/test/SemaSYCL/union-kernel-param.cpp index 4fbc4aaa56502..344aeb53c0693 100644 --- a/clang/test/SemaSYCL/union-kernel-param.cpp +++ b/clang/test/SemaSYCL/union-kernel-param.cpp @@ -10,7 +10,7 @@ union MyUnion { }; 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/union-kernel-param1.cpp b/clang/test/SemaSYCL/union-kernel-param1.cpp index c3567fb40cc8a..d24615ebc7421 100644 --- a/clang/test/SemaSYCL/union-kernel-param1.cpp +++ b/clang/test/SemaSYCL/union-kernel-param1.cpp @@ -4,7 +4,7 @@ // unions containing Arrays. 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/unique-stable-name-multiple-target-crash.cpp b/clang/test/SemaSYCL/unique-stable-name-multiple-target-crash.cpp index c169010be537a..5328c3ce907e2 100644 --- a/clang/test/SemaSYCL/unique-stable-name-multiple-target-crash.cpp +++ b/clang/test/SemaSYCL/unique-stable-name-multiple-target-crash.cpp @@ -6,7 +6,8 @@ // // expected-no-diagnostics class Empty {}; -template __attribute__((sycl_kernel)) void kernel(F) { +template __attribute__((sycl_kernel)) +void kernel(const F&) { __builtin_sycl_unique_stable_name(F); } diff --git a/clang/test/SemaSYCL/unique_stable_name.cpp b/clang/test/SemaSYCL/unique_stable_name.cpp index cc35c85fa5f22..3a14ef6094501 100644 --- a/clang/test/SemaSYCL/unique_stable_name.cpp +++ b/clang/test/SemaSYCL/unique_stable_name.cpp @@ -2,7 +2,7 @@ // RUN: %clang_cc1 %s -std=c++17 -triple x86_64-linux-gnu -Wno-sycl-2020-compat -fsycl-is-device -verify -fsyntax-only -Wno-unused template -[[clang::sycl_kernel]] void kernel_single_task(KernelType kernelFunc) { // #kernelSingleTask +[[clang::sycl_kernel]] void kernel_single_task(const KernelType &kernelFunc) { // #kernelSingleTask kernelFunc(); } diff --git a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_fixed_point/capability-arbitrary-precision-fixed-point-numbers.ll b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_fixed_point/capability-arbitrary-precision-fixed-point-numbers.ll index 3825ec232848d..c9ffa246a55b4 100644 --- a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_fixed_point/capability-arbitrary-precision-fixed-point-numbers.ll +++ b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_fixed_point/capability-arbitrary-precision-fixed-point-numbers.ll @@ -70,7 +70,7 @@ ; } ; template -; __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +; __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { ; kernelFunc(); ; } diff --git a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_floating_point/capability-arbitrary-precision-floating-point.ll b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_floating_point/capability-arbitrary-precision-floating-point.ll index 344aa26c5755c..bca8da2d3724a 100644 --- a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_floating_point/capability-arbitrary-precision-floating-point.ll +++ b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_floating_point/capability-arbitrary-precision-floating-point.ll @@ -349,7 +349,7 @@ ; } ; ; template -; __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +; __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { ; kernelFunc(); ; } ; diff --git a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttr.ll b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttr.ll index 0f822fb41c2e1..e8b6bf2bc1aef 100644 --- a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttr.ll +++ b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttr.ll @@ -54,7 +54,7 @@ ; } ; ; template -; __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +; __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { ; kernelFunc(); ; } ; diff --git a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttrOnClosure.ll b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttrOnClosure.ll index 9e8883e84f105..18abee701aeca 100644 --- a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttrOnClosure.ll +++ b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttrOnClosure.ll @@ -4,7 +4,7 @@ ; ; template ; __attribute__((sycl_kernel)) -; void kernel(Func f) { +; void kernel(const Func &f) { ; f(); ; } ; diff --git a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_accesses/IntelFPGAMemoryAccesses.ll b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_accesses/IntelFPGAMemoryAccesses.ll index 7405618de691e..be274c22cffbf 100644 --- a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_accesses/IntelFPGAMemoryAccesses.ll +++ b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_accesses/IntelFPGAMemoryAccesses.ll @@ -27,7 +27,7 @@ ; } ; ; template -; __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +; __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { ; kernelFunc(); ; } ; diff --git a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_attributes/IntelFPGAMemoryAttributesForStaticVar.ll b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_attributes/IntelFPGAMemoryAttributesForStaticVar.ll index b6d67a7e0377a..27e8120605cd1 100644 --- a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_attributes/IntelFPGAMemoryAttributesForStaticVar.ll +++ b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_attributes/IntelFPGAMemoryAttributesForStaticVar.ll @@ -22,7 +22,7 @@ ; } ; ; template -; __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { +; __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { ; kernelFunc(); ; } ; diff --git a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_kernel_attributes/intel_fpga_function_attributes.ll b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_kernel_attributes/intel_fpga_function_attributes.ll index b8e56a78edadd..1ae7bddc94d27 100644 --- a/llvm-spirv/test/extensions/INTEL/SPV_INTEL_kernel_attributes/intel_fpga_function_attributes.ll +++ b/llvm-spirv/test/extensions/INTEL/SPV_INTEL_kernel_attributes/intel_fpga_function_attributes.ll @@ -14,7 +14,7 @@ ;; }; ;; ;; template -;; __attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +;; __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { ;; kernelFunc(); ;; } ;; diff --git a/sycl/doc/design/CompilerAndRuntimeDesign.md b/sycl/doc/design/CompilerAndRuntimeDesign.md index 01af23623ca2b..15ac12a8ff2f8 100644 --- a/sycl/doc/design/CompilerAndRuntimeDesign.md +++ b/sycl/doc/design/CompilerAndRuntimeDesign.md @@ -146,7 +146,7 @@ previous code snippet above looks like this: ```C++ // SYCL kernel is defined in SYCL headers: template -__attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) { +__attribute__((sycl_kernel)) void sycl_kernel_function(const KernelType &KernelFuncObj) { // ... KernelFuncObj(); } diff --git a/sycl/test/check_device_code/esimd/NBarrierAttr.cpp b/sycl/test/check_device_code/esimd/NBarrierAttr.cpp index 174b870c9b965..064ee05b95559 100644 --- a/sycl/test/check_device_code/esimd/NBarrierAttr.cpp +++ b/sycl/test/check_device_code/esimd/NBarrierAttr.cpp @@ -12,7 +12,7 @@ using namespace sycl::ext::intel::esimd; using namespace sycl::ext::intel::experimental::esimd; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/sycl/test/check_device_code/esimd/dae.cpp b/sycl/test/check_device_code/esimd/dae.cpp index 222a8628647a8..1293e365a6f9a 100644 --- a/sycl/test/check_device_code/esimd/dae.cpp +++ b/sycl/test/check_device_code/esimd/dae.cpp @@ -9,7 +9,7 @@ using namespace sycl; template -__attribute__((sycl_kernel)) void my_kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void my_kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/sycl/test/check_device_code/esimd/dpas.cpp b/sycl/test/check_device_code/esimd/dpas.cpp index 7e2c2682d3988..e3a69717972c5 100644 --- a/sycl/test/check_device_code/esimd/dpas.cpp +++ b/sycl/test/check_device_code/esimd/dpas.cpp @@ -24,11 +24,11 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void xmx_func_end(); class EsimdFunctor { public: - void operator()() __attribute__((sycl_explicit_simd)) { xmx_func(); } + void operator()() const __attribute__((sycl_explicit_simd)) { xmx_func(); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/sycl/test/check_device_code/esimd/fp16_converts.cpp b/sycl/test/check_device_code/esimd/fp16_converts.cpp index 87ef8c54e952b..e122079bc3038 100644 --- a/sycl/test/check_device_code/esimd/fp16_converts.cpp +++ b/sycl/test/check_device_code/esimd/fp16_converts.cpp @@ -22,14 +22,14 @@ using bfloat16 = sycl::ext::oneapi::bfloat16; class EsimdFunctor { public: - void operator()() __attribute__((sycl_explicit_simd)) { + void operator()() const __attribute__((sycl_explicit_simd)) { bf16_vector(); bf16_scalar(); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/sycl/test/check_device_code/esimd/genx_func_attr.cpp b/sycl/test/check_device_code/esimd/genx_func_attr.cpp index f92c5f40c36a3..85f05604caff2 100644 --- a/sycl/test/check_device_code/esimd/genx_func_attr.cpp +++ b/sycl/test/check_device_code/esimd/genx_func_attr.cpp @@ -12,7 +12,7 @@ using namespace sycl::ext::intel::esimd; using namespace sycl::ext::intel::experimental::esimd; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/sycl/test/check_device_code/esimd/glob.cpp b/sycl/test/check_device_code/esimd/glob.cpp index 2f005a7c71340..5ebed0da9b40b 100644 --- a/sycl/test/check_device_code/esimd/glob.cpp +++ b/sycl/test/check_device_code/esimd/glob.cpp @@ -25,7 +25,7 @@ ESIMD_PRIVATE ESIMD_REGISTER(17 + VL) simd vc1; // "VCVolatile"{{.*}} } template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/sycl/test/check_device_code/esimd/intrins_trans.cpp b/sycl/test/check_device_code/esimd/intrins_trans.cpp index 35010dd1c3471..49c2835a069b9 100644 --- a/sycl/test/check_device_code/esimd/intrins_trans.cpp +++ b/sycl/test/check_device_code/esimd/intrins_trans.cpp @@ -28,11 +28,11 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd foo(); class EsimdFunctor { public: - void operator()() __attribute__((sycl_explicit_simd)) { foo(); } + void operator()() const __attribute__((sycl_explicit_simd)) { foo(); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } diff --git a/sycl/test/check_device_code/esimd/lsc.cpp b/sycl/test/check_device_code/esimd/lsc.cpp index 7e2255bdac453..d1ae6effff15e 100644 --- a/sycl/test/check_device_code/esimd/lsc.cpp +++ b/sycl/test/check_device_code/esimd/lsc.cpp @@ -21,25 +21,25 @@ using namespace sycl::ext::intel::experimental::esimd; using AccType = sycl::accessor; -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void foo(AccType &); +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void foo(const AccType &); class EsimdFunctor { public: AccType acc; - void operator()() __attribute__((sycl_explicit_simd)) { foo(acc); } + void operator()() const __attribute__((sycl_explicit_simd)) { foo(acc); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } -void bar(AccType &acc) { +void bar(const AccType &acc) { EsimdFunctor esimdf{acc}; kernel(esimdf); } -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void foo(AccType &acc) { +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void foo(const AccType &acc) { constexpr int VL = 4; int *ptr = 0; uintptr_t addr = reinterpret_cast(ptr); diff --git a/sycl/test/check_device_code/esimd/memory_properties_atomic_update.cpp b/sycl/test/check_device_code/esimd/memory_properties_atomic_update.cpp index a0b35444ea5fa..bf3bc3b29d71e 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_atomic_update.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_atomic_update.cpp @@ -23,8 +23,8 @@ using LocalAccType = sycl::local_accessor; using LocalAccTypeInt = sycl::local_accessor; SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void -test_atomic_update(AccType &, LocalAccTypeInt &, float *, int byte_offset32, - size_t byte_offset64); +test_atomic_update(const AccType &, const LocalAccTypeInt &, float *, + int byte_offset32, size_t byte_offset64); class EsimdFunctor { public: @@ -34,18 +34,19 @@ class EsimdFunctor { float *ptr; int byte_offset32; size_t byte_offset64; - void operator()() __attribute__((sycl_explicit_simd)) { + void operator()() const __attribute__((sycl_explicit_simd)) { test_atomic_update(acc, local_acc_int, ptr, byte_offset32, byte_offset64); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } -void bar(AccType &acc, LocalAccType &local_acc, LocalAccTypeInt &local_acc_int, - float *ptr, int byte_offset32, size_t byte_offset64) { +void bar(const AccType &acc, const LocalAccType &local_acc, + const LocalAccTypeInt &local_acc_int, float *ptr, int byte_offset32, + size_t byte_offset64) { EsimdFunctor esimdf{acc, local_acc, local_acc_int, ptr, byte_offset32, byte_offset64}; kernel(esimdf); @@ -53,8 +54,8 @@ void bar(AccType &acc, LocalAccType &local_acc, LocalAccTypeInt &local_acc_int, // CHECK-LABEL: define {{.*}} @_Z18test_atomic_update{{.*}} SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void -test_atomic_update(AccType &acc, LocalAccTypeInt local_acc, float *ptrf, - int byte_offset32, size_t byte_offset64) { +test_atomic_update(const AccType &acc, const LocalAccTypeInt local_acc, + float *ptrf, int byte_offset32, size_t byte_offset64) { constexpr int VL = 4; int *ptr = 0; uintptr_t addr = reinterpret_cast(ptr); diff --git a/sycl/test/check_device_code/esimd/memory_properties_copytocopyfrom.cpp b/sycl/test/check_device_code/esimd/memory_properties_copytocopyfrom.cpp index c1507a13b0f46..e62a16c611323 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_copytocopyfrom.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_copytocopyfrom.cpp @@ -22,15 +22,16 @@ using AccType = sycl::accessor; using LocalAccType = sycl::local_accessor; using LocalAccTypeInt = sycl::local_accessor; -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_copy_to(AccType &, LocalAccType &, +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_copy_to(const AccType &, + const LocalAccType &, float *, int byte_offset32, size_t byte_offset64); -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_copy_from(AccType &, LocalAccType &, - float *, - int byte_offset32, - size_t byte_offset64); -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_ctor(AccType &, LocalAccType &, - float *, int byte_offset32, +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void +test_copy_from(const AccType &, const LocalAccType &, float *, + int byte_offset32, size_t byte_offset64); +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_ctor(const AccType &, + const LocalAccType &, float *, + int byte_offset32, size_t byte_offset64); class EsimdFunctor { public: @@ -40,7 +41,7 @@ class EsimdFunctor { float *ptr; int byte_offset32; size_t byte_offset64; - void operator()() __attribute__((sycl_explicit_simd)) { + void operator()() const __attribute__((sycl_explicit_simd)) { test_copy_to(acc, local_acc, ptr, byte_offset32, byte_offset64); test_copy_from(acc, local_acc, ptr, byte_offset32, byte_offset64); test_ctor(acc, local_acc, ptr, byte_offset32, byte_offset64); @@ -48,7 +49,7 @@ class EsimdFunctor { }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } @@ -61,7 +62,7 @@ void bar(AccType &acc, LocalAccType &local_acc, LocalAccTypeInt &local_acc_int, // CHECK-LABEL: define {{.*}} @_Z12test_copy_to{{.*}} SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void -test_copy_to(AccType &acc, LocalAccType &local_acc, float *ptrf, +test_copy_to(const AccType &acc, const LocalAccType &local_acc, float *ptrf, int byte_offset32, size_t byte_offset64) { properties props_a{cache_hint_L1, cache_hint_L2, alignment<32>}; @@ -87,7 +88,7 @@ test_copy_to(AccType &acc, LocalAccType &local_acc, float *ptrf, // CHECK-LABEL: define {{.*}} @_Z14test_copy_from{{.*}} SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void -test_copy_from(AccType &acc, LocalAccType &local_acc, float *ptrf, +test_copy_from(const AccType &acc, const LocalAccType &local_acc, float *ptrf, int byte_offset32, size_t byte_offset64) { properties props_a{cache_hint_L1, cache_hint_L2, alignment<32>}; @@ -113,8 +114,8 @@ test_copy_from(AccType &acc, LocalAccType &local_acc, float *ptrf, } // CHECK-LABEL: define {{.*}} @_Z9test_ctor{{.*}} -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_ctor(AccType &acc, - LocalAccType &local_acc, +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_ctor(const AccType &acc, + const LocalAccType &local_acc, float *ptrf, int byte_offset32, size_t byte_offset64) { properties props_a{cache_hint_L1, diff --git a/sycl/test/check_device_code/esimd/memory_properties_gather.cpp b/sycl/test/check_device_code/esimd/memory_properties_gather.cpp index 2ebd4749a49c3..9669ceafbd79f 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_gather.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_gather.cpp @@ -21,7 +21,8 @@ using AccType = sycl::accessor; using LocalAccType = sycl::local_accessor; using LocalAccTypeInt = sycl::local_accessor; -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_gather(AccType &, LocalAccType &, +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_gather(const AccType &, + const LocalAccType &, float *, int byte_offset32, size_t byte_offset64); SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_slm_gather(int byte_offset32); @@ -34,14 +35,14 @@ class EsimdFunctor { float *ptr; int byte_offset32; size_t byte_offset64; - void operator()() __attribute__((sycl_explicit_simd)) { + void operator()() const __attribute__((sycl_explicit_simd)) { test_gather(acc, local_acc, ptr, byte_offset32, byte_offset64); test_slm_gather(byte_offset32); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } @@ -54,7 +55,7 @@ void bar(AccType &acc, LocalAccType &local_acc, LocalAccTypeInt &local_acc_int, // CHECK-LABEL: define {{.*}} @_Z11test_gather{{.*}} SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void -test_gather(AccType &acc, LocalAccType &local_acc, float *ptrf, +test_gather(const AccType &acc, const LocalAccType &local_acc, float *ptrf, int byte_offset32, size_t byte_offset64) { properties props_cache_load{cache_hint_L1, cache_hint_L2, diff --git a/sycl/test/check_device_code/esimd/memory_properties_load_store.cpp b/sycl/test/check_device_code/esimd/memory_properties_load_store.cpp index 520d25b324af7..864e81d1aed4a 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_load_store.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_load_store.cpp @@ -22,13 +22,12 @@ using AccType = sycl::accessor; using LocalAccType = sycl::local_accessor; using LocalAccTypeInt = sycl::local_accessor; -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_block_load(AccType &, - LocalAccType &, float *, - int byte_offset32, - size_t byte_offset64); SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void -test_block_store(AccType &, LocalAccType &local_acc, float *, int byte_offset32, - size_t byte_offset64); +test_block_load(const AccType &, const LocalAccType &, float *, + int byte_offset32, size_t byte_offset64); +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void +test_block_store(const AccType &, const LocalAccType &local_acc, float *, + int byte_offset32, size_t byte_offset64); class EsimdFunctor { public: @@ -38,19 +37,20 @@ class EsimdFunctor { float *ptr; int byte_offset32; size_t byte_offset64; - void operator()() __attribute__((sycl_explicit_simd)) { + void operator()() const __attribute__((sycl_explicit_simd)) { test_block_load(acc, local_acc, ptr, byte_offset32, byte_offset64); test_block_store(acc, local_acc, ptr, byte_offset32, byte_offset64); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } -void bar(AccType &acc, LocalAccType &local_acc, LocalAccTypeInt &local_acc_int, - float *ptr, int byte_offset32, size_t byte_offset64) { +void bar(const AccType &acc, const LocalAccType &local_acc, + const LocalAccTypeInt &local_acc_int, float *ptr, int byte_offset32, + size_t byte_offset64) { EsimdFunctor esimdf{acc, local_acc, local_acc_int, ptr, byte_offset32, byte_offset64}; kernel(esimdf); @@ -58,7 +58,7 @@ void bar(AccType &acc, LocalAccType &local_acc, LocalAccTypeInt &local_acc_int, // CHECK-LABEL: define {{.*}} @_Z15test_block_load{{.*}} SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void -test_block_load(AccType &acc, LocalAccType &local_acc, float *ptrf, +test_block_load(const AccType &acc, const LocalAccType &local_acc, float *ptrf, int byte_offset32, size_t byte_offset64) { properties props_a{cache_hint_L1, cache_hint_L2, alignment<16>}; @@ -257,7 +257,7 @@ test_block_load(AccType &acc, LocalAccType &local_acc, float *ptrf, // CHECK-LABEL: define {{.*}} @_Z16test_block_store{{.*}} SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void -test_block_store(AccType &acc, LocalAccType &local_acc, float *ptrf, +test_block_store(const AccType &acc, const LocalAccType &local_acc, float *ptrf, int byte_offset32, size_t byte_offset64) { // Test USM block store constexpr int N = 4; diff --git a/sycl/test/check_device_code/esimd/memory_properties_prefetch_2d.cpp b/sycl/test/check_device_code/esimd/memory_properties_prefetch_2d.cpp index 8638deae418af..bdec6282aaff5 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_prefetch_2d.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_prefetch_2d.cpp @@ -22,8 +22,9 @@ using AccType = sycl::accessor; using LocalAccType = sycl::local_accessor; using LocalAccTypeInt = sycl::local_accessor; -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void -test_prefetch(AccType &, float *, int byte_offset32, size_t byte_offset64); +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_prefetch(const AccType &, float *, + int byte_offset32, + size_t byte_offset64); SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_2d(float *); class EsimdFunctor { @@ -34,26 +35,28 @@ class EsimdFunctor { float *ptr; int byte_offset32; size_t byte_offset64; - void operator()() __attribute__((sycl_explicit_simd)) { + void operator()() const __attribute__((sycl_explicit_simd)) { test_prefetch(acc, ptr, byte_offset32, byte_offset64); test_2d(ptr); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } -void bar(AccType &acc, LocalAccType &local_acc, LocalAccTypeInt &local_acc_int, - float *ptr, int byte_offset32, size_t byte_offset64) { +void bar(const AccType &acc, const LocalAccType &local_acc, + const LocalAccTypeInt &local_acc_int, float *ptr, int byte_offset32, + size_t byte_offset64) { EsimdFunctor esimdf{acc, local_acc, local_acc_int, ptr, byte_offset32, byte_offset64}; kernel(esimdf); } // CHECK-LABEL: define {{.*}} @_Z13test_prefetch{{.*}} -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_prefetch(AccType &acc, float *ptrf, +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_prefetch(const AccType &acc, + float *ptrf, int byte_offset32, size_t byte_offset64) { properties props_cache_load{cache_hint_L1, diff --git a/sycl/test/check_device_code/esimd/memory_properties_scatter.cpp b/sycl/test/check_device_code/esimd/memory_properties_scatter.cpp index 68f1d91078071..843210536c0b9 100644 --- a/sycl/test/check_device_code/esimd/memory_properties_scatter.cpp +++ b/sycl/test/check_device_code/esimd/memory_properties_scatter.cpp @@ -21,7 +21,8 @@ using AccType = sycl::accessor; using LocalAccType = sycl::local_accessor; using LocalAccTypeInt = sycl::local_accessor; -SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_scatter(AccType &, LocalAccType &, +SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_scatter(const AccType &, + const LocalAccType &, float *, int byte_offset32, size_t byte_offset64); SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void test_slm_scatter(int byte_offset32); @@ -34,19 +35,20 @@ class EsimdFunctor { float *ptr; int byte_offset32; size_t byte_offset64; - void operator()() __attribute__((sycl_explicit_simd)) { + void operator()() const __attribute__((sycl_explicit_simd)) { test_scatter(acc, local_acc, ptr, byte_offset32, byte_offset64); test_slm_scatter(byte_offset32); } }; template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } -void bar(AccType &acc, LocalAccType &local_acc, LocalAccTypeInt &local_acc_int, - float *ptr, int byte_offset32, size_t byte_offset64) { +void bar(const AccType &acc, const LocalAccType &local_acc, + const LocalAccTypeInt &local_acc_int, float *ptr, int byte_offset32, + size_t byte_offset64) { EsimdFunctor esimdf{acc, local_acc, local_acc_int, ptr, byte_offset32, byte_offset64}; kernel(esimdf); @@ -54,7 +56,7 @@ void bar(AccType &acc, LocalAccType &local_acc, LocalAccTypeInt &local_acc_int, // CHECK-LABEL: define {{.*}} @_Z12test_scatter{{.*}} SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void -test_scatter(AccType &acc, LocalAccType &local_acc, float *ptrf, +test_scatter(const AccType &acc, const LocalAccType &local_acc, float *ptrf, int byte_offset32, size_t byte_offset64) { properties props_cache_load{cache_hint_L1, cache_hint_L2,