From 904b81a43fec7e06570917ee32fa689abdcba143 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Wed, 7 Apr 2021 23:38:18 -0700 Subject: [PATCH 1/7] [ESIMD] SYCL device function cannot be called from ESIMD context There are two subsequent changes I'm planning to make: 1. Mark all the ESIMD intrinsics with SYCL_ESIMD_FUNCTION. This change will change many lines, so I will submit it as a separate NFC patch. 2. Update the documentation to reflect this change. --- .../clang/Basic/DiagnosticSemaKinds.td | 2 ++ clang/lib/Sema/SemaExpr.cpp | 13 +++++++++ clang/test/CodeGenSYCL/esimd_metadata2.cpp | 4 +-- .../SemaSYCL/esimd-sycl-context-switch.cpp | 27 +++++++++++++++++++ sycl/include/CL/sycl/INTEL/esimd.hpp | 14 +++++----- .../CL/sycl/INTEL/esimd/esimd_memory.hpp | 2 +- sycl/test/esimd/global_var.cpp | 3 --- 7 files changed, 52 insertions(+), 13 deletions(-) create mode 100644 clang/test/SemaSYCL/esimd-sycl-context-switch.cpp diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 682cc7d7be535..06e0551ce1747 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11219,6 +11219,8 @@ def err_esimd_glob_cant_init : Error< "SYCL explicit SIMD does not permit private global variable to have an initializer">; def err_esimd_global_in_sycl_context : Error< "ESIMD globals cannot be used in a SYCL context">; +def err_sycl_device_function_is_called_from_esimd : Error< + "SYCL device function cannot be called from ESIMD context">; def err_nullptr_t_type_in_sycl_kernel : Error<"%0 is an invalid kernel name, " "'std::nullptr_t' is declared in the 'std' namespace ">; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 44c99f7e7b4e0..e0d021a38ec64 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -241,6 +241,19 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, !checkAllowedSYCLInitializer(VD, /*CheckValueDependent =*/true)) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelConstStaticVariable; + } else if (isa(D)) { + // SYCL device function cannot be called from ESIMD context. However, + // there are some device function declarations that are shared between + // SYCL and ESIMD, e.g. spirv builtins. Those are reserved functions + // and we allow to call them from ESIMD context. + FunctionDecl *FDecl = cast(D); + const IdentifierInfo *Id = FDecl->getIdentifier(); + if ((getEmissionReason(FDecl) == Sema::DeviceDiagnosticReason::Sycl) && + Id && !Id->isReservedName(/*doubleUnderscoreOnly=*/true)) { + SYCLDiagIfDeviceCode( + *Locs.begin(), diag::err_sycl_device_function_is_called_from_esimd, + Sema::DeviceDiagnosticReason::Esimd); + } } } diff --git a/clang/test/CodeGenSYCL/esimd_metadata2.cpp b/clang/test/CodeGenSYCL/esimd_metadata2.cpp index 5c8ebb09f26b0..b41ae63bd0159 100644 --- a/clang/test/CodeGenSYCL/esimd_metadata2.cpp +++ b/clang/test/CodeGenSYCL/esimd_metadata2.cpp @@ -3,8 +3,8 @@ // This test checks that attribute !intel_reqd_sub_group_size !1 // is added for kernels with !sycl_explicit_simd -__attribute__((sycl_device)) void shared_func_decl(); -__attribute__((sycl_device)) void shared_func() { shared_func_decl(); } +void shared_func_decl(); +void shared_func() { shared_func_decl(); } __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func() { shared_func(); } diff --git a/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp b/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp new file mode 100644 index 0000000000000..e48aac050cf41 --- /dev/null +++ b/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp @@ -0,0 +1,27 @@ +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s + +// This test checks that SYCL device functions cannot be called from ESIMD context. + +__attribute__((sycl_device)) void sycl_func() {} +__attribute__((sycl_device)) void __reserved_func() {} + +// Immediate diagnostic +__attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func1() { + // expected-error@+1{{SYCL device function cannot be called from ESIMD context}} + sycl_func(); + // Reserved functions are allowed + __reserved_func(); +} + +// Deffered diagnostic +void foo() { + // expected-error@+1{{SYCL device function cannot be called from ESIMD context}} + sycl_func(); + // Reserved functions are allowed + __reserved_func(); +} + +__attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func2() { + // expected-note@+1{{called by}} + foo(); +} diff --git a/sycl/include/CL/sycl/INTEL/esimd.hpp b/sycl/include/CL/sycl/INTEL/esimd.hpp index 1e79f875df4d7..582c2f2dfbca3 100644 --- a/sycl/include/CL/sycl/INTEL/esimd.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd.hpp @@ -10,13 +10,6 @@ #pragma once -/// \defgroup sycl_esimd DPC++ Explicit SIMD API - -#include -#include -#include -#include - #ifdef __SYCL_DEVICE_ONLY__ #define SYCL_ESIMD_KERNEL __attribute__((sycl_explicit_simd)) #define SYCL_ESIMD_FUNCTION __attribute__((sycl_explicit_simd)) @@ -24,3 +17,10 @@ #define SYCL_ESIMD_KERNEL #define SYCL_ESIMD_FUNCTION #endif + +/// \defgroup sycl_esimd DPC++ Explicit SIMD API + +#include +#include +#include +#include diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp index dba3ebf5140b6..a4f072c317af9 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp @@ -639,7 +639,7 @@ inline ESIMD_NODEBUG void esimd_sbarrier(EsimdSbarrierType flag) { /// @{ /// Declare per-work-group slm size. -SYCL_EXTERNAL void slm_init(uint32_t size); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void slm_init(uint32_t size); /// SLM gather. /// diff --git a/sycl/test/esimd/global_var.cpp b/sycl/test/esimd/global_var.cpp index 1503159c4b09d..93ddc8fa53541 100644 --- a/sycl/test/esimd/global_var.cpp +++ b/sycl/test/esimd/global_var.cpp @@ -28,8 +28,6 @@ SYCL_EXTERNAL void init_vc_sycl(int x) { vc = x; } -SYCL_EXTERNAL void foo() {} - void kernel_call() { queue q; q.submit([&](cl::sycl::handler &cgh) { @@ -38,7 +36,6 @@ void kernel_call() { // ESIMD kernel is allowed to use ESIMD // global vc = 0; - foo(); func_that_uses_esimd_glob(); }); }); From 16a1d387cd24a20f28c6a4a6f57fd2866f755097 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Wed, 7 Apr 2021 23:58:17 -0700 Subject: [PATCH 2/7] Fixed failing test --- clang/test/CodeGenSYCL/esimd_metadata2.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/clang/test/CodeGenSYCL/esimd_metadata2.cpp b/clang/test/CodeGenSYCL/esimd_metadata2.cpp index b41ae63bd0159..102ec1323773a 100644 --- a/clang/test/CodeGenSYCL/esimd_metadata2.cpp +++ b/clang/test/CodeGenSYCL/esimd_metadata2.cpp @@ -3,8 +3,7 @@ // This test checks that attribute !intel_reqd_sub_group_size !1 // is added for kernels with !sycl_explicit_simd -void shared_func_decl(); -void shared_func() { shared_func_decl(); } +void shared_func() { } __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func() { shared_func(); } @@ -12,7 +11,6 @@ __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func // CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}esimd_funcv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} { // 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: declare spir_func void @{{.*}}shared_func_declv() #{{[0-9]+}} class ESIMDFunctor { public: From 0253b801dc6f9c8edf2e587f07cbaa5627f774c5 Mon Sep 17 00:00:00 2001 From: DenisBakhvalov <61807338+DenisBakhvalov@users.noreply.github.com> Date: Thu, 8 Apr 2021 13:51:12 -0700 Subject: [PATCH 3/7] Apply suggestions from code review Co-authored-by: Aaron Ballman --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/lib/Sema/SemaExpr.cpp | 3 +-- clang/test/SemaSYCL/esimd-sycl-context-switch.cpp | 2 +- 3 files changed, 3 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 06e0551ce1747..07b301c165191 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11220,7 +11220,7 @@ def err_esimd_glob_cant_init : Error< def err_esimd_global_in_sycl_context : Error< "ESIMD globals cannot be used in a SYCL context">; def err_sycl_device_function_is_called_from_esimd : Error< - "SYCL device function cannot be called from ESIMD context">; + "SYCL device function cannot be called from an ESIMD context">; def err_nullptr_t_type_in_sycl_kernel : Error<"%0 is an invalid kernel name, " "'std::nullptr_t' is declared in the 'std' namespace ">; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index e0d021a38ec64..2feeef1827d50 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -241,12 +241,11 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, !checkAllowedSYCLInitializer(VD, /*CheckValueDependent =*/true)) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelConstStaticVariable; - } else if (isa(D)) { + } else if (auto *FDecl = dyn_cast(D)) { // SYCL device function cannot be called from ESIMD context. However, // there are some device function declarations that are shared between // SYCL and ESIMD, e.g. spirv builtins. Those are reserved functions // and we allow to call them from ESIMD context. - FunctionDecl *FDecl = cast(D); const IdentifierInfo *Id = FDecl->getIdentifier(); if ((getEmissionReason(FDecl) == Sema::DeviceDiagnosticReason::Sycl) && Id && !Id->isReservedName(/*doubleUnderscoreOnly=*/true)) { diff --git a/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp b/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp index e48aac050cf41..4ebaf81eb6bdf 100644 --- a/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp +++ b/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp @@ -13,7 +13,7 @@ __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func __reserved_func(); } -// Deffered diagnostic +// Deferred diagnostic void foo() { // expected-error@+1{{SYCL device function cannot be called from ESIMD context}} sycl_func(); From ba430cee055806c17610b7ec67daa4a7db5966e7 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Thu, 8 Apr 2021 16:10:27 -0700 Subject: [PATCH 4/7] Limit the set of allowed SYCL_EXTERNAL functions --- clang/lib/Sema/SemaExpr.cpp | 3 +- .../SemaSYCL/esimd-sycl-context-switch.cpp | 17 ++- .../sycl/INTEL/esimd/detail/esimd_intrin.hpp | 29 ++-- .../INTEL/esimd/detail/esimd_math_intrin.hpp | 140 +++++++++--------- .../esimd/detail/esimd_memory_intrin.hpp | 107 +++++++------ sycl/test/esimd/esimd-util-compiler-eval.cpp | 2 +- 6 files changed, 158 insertions(+), 140 deletions(-) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 2feeef1827d50..8270a092be021 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -248,7 +248,8 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, // and we allow to call them from ESIMD context. const IdentifierInfo *Id = FDecl->getIdentifier(); if ((getEmissionReason(FDecl) == Sema::DeviceDiagnosticReason::Sycl) && - Id && !Id->isReservedName(/*doubleUnderscoreOnly=*/true)) { + Id && !Id->getName().startswith("__spirv_") && + !Id->getName().startswith("__sycl_")) { SYCLDiagIfDeviceCode( *Locs.begin(), diag::err_sycl_device_function_is_called_from_esimd, Sema::DeviceDiagnosticReason::Esimd); diff --git a/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp b/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp index 4ebaf81eb6bdf..f67d5990c8216 100644 --- a/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp +++ b/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp @@ -3,22 +3,25 @@ // This test checks that SYCL device functions cannot be called from ESIMD context. __attribute__((sycl_device)) void sycl_func() {} -__attribute__((sycl_device)) void __reserved_func() {} +__attribute__((sycl_device)) void __spirv_reserved_func() {} +__attribute__((sycl_device)) void __sycl_reserved_func() {} +__attribute__((sycl_device)) void __other_reserved_func() {} // Immediate diagnostic __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func1() { - // expected-error@+1{{SYCL device function cannot be called from ESIMD context}} + // expected-error@+1{{SYCL device function cannot be called from an ESIMD context}} sycl_func(); - // Reserved functions are allowed - __reserved_func(); + // Reserved SPIRV and SYCL functions are allowed + __spirv_reserved_func(); + __sycl_reserved_func(); + // expected-error@+1{{SYCL device function cannot be called from an ESIMD context}} + __other_reserved_func(); } // Deferred diagnostic void foo() { - // expected-error@+1{{SYCL device function cannot be called from ESIMD context}} + // expected-error@+1{{SYCL device function cannot be called from an ESIMD context}} sycl_func(); - // Reserved functions are allowed - __reserved_func(); } __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func2() { diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp index 330eefd047605..b871589fbed95 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp @@ -64,11 +64,11 @@ // template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_rdregion(__SIGD::vector_type_t Input, uint16_t Offset); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_rdindirect(__SIGD::vector_type_t Input, __SIGD::vector_type_t Offset); @@ -121,13 +121,13 @@ __esimd_rdindirect(__SIGD::vector_type_t Input, // template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_wrregion(__SIGD::vector_type_t OldVal, __SIGD::vector_type_t NewVal, uint16_t Offset, sycl::INTEL::gpu::mask_type_t Mask = 1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_wrindirect(__SIGD::vector_type_t OldVal, __SIGD::vector_type_t NewVal, __SIGD::vector_type_t Offset, @@ -217,7 +217,7 @@ readRegion(const __SIGD::vector_type_t &Base, std::pair Region) { // optimization on simd object // template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_vload(const __SIGD::vector_type_t *ptr); // vstore @@ -225,21 +225,24 @@ __esimd_vload(const __SIGD::vector_type_t *ptr); // map to the backend vstore intrinsic, used by compiler to control // optimization on simd object template -SYCL_EXTERNAL void __esimd_vstore(__SIGD::vector_type_t *ptr, - __SIGD::vector_type_t vals); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void +__esimd_vstore(__SIGD::vector_type_t *ptr, + __SIGD::vector_type_t vals); template -SYCL_EXTERNAL uint16_t __esimd_any(__SIGD::vector_type_t src); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION uint16_t +__esimd_any(__SIGD::vector_type_t src); template -SYCL_EXTERNAL uint16_t __esimd_all(__SIGD::vector_type_t src); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION uint16_t +__esimd_all(__SIGD::vector_type_t src); #ifndef __SYCL_DEVICE_ONLY__ // Implementations of ESIMD intrinsics for the SYCL host device template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_rdregion(__SIGD::vector_type_t Input, uint16_t Offset) { uint16_t EltOffset = Offset / sizeof(T); assert(Offset % sizeof(T) == 0); @@ -258,7 +261,7 @@ __esimd_rdregion(__SIGD::vector_type_t Input, uint16_t Offset) { } template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_rdindirect(__SIGD::vector_type_t Input, __SIGD::vector_type_t Offset) { __SIGD::vector_type_t Result; @@ -273,7 +276,7 @@ __esimd_rdindirect(__SIGD::vector_type_t Input, template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_wrregion(__SIGD::vector_type_t OldVal, __SIGD::vector_type_t NewVal, uint16_t Offset, sycl::INTEL::gpu::mask_type_t Mask) { @@ -296,7 +299,7 @@ __esimd_wrregion(__SIGD::vector_type_t OldVal, } template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_wrindirect(__SIGD::vector_type_t OldVal, __SIGD::vector_type_t NewVal, __SIGD::vector_type_t Offset, diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_math_intrin.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_math_intrin.hpp index 5d2fb0f856fa5..a91d5fd863254 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_math_intrin.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_math_intrin.hpp @@ -20,298 +20,298 @@ // saturation intrinsics template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_satf(__SIGD::vector_type_t src); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_fptoui_sat(__SIGD::vector_type_t src); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_fptosi_sat(__SIGD::vector_type_t src); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_uutrunc_sat(__SIGD::vector_type_t src); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_ustrunc_sat(__SIGD::vector_type_t src); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sutrunc_sat(__SIGD::vector_type_t src); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sstrunc_sat(__SIGD::vector_type_t src); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_abs(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_ssshl(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sushl(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_usshl(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_uushl(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_ssshl_sat(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sushl_sat(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_usshl_sat(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_uushl_sat(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_rol(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_ror(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_umulh(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_smulh(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_frc(__SIGD::vector_type_t src0); /// 3 kinds of max template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_fmax(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_umax(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_smax(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_lzd(__SIGD::vector_type_t src0); /// 3 kinds of min template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_fmin(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_umin(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_smin(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_bfrev(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_cbit(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t __esimd_bfins( +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_bfins( __SIGD::vector_type_t src0, __SIGD::vector_type_t src1, __SIGD::vector_type_t src2, __SIGD::vector_type_t src3); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_bfext(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_fbl(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sfbh(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_ufbh(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_inv(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_log(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_exp(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sqrt(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sqrt_ieee(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_rsqrt(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sin(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_cos(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_pow(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_div_ieee(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_rndd(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_rndu(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_rnde(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_rndz(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sqrt_ieee(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_div_ieee(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1); template -SYCL_EXTERNAL uint32_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION uint32_t __esimd_pack_mask(__SIGD::vector_type_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_unpack_mask(uint32_t src0); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_uudp4a(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_usdp4a(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sudp4a(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_ssdp4a(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_uudp4a_sat(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_usdp4a_sat(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_sudp4a_sat(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_ssdp4a_sat(__SIGD::vector_type_t src0, __SIGD::vector_type_t src1, __SIGD::vector_type_t src2); // Reduction functions template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_reduced_fmax(__SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_reduced_umax(__SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_reduced_smax(__SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_reduced_fmin(__SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_reduced_umin(__SIGD::vector_type_t src1, __SIGD::vector_type_t src2); template -__SIGD::vector_type_t - SYCL_EXTERNAL __esimd_reduced_smin(__SIGD::vector_type_t src1, - __SIGD::vector_type_t src2); +__SIGD::vector_type_t SYCL_EXTERNAL SYCL_ESIMD_FUNCTION +__esimd_reduced_smin(__SIGD::vector_type_t src1, + __SIGD::vector_type_t src2); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_dp4(__SIGD::vector_type_t v1, __SIGD::vector_type_t v2); #ifndef __SYCL_DEVICE_ONLY__ diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_memory_intrin.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_memory_intrin.hpp index 60ed5eaefb687..1a535a2cfacd2 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_memory_intrin.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_memory_intrin.hpp @@ -74,7 +74,7 @@ constexpr unsigned int ElemsPerAddrDecoding(unsigned int ElemsPerAddrEncoded) { template -SYCL_EXTERNAL +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_flat_read(__SIGD::vector_type_t addrs, int ElemsPerAddr = NumBlk, @@ -84,7 +84,7 @@ SYCL_EXTERNAL template -SYCL_EXTERNAL void __esimd_flat_write( +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_flat_write( __SIGD::vector_type_t addrs, __SIGD::vector_type_t vals, int ElemsPerAddr = NumBlk, __SIGD::vector_type_t pred = 1); @@ -93,39 +93,41 @@ SYCL_EXTERNAL void __esimd_flat_write( template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_flat_block_read_unaligned(uint64_t addr); // flat_block_write writes a block of data using one flat address template -SYCL_EXTERNAL void __esimd_flat_block_write(uint64_t addr, - __SIGD::vector_type_t vals); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void +__esimd_flat_block_write(uint64_t addr, __SIGD::vector_type_t vals); // Reads a block of data from given surface at given offset. template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_block_read(SurfIndAliasTy surf_ind, uint32_t offset); // Writes given block of data to a surface with given index at given offset. template -SYCL_EXTERNAL void __esimd_block_write(SurfIndAliasTy surf_ind, uint32_t offset, - __SIGD::vector_type_t vals); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void +__esimd_block_write(SurfIndAliasTy surf_ind, uint32_t offset, + __SIGD::vector_type_t vals); // flat_read4 does flat-address gather4 template -__SIGD::vector_type_t SYCL_EXTERNAL -__esimd_flat_read4(__SIGD::vector_type_t addrs, - __SIGD::vector_type_t pred = 1); +__SIGD::vector_type_t + SYCL_EXTERNAL SYCL_ESIMD_FUNCTION + __esimd_flat_read4(__SIGD::vector_type_t addrs, + __SIGD::vector_type_t pred = 1); // flat_write does flat-address scatter template -SYCL_EXTERNAL void +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_flat_write4(__SIGD::vector_type_t addrs, __SIGD::vector_type_t vals, __SIGD::vector_type_t pred = 1); @@ -154,7 +156,7 @@ __esimd_flat_write4(__SIGD::vector_type_t addrs, template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_surf_read(int16_t scale, SurfIndAliasTy surf_ind, uint32_t global_offset, __SIGD::vector_type_t elem_offsets) @@ -196,7 +198,7 @@ __esimd_surf_read(int16_t scale, SurfIndAliasTy surf_ind, template -SYCL_EXTERNAL void +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_surf_write(__SIGD::vector_type_t pred, int16_t scale, SurfIndAliasTy surf_ind, uint32_t global_offset, __SIGD::vector_type_t elem_offsets, @@ -219,14 +221,14 @@ __esimd_surf_write(__SIGD::vector_type_t pred, int16_t scale, template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_flat_atomic0(__SIGD::vector_type_t addrs, __SIGD::vector_type_t pred); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_flat_atomic1(__SIGD::vector_type_t addrs, __SIGD::vector_type_t src0, __SIGD::vector_type_t pred); @@ -234,71 +236,77 @@ __esimd_flat_atomic1(__SIGD::vector_type_t addrs, template -SYCL_EXTERNAL __SIGD::vector_type_t __esimd_flat_atomic2( - __SIGD::vector_type_t addrs, __SIGD::vector_type_t src0, - __SIGD::vector_type_t src1, __SIGD::vector_type_t pred); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t +__esimd_flat_atomic2(__SIGD::vector_type_t addrs, + __SIGD::vector_type_t src0, + __SIGD::vector_type_t src1, + __SIGD::vector_type_t pred); // esimd_barrier, generic group barrier -SYCL_EXTERNAL void __esimd_barrier(); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_barrier(); // generic work-group split barrier -SYCL_EXTERNAL void __esimd_sbarrier(sycl::INTEL::gpu::EsimdSbarrierType flag); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void +__esimd_sbarrier(sycl::INTEL::gpu::EsimdSbarrierType flag); // slm_fence sets the SLM read/write order -SYCL_EXTERNAL void __esimd_slm_fence(uint8_t cntl); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_slm_fence(uint8_t cntl); // slm_read does SLM gather template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_slm_read(__SIGD::vector_type_t addrs, __SIGD::vector_type_t pred = 1); // slm_write does SLM scatter template -SYCL_EXTERNAL void +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_slm_write(__SIGD::vector_type_t addrs, __SIGD::vector_type_t vals, __SIGD::vector_type_t pred = 1); // slm_block_read reads a block of data from SLM template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_slm_block_read(uint32_t addr); // slm_block_write writes a block of data to SLM template -SYCL_EXTERNAL void __esimd_slm_block_write(uint32_t addr, - __SIGD::vector_type_t vals); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void +__esimd_slm_block_write(uint32_t addr, __SIGD::vector_type_t vals); // slm_read4 does SLM gather4 template -SYCL_EXTERNAL __SIGD::vector_type_t -__esimd_slm_read4(__SIGD::vector_type_t addrs, - __SIGD::vector_type_t pred = 1); +SYCL_EXTERNAL + SYCL_ESIMD_FUNCTION __SIGD::vector_type_t + __esimd_slm_read4(__SIGD::vector_type_t addrs, + __SIGD::vector_type_t pred = 1); // slm_write4 does SLM scatter4 template -SYCL_EXTERNAL void +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_slm_write4(__SIGD::vector_type_t addrs, __SIGD::vector_type_t vals, __SIGD::vector_type_t pred = 1); // slm_atomic: SLM atomic template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_slm_atomic0(__SIGD::vector_type_t addrs, __SIGD::vector_type_t pred); template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_slm_atomic1(__SIGD::vector_type_t addrs, __SIGD::vector_type_t src0, __SIGD::vector_type_t pred); template -SYCL_EXTERNAL __SIGD::vector_type_t __esimd_slm_atomic2( - __SIGD::vector_type_t addrs, __SIGD::vector_type_t src0, - __SIGD::vector_type_t src1, __SIGD::vector_type_t pred); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t +__esimd_slm_atomic2(__SIGD::vector_type_t addrs, + __SIGD::vector_type_t src0, + __SIGD::vector_type_t src1, + __SIGD::vector_type_t pred); // Media block load // @@ -325,7 +333,7 @@ SYCL_EXTERNAL __SIGD::vector_type_t __esimd_slm_atomic2( // @return the linearized 2D block data read from surface. // template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_media_block_load(unsigned modififer, TACC handle, unsigned plane, unsigned width, unsigned x, unsigned y); @@ -354,7 +362,7 @@ __esimd_media_block_load(unsigned modififer, TACC handle, unsigned plane, // @param vals the linearized 2D block data to be written to surface. // template -SYCL_EXTERNAL void +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_media_block_store(unsigned modififer, TACC handle, unsigned plane, unsigned width, unsigned x, unsigned y, __SIGD::vector_type_t vals); @@ -366,7 +374,8 @@ __esimd_media_block_store(unsigned modififer, TACC handle, unsigned plane, /// Returns the binding table index value. /// template -SYCL_EXTERNAL uint32_t __esimd_get_value(SurfIndAliasTy sid); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION uint32_t +__esimd_get_value(SurfIndAliasTy sid); /// \brief Raw sends load. /// @@ -401,12 +410,14 @@ SYCL_EXTERNAL uint32_t __esimd_get_value(SurfIndAliasTy sid); /// template -SYCL_EXTERNAL __SIGD::vector_type_t __esimd_raw_sends_load( - uint8_t modifier, uint8_t execSize, __SIGD::vector_type_t pred, - uint8_t numSrc0, uint8_t numSrc1, uint8_t numDst, uint8_t sfid, - uint32_t exDesc, uint32_t msgDesc, __SIGD::vector_type_t msgSrc0, - __SIGD::vector_type_t msgSrc1, - __SIGD::vector_type_t msgDst); +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t +__esimd_raw_sends_load(uint8_t modifier, uint8_t execSize, + __SIGD::vector_type_t pred, uint8_t numSrc0, + uint8_t numSrc1, uint8_t numDst, uint8_t sfid, + uint32_t exDesc, uint32_t msgDesc, + __SIGD::vector_type_t msgSrc0, + __SIGD::vector_type_t msgSrc1, + __SIGD::vector_type_t msgDst); /// \brief Raw send load. /// @@ -435,7 +446,7 @@ SYCL_EXTERNAL __SIGD::vector_type_t __esimd_raw_sends_load( /// Returns a simd vector of type Ty1 and size N1. /// template -SYCL_EXTERNAL __SIGD::vector_type_t +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t __esimd_raw_send_load(uint8_t modifier, uint8_t execSize, __SIGD::vector_type_t pred, uint8_t numSrc0, uint8_t numDst, uint8_t sfid, uint32_t exDesc, @@ -467,7 +478,7 @@ __esimd_raw_send_load(uint8_t modifier, uint8_t execSize, /// @param msgSrc1 the second source operand of send message. /// template -SYCL_EXTERNAL void __esimd_raw_sends_store( +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_raw_sends_store( uint8_t modifier, uint8_t execSize, __SIGD::vector_type_t pred, uint8_t numSrc0, uint8_t numSrc1, uint8_t sfid, uint32_t exDesc, uint32_t msgDesc, __SIGD::vector_type_t msgSrc0, @@ -493,7 +504,7 @@ SYCL_EXTERNAL void __esimd_raw_sends_store( /// @param msgSrc0 the first source operand of send message. /// template -SYCL_EXTERNAL void +SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void __esimd_raw_send_store(uint8_t modifier, uint8_t execSize, __SIGD::vector_type_t pred, uint8_t numSrc0, uint8_t sfid, uint32_t exDesc, uint32_t msgDesc, diff --git a/sycl/test/esimd/esimd-util-compiler-eval.cpp b/sycl/test/esimd/esimd-util-compiler-eval.cpp index d4652bc8d504f..ae769acd5caf5 100644 --- a/sycl/test/esimd/esimd-util-compiler-eval.cpp +++ b/sycl/test/esimd/esimd-util-compiler-eval.cpp @@ -2,7 +2,7 @@ // This test checks compile-time evaluation of functions from esimd_util.hpp #include "CL/sycl.hpp" -#include "CL/sycl/INTEL/esimd/esimd.hpp" +#include "CL/sycl/INTEL/esimd.hpp" static_assert(sycl::INTEL::gpu::detail::getNextPowerOf2<0>() == 0, ""); static_assert(sycl::INTEL::gpu::detail::getNextPowerOf2<1>() == 1, ""); From 47e5030346bee7cdb1cba298f5feacbc385012a8 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Thu, 8 Apr 2021 16:18:20 -0700 Subject: [PATCH 5/7] Fixed a comment --- clang/lib/Sema/SemaExpr.cpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 8270a092be021..6df81279e485d 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -242,10 +242,9 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelConstStaticVariable; } else if (auto *FDecl = dyn_cast(D)) { - // SYCL device function cannot be called from ESIMD context. However, - // there are some device function declarations that are shared between - // SYCL and ESIMD, e.g. spirv builtins. Those are reserved functions - // and we allow to call them from ESIMD context. + // SYCL device function cannot be called from an ESIMD context. However, + // funcitons that start with '__spirv_' or '__sycl_' are exceptions to + // this rule. const IdentifierInfo *Id = FDecl->getIdentifier(); if ((getEmissionReason(FDecl) == Sema::DeviceDiagnosticReason::Sycl) && Id && !Id->getName().startswith("__spirv_") && From 9372dc548cd20aceb74e0ab9f8c624548a84d391 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Fri, 9 Apr 2021 09:54:49 -0700 Subject: [PATCH 6/7] Added more tests --- .../SemaSYCL/esimd-sycl-context-switch.cpp | 61 ++++++++++++++++++- 1 file changed, 59 insertions(+), 2 deletions(-) diff --git a/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp b/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp index f67d5990c8216..d3851f34c4381 100644 --- a/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp +++ b/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp @@ -7,7 +7,7 @@ __attribute__((sycl_device)) void __spirv_reserved_func() {} __attribute__((sycl_device)) void __sycl_reserved_func() {} __attribute__((sycl_device)) void __other_reserved_func() {} -// Immediate diagnostic +// -- Immediate diagnostic __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func1() { // expected-error@+1{{SYCL device function cannot be called from an ESIMD context}} sycl_func(); @@ -18,7 +18,7 @@ __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func __other_reserved_func(); } -// Deferred diagnostic +// -- Deferred diagnostic void foo() { // expected-error@+1{{SYCL device function cannot be called from an ESIMD context}} sycl_func(); @@ -28,3 +28,60 @@ __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func // expected-note@+1{{called by}} foo(); } + +// -- Class method +struct S { + __attribute__((sycl_device)) void sycl_func() {} +}; + +__attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func3() { + S s; + // expected-error@+1{{SYCL device function cannot be called from an ESIMD context}} + s.sycl_func(); +} + +// -- Template function +template +__attribute__((sycl_device)) void sycl_func() {} + +__attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func4() { + // expected-error@+1{{SYCL device function cannot be called from an ESIMD context}} + sycl_func(); +} + +// -- std::function +namespace std { +template +_Tp declval(); + +template +struct __res { + template + static decltype(declval<_Functor>()(_Args()...)) _S_test(int); + + template + static void _S_test(...); + + typedef decltype(_S_test<_ArgTypes...>(0)) type; +}; + +template +struct function; + +template +struct function<_R(_ArgTypes...)> { + template ::type> + __attribute__((sycl_device, sycl_explicit_simd)) function(_Functor) {} + __attribute__((sycl_device, sycl_explicit_simd)) _R operator()(_ArgTypes...) const; +}; +} // namespace std + +__attribute__((sycl_device)) void sycl_func1() {} + +__attribute__((sycl_device, sycl_explicit_simd)) void passthrough(std::function &&C) { C(); } + +__attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func5() { + // expected-error@+1{{SYCL device function cannot be called from an ESIMD context}} + passthrough(sycl_func1); +} \ No newline at end of file From 29c367b121f91b0fe596a343b9ba6f85fcba50a3 Mon Sep 17 00:00:00 2001 From: DenisBakhvalov <61807338+DenisBakhvalov@users.noreply.github.com> Date: Mon, 12 Apr 2021 10:04:22 -0700 Subject: [PATCH 7/7] Update esimd-sycl-context-switch.cpp Added a new line --- clang/test/SemaSYCL/esimd-sycl-context-switch.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp b/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp index d3851f34c4381..4b04ef44c040f 100644 --- a/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp +++ b/clang/test/SemaSYCL/esimd-sycl-context-switch.cpp @@ -84,4 +84,4 @@ __attribute__((sycl_device, sycl_explicit_simd)) void passthrough(std::function< __attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func5() { // expected-error@+1{{SYCL device function cannot be called from an ESIMD context}} passthrough(sycl_func1); -} \ No newline at end of file +}