Skip to content

[ESIMD] SYCL device function cannot be called from ESIMD context #3511

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
kbobrovs merged 7 commits into from
Apr 21, 2021
Merged

[ESIMD] SYCL device function cannot be called from ESIMD context #3511

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
904b81a
[ESIMD] SYCL device function cannot be called from ESIMD context
DenisBakhvalov Apr 8, 2021
16a1d38
Fixed failing test
DenisBakhvalov Apr 8, 2021
0253b80
Apply suggestions from code review
DenisBakhvalov Apr 8, 2021
ba430ce
Limit the set of allowed SYCL_EXTERNAL functions
DenisBakhvalov Apr 8, 2021
47e5030
Fixed a comment
DenisBakhvalov Apr 8, 2021
9372dc5
Added more tests
DenisBakhvalov Apr 9, 2021
29c367b
Update esimd-sycl-context-switch.cpp
DenisBakhvalov Apr 12, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 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 ">;
Expand Down
12 changes: 12 additions & 0 deletions clang/lib/Sema/SemaExpr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -241,6 +241,18 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
!checkAllowedSYCLInitializer(VD, /*CheckValueDependent =*/true))
SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
<< Sema::KernelConstStaticVariable;
} else if (auto *FDecl = dyn_cast<FunctionDecl>(D)) {
// 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_") &&
!Id->getName().startswith("__sycl_")) {
SYCLDiagIfDeviceCode(
*Locs.begin(), diag::err_sycl_device_function_is_called_from_esimd,
Sema::DeviceDiagnosticReason::Esimd);
}
}
}

Expand Down
4 changes: 1 addition & 3 deletions clang/test/CodeGenSYCL/esimd_metadata2.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,16 +3,14 @@
// 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() { }

__attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func() { shared_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]+}} {
// 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:
Expand Down
87 changes: 87 additions & 0 deletions clang/test/SemaSYCL/esimd-sycl-context-switch.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,87 @@
// 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 __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 an ESIMD context}}
sycl_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 an ESIMD context}}
sycl_func();
}

__attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func2() {
// 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 <typename Ty>
__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<int>();
}

// -- std::function
namespace std {
template <typename _Tp>
_Tp declval();

template <typename _Functor, typename... _ArgTypes>
struct __res {
template <typename... _Args>
static decltype(declval<_Functor>()(_Args()...)) _S_test(int);

template <typename...>
static void _S_test(...);

typedef decltype(_S_test<_ArgTypes...>(0)) type;
};

template <typename>
struct function;

template <typename _R, typename... _ArgTypes>
struct function<_R(_ArgTypes...)> {
template <typename _Functor,
typename = typename __res<_Functor, _ArgTypes...>::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<void(void)> &&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);
}
14 changes: 7 additions & 7 deletions sycl/include/CL/sycl/INTEL/esimd.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,17 +10,17 @@

#pragma once

/// \defgroup sycl_esimd DPC++ Explicit SIMD API

#include <CL/sycl/INTEL/esimd/esimd.hpp>
#include <CL/sycl/INTEL/esimd/esimd_math.hpp>
#include <CL/sycl/INTEL/esimd/esimd_memory.hpp>
#include <CL/sycl/INTEL/esimd/esimd_view.hpp>

#ifdef __SYCL_DEVICE_ONLY__
#define SYCL_ESIMD_KERNEL __attribute__((sycl_explicit_simd))
#define SYCL_ESIMD_FUNCTION __attribute__((sycl_explicit_simd))
#else
#define SYCL_ESIMD_KERNEL
#define SYCL_ESIMD_FUNCTION
#endif

Copy link
Contributor Author

@DenisBakhvalov DenisBakhvalov Apr 8, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This change is needed because SYCL_ESIMD_FUNCTION is used in the headers included below.

/// \defgroup sycl_esimd DPC++ Explicit SIMD API

#include <CL/sycl/INTEL/esimd/esimd.hpp>
#include <CL/sycl/INTEL/esimd/esimd_math.hpp>
#include <CL/sycl/INTEL/esimd/esimd_memory.hpp>
#include <CL/sycl/INTEL/esimd/esimd_view.hpp>
29 changes: 16 additions & 13 deletions sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -64,11 +64,11 @@
//
template <typename T, int N, int M, int VStride, int Width, int Stride,
int ParentWidth = 0>
SYCL_EXTERNAL __SIGD::vector_type_t<T, M>
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t<T, M>
Copy link
Contributor

@kbobrovs kbobrovs Apr 20, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think the suggested way to add function attributes is after the last ')' - @erichkeane, is it correct?

Copy link
Contributor

@AaronBallman AaronBallman Apr 20, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

No, that adds the attribute to the function type and not the function declaration.

Copy link
Contributor

@erichkeane erichkeane Apr 20, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What Aaron said! For an attribute to apply to the function declaration it needs to go on the left. See SYCL_EXTERNAL, which is a function declaration attribute.

Copy link
Contributor

@kbobrovs kbobrovs Apr 21, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see, thanks. We'll need to refactor examples and tests.
What would be the right place for the attribute to apply to a lambda?
single_task<class test>([](item<1> i) {}); ?

Copy link
Contributor

@AaronBallman AaronBallman Apr 21, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In a lambda, the only place you can currently put an attribute is in the type position (this was resolved for C++23 but SYCL can't require use of that yet, for obvious reasons). So:

single_task<class test>([](item<1> i) [[attr]] {});

is the only place for it to go on a lambda for the moment.

__esimd_rdregion(__SIGD::vector_type_t<T, N> Input, uint16_t Offset);

template <typename T, int N, int M, int ParentWidth = 0>
SYCL_EXTERNAL __SIGD::vector_type_t<T, M>
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t<T, M>
__esimd_rdindirect(__SIGD::vector_type_t<T, N> Input,
__SIGD::vector_type_t<uint16_t, M> Offset);

Expand Down Expand Up @@ -121,13 +121,13 @@ __esimd_rdindirect(__SIGD::vector_type_t<T, N> Input,
//
template <typename T, int N, int M, int VStride, int Width, int Stride,
int ParentWidth = 0>
SYCL_EXTERNAL __SIGD::vector_type_t<T, N>
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t<T, N>
__esimd_wrregion(__SIGD::vector_type_t<T, N> OldVal,
__SIGD::vector_type_t<T, M> NewVal, uint16_t Offset,
sycl::INTEL::gpu::mask_type_t<M> Mask = 1);

template <typename T, int N, int M, int ParentWidth = 0>
SYCL_EXTERNAL __SIGD::vector_type_t<T, N>
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t<T, N>
__esimd_wrindirect(__SIGD::vector_type_t<T, N> OldVal,
__SIGD::vector_type_t<T, M> NewVal,
__SIGD::vector_type_t<uint16_t, M> Offset,
Expand Down Expand Up @@ -217,29 +217,32 @@ readRegion(const __SIGD::vector_type_t<BT, BN> &Base, std::pair<T, U> Region) {
// optimization on simd object
//
template <typename T, int N>
SYCL_EXTERNAL __SIGD::vector_type_t<T, N>
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t<T, N>
__esimd_vload(const __SIGD::vector_type_t<T, N> *ptr);

// vstore
//
// map to the backend vstore intrinsic, used by compiler to control
// optimization on simd object
template <typename T, int N>
SYCL_EXTERNAL void __esimd_vstore(__SIGD::vector_type_t<T, N> *ptr,
__SIGD::vector_type_t<T, N> vals);
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION void
__esimd_vstore(__SIGD::vector_type_t<T, N> *ptr,
__SIGD::vector_type_t<T, N> vals);

template <typename T, int N>
SYCL_EXTERNAL uint16_t __esimd_any(__SIGD::vector_type_t<T, N> src);
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION uint16_t
__esimd_any(__SIGD::vector_type_t<T, N> src);

template <typename T, int N>
SYCL_EXTERNAL uint16_t __esimd_all(__SIGD::vector_type_t<T, N> src);
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION uint16_t
__esimd_all(__SIGD::vector_type_t<T, N> src);

#ifndef __SYCL_DEVICE_ONLY__

// Implementations of ESIMD intrinsics for the SYCL host device
template <typename T, int N, int M, int VStride, int Width, int Stride,
int ParentWidth>
SYCL_EXTERNAL __SIGD::vector_type_t<T, M>
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t<T, M>
__esimd_rdregion(__SIGD::vector_type_t<T, N> Input, uint16_t Offset) {
uint16_t EltOffset = Offset / sizeof(T);
assert(Offset % sizeof(T) == 0);
Expand All @@ -258,7 +261,7 @@ __esimd_rdregion(__SIGD::vector_type_t<T, N> Input, uint16_t Offset) {
}

template <typename T, int N, int M, int ParentWidth>
SYCL_EXTERNAL __SIGD::vector_type_t<T, M>
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t<T, M>
__esimd_rdindirect(__SIGD::vector_type_t<T, N> Input,
__SIGD::vector_type_t<uint16_t, M> Offset) {
__SIGD::vector_type_t<T, M> Result;
Expand All @@ -273,7 +276,7 @@ __esimd_rdindirect(__SIGD::vector_type_t<T, N> Input,

template <typename T, int N, int M, int VStride, int Width, int Stride,
int ParentWidth>
SYCL_EXTERNAL __SIGD::vector_type_t<T, N>
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t<T, N>
__esimd_wrregion(__SIGD::vector_type_t<T, N> OldVal,
__SIGD::vector_type_t<T, M> NewVal, uint16_t Offset,
sycl::INTEL::gpu::mask_type_t<M> Mask) {
Expand All @@ -296,7 +299,7 @@ __esimd_wrregion(__SIGD::vector_type_t<T, N> OldVal,
}

template <typename T, int N, int M, int ParentWidth>
SYCL_EXTERNAL __SIGD::vector_type_t<T, N>
SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SIGD::vector_type_t<T, N>
__esimd_wrindirect(__SIGD::vector_type_t<T, N> OldVal,
__SIGD::vector_type_t<T, M> NewVal,
__SIGD::vector_type_t<uint16_t, M> Offset,
Expand Down
Loading