diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 41fd64edb9b69..495de35db6914 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -5302,10 +5302,17 @@ bool Util::isSyclFunction(const FunctionDecl *FD, StringRef Name) { if (DC->isTranslationUnit()) return false; - std::array Scopes = { + std::array ScopesSycl = { Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"), Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl")}; - return matchContext(DC, Scopes); + std::array 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) { diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 5fed15a1a6acc..cd0dfa8b9b71a 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -120,13 +120,6 @@ struct no_alias { } // namespace oneapi } // namespace ext -namespace ext { -namespace oneapi { -template -class accessor_property_list {}; -} // namespace oneapi -} // namespace ext - template struct id { template @@ -146,6 +139,20 @@ template struct item { int Data; }; +namespace ext { +namespace oneapi { +template +class accessor_property_list {}; +namespace experimental { +template item +this_item() { return item{}; } + +template id +this_id() { return id{}; } +} // namespace experimental +} // namespace oneapi +} // namespace ext + template item this_item() { return item{}; } diff --git a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp index 4e716b1411d1a..cc6507b20bb34 100755 --- a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp +++ b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp @@ -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 { @@ -50,6 +54,22 @@ // CHECK-NEXT: __SYCL_DLL_LOCAL // CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; } // CHECK-NEXT:}; +// CHECK-NEXT:template <> struct KernelInfo { +// 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 { // CHECK-NEXT: __SYCL_DLL_LOCAL // CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3OWL"; } @@ -82,6 +102,22 @@ // CHECK-NEXT: __SYCL_DLL_LOCAL // CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; } // CHECK-NEXT:}; +// CHECK-NEXT:template <> struct KernelInfo { +// 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 { // CHECK-NEXT: __SYCL_DLL_LOCAL // CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3FOX"; } @@ -98,6 +134,22 @@ // CHECK-NEXT: __SYCL_DLL_LOCAL // CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; } // CHECK-NEXT:}; +// CHECK-NEXT:template <> struct KernelInfo { +// 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 { // CHECK-NEXT: __SYCL_DLL_LOCAL // CHECK-NEXT: static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3BEE"; } @@ -114,6 +166,22 @@ // CHECK-NEXT: __SYCL_DLL_LOCAL // CHECK-NEXT: static constexpr bool callsAnyThisFreeFunction() { return 1; } // CHECK-NEXT:}; +// CHECK-NEXT:template <> struct KernelInfo { +// 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" @@ -121,6 +189,8 @@ 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 item this_item(int i) { return item<1>{i}; } @@ -142,6 +212,11 @@ int main() { cgh.parallel_for(range<1>(1), [=](::item<1> I) { this_item<1>(); }); + // This kernel calls sycl::ext::oneapi::experimental::this_item + cgh.parallel_for(range<1>(1), [=](::item<1> I) { + ext::oneapi::experimental::this_item<1>(); + }); + // This kernel does not call sycl::this_item cgh.parallel_for(range<1>(1), [=](id<1> I) { class C c; @@ -151,11 +226,24 @@ int main() { // This kernel calls sycl::this_item cgh.parallel_for(range<1>(1), [=](id<1> I) { f(); }); + // This kernel calls sycl::ext::oneapi::experimental::this_item + cgh.parallel_for(range<1>(1), [=](id<1> I) { h(); }); + // This kernel does not call sycl::this_item, but does call this_id cgh.parallel_for(range<1>(1), [=](id<1> I) { this_id<1>(); }); + // This kernel calls sycl::ext::oneapi::experimental::this_id + cgh.parallel_for(range<1>(1), [=](id<1> I) { + ext::oneapi::experimental::this_id<1>(); + }); + // This kernel calls sycl::this_item cgh.parallel_for(range<1>(1), [=](auto I) { this_item<1>(); }); + + // This kernel calls sycl::ext::oneapi::experimental::this_item + cgh.parallel_for(range<1>(1), [=](auto I) { + ext::oneapi::experimental::this_item<1>(); + }); }); return 0;