Skip to content

Commit 65436fb

Browse files
author
Alexander Batashev
authored
[SYCL] Fix free function queries for host device (#4365)
#4090 moved free function queries to a different namespace. Host device implementation actually relies on integration header to extract information about free function queries and set the appropriate values.
1 parent 3d96e1d commit 65436fb

File tree

3 files changed

+112
-10
lines changed

3 files changed

+112
-10
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 9 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -5302,10 +5302,17 @@ bool Util::isSyclFunction(const FunctionDecl *FD, StringRef Name) {
53025302
if (DC->isTranslationUnit())
53035303
return false;
53045304

5305-
std::array<DeclContextDesc, 2> Scopes = {
5305+
std::array<DeclContextDesc, 2> ScopesSycl = {
53065306
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
53075307
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl")};
5308-
return matchContext(DC, Scopes);
5308+
std::array<DeclContextDesc, 5> ScopesOneapiExp = {
5309+
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
5310+
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
5311+
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
5312+
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
5313+
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "experimental")};
5314+
5315+
return matchContext(DC, ScopesSycl) || matchContext(DC, ScopesOneapiExp);
53095316
}
53105317

53115318
bool Util::isAccessorPropertyListType(QualType Ty) {

‎clang/test/CodeGenSYCL/Inputs/sycl.hpp

Lines changed: 14 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -120,13 +120,6 @@ struct no_alias {
120120
} // namespace oneapi
121121
} // namespace ext
122122

123-
namespace ext {
124-
namespace oneapi {
125-
template <typename... properties>
126-
class accessor_property_list {};
127-
} // namespace oneapi
128-
} // namespace ext
129-
130123
template <int dim>
131124
struct id {
132125
template <typename... T>
@@ -146,6 +139,20 @@ template <int dim> struct item {
146139
int Data;
147140
};
148141

142+
namespace ext {
143+
namespace oneapi {
144+
template <typename... properties>
145+
class accessor_property_list {};
146+
namespace experimental {
147+
template <int Dims> item<Dims>
148+
this_item() { return item<Dims>{}; }
149+
150+
template <int Dims> id<Dims>
151+
this_id() { return id<Dims>{}; }
152+
} // namespace experimental
153+
} // namespace oneapi
154+
} // namespace ext
155+
149156
template <int Dims> item<Dims>
150157
this_item() { return item<Dims>{}; }
151158

‎clang/test/CodeGenSYCL/parallel_for_this_item.cpp

Lines changed: 89 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12,10 +12,14 @@
1212
// CHECK-NEXT: const char* const kernel_names[] = {
1313
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3GNU",
1414
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3EMU",
15+
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3COW",
1516
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3OWL",
1617
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3RAT",
18+
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3CAT",
1719
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3FOX",
18-
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3BEE"
20+
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3PIG",
21+
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3BEE",
22+
// CHECK-NEXT: "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3DOG"
1923
// CHECK-NEXT: };
2024

2125
// CHECK:template <> struct KernelInfo<GNU> {
@@ -50,6 +54,22 @@
5054
// CHECK-NEXT: __SYCL_DLL_LOCAL
5155
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
5256
// CHECK-NEXT:};
57+
// CHECK-NEXT:template <> struct KernelInfo<COW> {
58+
// CHECK-NEXT: __SYCL_DLL_LOCAL
59+
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3COW"; }
60+
// CHECK-NEXT: __SYCL_DLL_LOCAL
61+
// CHECK-NEXT: static constexpr unsigned getNumParams() { return 0; }
62+
// CHECK-NEXT: __SYCL_DLL_LOCAL
63+
// CHECK-NEXT: static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
64+
// CHECK-NEXT: return kernel_signatures[i+0];
65+
// CHECK-NEXT: }
66+
// CHECK-NEXT: __SYCL_DLL_LOCAL
67+
// CHECK-NEXT: static constexpr bool isESIMD() { return 0; }
68+
// CHECK-NEXT: __SYCL_DLL_LOCAL
69+
// CHECK-NEXT: static constexpr bool callsThisItem() { return 1; }
70+
// CHECK-NEXT: __SYCL_DLL_LOCAL
71+
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
72+
// CHECK-NEXT:};
5373
// CHECK-NEXT:template <> struct KernelInfo<OWL> {
5474
// CHECK-NEXT: __SYCL_DLL_LOCAL
5575
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3OWL"; }
@@ -82,6 +102,22 @@
82102
// CHECK-NEXT: __SYCL_DLL_LOCAL
83103
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
84104
// CHECK-NEXT:};
105+
// CHECK-NEXT:template <> struct KernelInfo<CAT> {
106+
// CHECK-NEXT: __SYCL_DLL_LOCAL
107+
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3CAT"; }
108+
// CHECK-NEXT: __SYCL_DLL_LOCAL
109+
// CHECK-NEXT: static constexpr unsigned getNumParams() { return 0; }
110+
// CHECK-NEXT: __SYCL_DLL_LOCAL
111+
// CHECK-NEXT: static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
112+
// CHECK-NEXT: return kernel_signatures[i+0];
113+
// CHECK-NEXT: }
114+
// CHECK-NEXT: __SYCL_DLL_LOCAL
115+
// CHECK-NEXT: static constexpr bool isESIMD() { return 0; }
116+
// CHECK-NEXT: __SYCL_DLL_LOCAL
117+
// CHECK-NEXT: static constexpr bool callsThisItem() { return 1; }
118+
// CHECK-NEXT: __SYCL_DLL_LOCAL
119+
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
120+
// CHECK-NEXT:};
85121
// CHECK-NEXT:template <> struct KernelInfo<FOX> {
86122
// CHECK-NEXT: __SYCL_DLL_LOCAL
87123
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3FOX"; }
@@ -98,6 +134,22 @@
98134
// CHECK-NEXT: __SYCL_DLL_LOCAL
99135
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
100136
// CHECK-NEXT:};
137+
// CHECK-NEXT:template <> struct KernelInfo<PIG> {
138+
// CHECK-NEXT: __SYCL_DLL_LOCAL
139+
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3PIG"; }
140+
// CHECK-NEXT: __SYCL_DLL_LOCAL
141+
// CHECK-NEXT: static constexpr unsigned getNumParams() { return 0; }
142+
// CHECK-NEXT: __SYCL_DLL_LOCAL
143+
// CHECK-NEXT: static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
144+
// CHECK-NEXT: return kernel_signatures[i+0];
145+
// CHECK-NEXT: }
146+
// CHECK-NEXT: __SYCL_DLL_LOCAL
147+
// CHECK-NEXT: static constexpr bool isESIMD() { return 0; }
148+
// CHECK-NEXT: __SYCL_DLL_LOCAL
149+
// CHECK-NEXT: static constexpr bool callsThisItem() { return 0; }
150+
// CHECK-NEXT: __SYCL_DLL_LOCAL
151+
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
152+
// CHECK-NEXT:};
101153
// CHECK-NEXT:template <> struct KernelInfo<BEE> {
102154
// CHECK-NEXT: __SYCL_DLL_LOCAL
103155
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3BEE"; }
@@ -114,13 +166,31 @@
114166
// CHECK-NEXT: __SYCL_DLL_LOCAL
115167
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
116168
// CHECK-NEXT:};
169+
// CHECK-NEXT:template <> struct KernelInfo<DOG> {
170+
// CHECK-NEXT: __SYCL_DLL_LOCAL
171+
// CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3DOG"; }
172+
// CHECK-NEXT: __SYCL_DLL_LOCAL
173+
// CHECK-NEXT: static constexpr unsigned getNumParams() { return 0; }
174+
// CHECK-NEXT: __SYCL_DLL_LOCAL
175+
// CHECK-NEXT: static constexpr const kernel_param_desc_t& getParamDesc(unsigned i) {
176+
// CHECK-NEXT: return kernel_signatures[i+0];
177+
// CHECK-NEXT: }
178+
// CHECK-NEXT: __SYCL_DLL_LOCAL
179+
// CHECK-NEXT: static constexpr bool isESIMD() { return 0; }
180+
// CHECK-NEXT: __SYCL_DLL_LOCAL
181+
// CHECK-NEXT: static constexpr bool callsThisItem() { return 1; }
182+
// CHECK-NEXT: __SYCL_DLL_LOCAL
183+
// CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; }
184+
// CHECK-NEXT:};
117185

118186
#include "sycl.hpp"
119187

120188
using namespace cl::sycl;
121189

122190
SYCL_EXTERNAL item<1> g() { return this_item<1>(); }
123191
SYCL_EXTERNAL item<1> f() { return g(); }
192+
SYCL_EXTERNAL item<1> s() { return ext::oneapi::experimental::this_item<1>(); }
193+
SYCL_EXTERNAL item<1> h() { return s(); }
124194

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

215+
// This kernel calls sycl::ext::oneapi::experimental::this_item
216+
cgh.parallel_for<class COW>(range<1>(1), [=](::item<1> I) {
217+
ext::oneapi::experimental::this_item<1>();
218+
});
219+
145220
// This kernel does not call sycl::this_item
146221
cgh.parallel_for<class OWL>(range<1>(1), [=](id<1> I) {
147222
class C c;
@@ -151,11 +226,24 @@ int main() {
151226
// This kernel calls sycl::this_item
152227
cgh.parallel_for<class RAT>(range<1>(1), [=](id<1> I) { f(); });
153228

229+
// This kernel calls sycl::ext::oneapi::experimental::this_item
230+
cgh.parallel_for<class CAT>(range<1>(1), [=](id<1> I) { h(); });
231+
154232
// This kernel does not call sycl::this_item, but does call this_id
155233
cgh.parallel_for<class FOX>(range<1>(1), [=](id<1> I) { this_id<1>(); });
156234

235+
// This kernel calls sycl::ext::oneapi::experimental::this_id
236+
cgh.parallel_for<class PIG>(range<1>(1), [=](id<1> I) {
237+
ext::oneapi::experimental::this_id<1>();
238+
});
239+
157240
// This kernel calls sycl::this_item
158241
cgh.parallel_for<class BEE>(range<1>(1), [=](auto I) { this_item<1>(); });
242+
243+
// This kernel calls sycl::ext::oneapi::experimental::this_item
244+
cgh.parallel_for<class DOG>(range<1>(1), [=](auto I) {
245+
ext::oneapi::experimental::this_item<1>();
246+
});
159247
});
160248

161249
return 0;

0 commit comments

Comments
 (0)