Skip to content

[SYCL] Fix free function queries for host device #4365

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
merged 2 commits into from
Aug 19, 2021
Merged

[SYCL] Fix free function queries for host device #4365

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
07b75a9
[SYCL] Fix free function queries for host device
Aug 18, 2021
576ea2d
fix test
Aug 18, 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
11 changes: 9 additions & 2 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5302,10 +5302,17 @@ bool Util::isSyclFunction(const FunctionDecl *FD, StringRef Name) {
if (DC->isTranslationUnit())
return false;

std::array<DeclContextDesc, 2> Scopes = {
std::array<DeclContextDesc, 2> ScopesSycl = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl")};
return matchContext(DC, Scopes);
std::array<DeclContextDesc, 5> ScopesOneapiExp = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "experimental")};

return matchContext(DC, ScopesSycl) || matchContext(DC, ScopesOneapiExp);
}

bool Util::isAccessorPropertyListType(QualType Ty) {
Expand Down
21 changes: 14 additions & 7 deletions clang/test/CodeGenSYCL/Inputs/sycl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -120,13 +120,6 @@ struct no_alias {
} // namespace oneapi
} // namespace ext

namespace ext {
namespace oneapi {
template <typename... properties>
class accessor_property_list {};
} // namespace oneapi
} // namespace ext

template <int dim>
struct id {
template <typename... T>
Expand All @@ -146,6 +139,20 @@ template <int dim> struct item {
int Data;
};

namespace ext {
namespace oneapi {
template <typename... properties>
class accessor_property_list {};
namespace experimental {
template <int Dims> item<Dims>
this_item() { return item<Dims>{}; }

template <int Dims> id<Dims>
this_id() { return id<Dims>{}; }
} // namespace experimental
} // namespace oneapi
} // namespace ext

template <int Dims> item<Dims>
this_item() { return item<Dims>{}; }

Expand Down
90 changes: 89 additions & 1 deletion clang/test/CodeGenSYCL/parallel_for_this_item.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,10 +12,14 @@
// CHECK-NEXT: const char* const kernel_names[] = {
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3GNU",
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3EMU",
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3COW",
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3OWL",
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3RAT",
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3CAT",
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3FOX",
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3BEE"
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3PIG",
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3BEE",
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3DOG"
// CHECK-NEXT: };

// CHECK:template <> struct KernelInfo<GNU> {
Expand Down Expand Up @@ -50,6 +54,22 @@
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
// CHECK-NEXT:};
// CHECK-NEXT:template <> struct KernelInfo<COW> {
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3COW"; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr unsigned getNumParams() { return 0; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
// CHECK-NEXT: return kernel_signatures[i+0];
// CHECK-NEXT: }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool isESIMD() { return 0; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsThisItem() { return 1; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
// CHECK-NEXT:};
// CHECK-NEXT:template <> struct KernelInfo<OWL> {
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3OWL"; }
Expand Down Expand Up @@ -82,6 +102,22 @@
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
// CHECK-NEXT:};
// CHECK-NEXT:template <> struct KernelInfo<CAT> {
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3CAT"; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr unsigned getNumParams() { return 0; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
// CHECK-NEXT: return kernel_signatures[i+0];
// CHECK-NEXT: }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool isESIMD() { return 0; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsThisItem() { return 1; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
// CHECK-NEXT:};
// CHECK-NEXT:template <> struct KernelInfo<FOX> {
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3FOX"; }
Expand All @@ -98,6 +134,22 @@
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
// CHECK-NEXT:};
// CHECK-NEXT:template <> struct KernelInfo<PIG> {
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3PIG"; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr unsigned getNumParams() { return 0; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
// CHECK-NEXT: return kernel_signatures[i+0];
// CHECK-NEXT: }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool isESIMD() { return 0; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsThisItem() { return 0; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
// CHECK-NEXT:};
// CHECK-NEXT:template <> struct KernelInfo<BEE> {
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3BEE"; }
Expand All @@ -114,13 +166,31 @@
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
// CHECK-NEXT:};
// CHECK-NEXT:template <> struct KernelInfo<DOG> {
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3DOG"; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr unsigned getNumParams() { return 0; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
// CHECK-NEXT: return kernel_signatures[i+0];
// CHECK-NEXT: }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool isESIMD() { return 0; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsThisItem() { return 1; }
// CHECK-NEXT: __SYCL_DLL_LOCAL
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
// CHECK-NEXT:};

#include "sycl.hpp"

using namespace cl::sycl;

SYCL_EXTERNAL item<1> g() { return this_item<1>(); }
SYCL_EXTERNAL item<1> f() { return g(); }
SYCL_EXTERNAL item<1> s() { return ext::oneapi::experimental::this_item<1>(); }
SYCL_EXTERNAL item<1> h() { return s(); }

// This is a similar-looking this_item function but not the real one.
template <int Dims> item<Dims> this_item(int i) { return item<1>{i}; }
Expand All @@ -142,6 +212,11 @@ int main() {
cgh.parallel_for<class EMU>(range<1>(1),
[=](::item<1> I) { this_item<1>(); });

// This kernel calls sycl::ext::oneapi::experimental::this_item
cgh.parallel_for<class COW>(range<1>(1), [=](::item<1> I) {
ext::oneapi::experimental::this_item<1>();
});

// This kernel does not call sycl::this_item
cgh.parallel_for<class OWL>(range<1>(1), [=](id<1> I) {
class C c;
Expand All @@ -151,11 +226,24 @@ int main() {
// This kernel calls sycl::this_item
cgh.parallel_for<class RAT>(range<1>(1), [=](id<1> I) { f(); });

// This kernel calls sycl::ext::oneapi::experimental::this_item
cgh.parallel_for<class CAT>(range<1>(1), [=](id<1> I) { h(); });

// This kernel does not call sycl::this_item, but does call this_id
cgh.parallel_for<class FOX>(range<1>(1), [=](id<1> I) { this_id<1>(); });

// This kernel calls sycl::ext::oneapi::experimental::this_id
cgh.parallel_for<class PIG>(range<1>(1), [=](id<1> I) {
ext::oneapi::experimental::this_id<1>();
});

// This kernel calls sycl::this_item
cgh.parallel_for<class BEE>(range<1>(1), [=](auto I) { this_item<1>(); });

// This kernel calls sycl::ext::oneapi::experimental::this_item
cgh.parallel_for<class DOG>(range<1>(1), [=](auto I) {
ext::oneapi::experimental::this_item<1>();
});
});

return 0;
Expand Down