From e574aa7c56113caa9b3c767c61cd915bba172655 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 6 Nov 2024 09:06:20 -0800 Subject: [PATCH 01/24] Add indeterminate constructor to work group memory interface --- .../ext/oneapi/experimental/work_group_memory.hpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index 7870ebd3ca73e..8fff0a7df3c23 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -36,6 +36,9 @@ class work_group_memory_impl { } // namespace detail namespace ext::oneapi::experimental { +struct indeterminate_t {}; +inline constexpr indeterminate_t indeterminate; + template class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory : sycl::detail::work_group_memory_impl { @@ -47,7 +50,13 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory value_type, access::address_space::local_space>::type *; public: - work_group_memory() = default; +// Frontend requries special types to have a default constructor in device +// compilation mode in order to have a uniform way of initializing an object of +// special type to then call the __init method on it. This is purely an +// implementation detail and not part of the spec. +#ifdef __SYCL_DEVICE_ONLY__ + work_group_memory(const indeterminate_t &); +#endif work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Wed, 6 Nov 2024 09:11:34 -0800 Subject: [PATCH 02/24] Update tests to use indeterminate constructor instead of default one --- sycl/test-e2e/WorkGroupMemory/swap_test.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/WorkGroupMemory/swap_test.cpp b/sycl/test-e2e/WorkGroupMemory/swap_test.cpp index 13fbde212a47d..e70c7ce478df6 100644 --- a/sycl/test-e2e/WorkGroupMemory/swap_test.cpp +++ b/sycl/test-e2e/WorkGroupMemory/swap_test.cpp @@ -49,7 +49,7 @@ template void swap_scalar(T &a, T &b) { syclexp::work_group_memory temp{cgh}; sycl::nd_range<1> ndr{size, wgsize}; cgh.parallel_for(ndr, [=](sycl::nd_item<1> it) { - syclexp::work_group_memory temp2; + syclexp::work_group_memory temp2{ syclexp::indeterminate }; temp2 = temp; // temp and temp2 have the same underlying data temp = acc_a[0]; acc_a[0] = acc_b[0]; @@ -264,7 +264,7 @@ void swap_array_2d(T (&a)[N][N], T (&b)[N][N], size_t batch_size) { const auto j = it.get_global_id()[1]; temp[i][j] = acc_a[i][j]; acc_a[i][j] = acc_b[i][j]; - syclexp::work_group_memory temp2; + syclexp::work_group_memory temp2{ syclexp::indeterminate }; temp2 = temp; acc_b[i][j] = temp2[i][j]; }); From 543aa6f4e2d94530d8171c10e008ef83b5908fca Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 6 Nov 2024 09:13:53 -0800 Subject: [PATCH 03/24] Add indeterminate constructor to work group memory interface --- .../include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index 8fff0a7df3c23..cb89e676de314 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -55,8 +55,9 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory // special type to then call the __init method on it. This is purely an // implementation detail and not part of the spec. #ifdef __SYCL_DEVICE_ONLY__ - work_group_memory(const indeterminate_t &); + work_group_memory() = default; #endif + work_group_memory(const indeterminate_t &); work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Wed, 6 Nov 2024 12:14:24 -0500 Subject: [PATCH 04/24] Fix typo --- sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index cb89e676de314..d4ff84d278eb7 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -50,7 +50,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory value_type, access::address_space::local_space>::type *; public: -// Frontend requries special types to have a default constructor in device +// Frontend requires special types to have a default constructor in device // compilation mode in order to have a uniform way of initializing an object of // special type to then call the __init method on it. This is purely an // implementation detail and not part of the spec. From 65a8e239c2b820356dc67f3102ee6031573ac8fc Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 6 Nov 2024 12:15:08 -0500 Subject: [PATCH 05/24] Fix compiler error --- sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index d4ff84d278eb7..0f49792931164 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -57,7 +57,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory #ifdef __SYCL_DEVICE_ONLY__ work_group_memory() = default; #endif - work_group_memory(const indeterminate_t &); + work_group_memory(const indeterminate_t &indeterminate) {}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Wed, 6 Nov 2024 12:20:27 -0500 Subject: [PATCH 06/24] Formatting changes --- sycl/test-e2e/WorkGroupMemory/swap_test.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/WorkGroupMemory/swap_test.cpp b/sycl/test-e2e/WorkGroupMemory/swap_test.cpp index e70c7ce478df6..7552774edcbbf 100644 --- a/sycl/test-e2e/WorkGroupMemory/swap_test.cpp +++ b/sycl/test-e2e/WorkGroupMemory/swap_test.cpp @@ -49,7 +49,7 @@ template void swap_scalar(T &a, T &b) { syclexp::work_group_memory temp{cgh}; sycl::nd_range<1> ndr{size, wgsize}; cgh.parallel_for(ndr, [=](sycl::nd_item<1> it) { - syclexp::work_group_memory temp2{ syclexp::indeterminate }; + syclexp::work_group_memory temp2{syclexp::indeterminate}; temp2 = temp; // temp and temp2 have the same underlying data temp = acc_a[0]; acc_a[0] = acc_b[0]; @@ -264,7 +264,7 @@ void swap_array_2d(T (&a)[N][N], T (&b)[N][N], size_t batch_size) { const auto j = it.get_global_id()[1]; temp[i][j] = acc_a[i][j]; acc_a[i][j] = acc_b[i][j]; - syclexp::work_group_memory temp2{ syclexp::indeterminate }; + syclexp::work_group_memory temp2{syclexp::indeterminate}; temp2 = temp; acc_b[i][j] = temp2[i][j]; }); From d6709edf676dadd6536b07fb530d5061355b0877 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 6 Nov 2024 12:39:23 -0500 Subject: [PATCH 07/24] Fix unused variable warning --- sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index 0f49792931164..ed99115e0a391 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -57,7 +57,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory #ifdef __SYCL_DEVICE_ONLY__ work_group_memory() = default; #endif - work_group_memory(const indeterminate_t &indeterminate) {}; + work_group_memory(const indeterminate_t &) {}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Fri, 8 Nov 2024 14:40:16 -0800 Subject: [PATCH 08/24] Modify handling of SYCL special types to account for default constructor access specifier --- clang/lib/Sema/SemaSYCL.cpp | 19 +++++++++++++++++-- 1 file changed, 17 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e45b038273d77..ac5ef099be76f 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3950,13 +3950,28 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } // Default inits the type, then calls the init-method in the body. + // A type may not have a public private default constructor as per its spec so + // typically the default constructor will be private and in such cases we must + // manually override the access specifier from private to public just for the + // duration of this default initialization. bool handleSpecialType(FieldDecl *FD, QualType Ty) { + const auto *RecordDecl = Ty->getAsCXXRecordDecl(); + AccessSpecifier default_constructor_access; + CXXConstructorDecl *default_constructor; + std::for_each(RecordDecl->ctor_begin(), RecordDecl->ctor_end(), + [&](auto elem) { + if (elem->isDefaultConstructor()) { + default_constructor_access = elem->getAccess(); + elem->setAccess(AS_public); + default_constructor = elem; + } + }); + addFieldInit(FD, Ty, std::nullopt, InitializationKind::CreateDefault(KernelCallerSrcLoc)); - + default_constructor->setAccess(default_constructor_access); addFieldMemberExpr(FD, Ty); - const auto *RecordDecl = Ty->getAsCXXRecordDecl(); createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); CXXMethodDecl *FinalizeMethod = getMethodByName(RecordDecl, FinalizeMethodName); From f551ba4de3add562bfb8da4c47bcbfc235186e6f Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 8 Nov 2024 14:41:35 -0800 Subject: [PATCH 09/24] Modify handling of SYCL special types to account for default constructor access specifier --- clang/lib/Sema/SemaSYCL.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ac5ef099be76f..86a2d8a7d8edd 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3950,10 +3950,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } // Default inits the type, then calls the init-method in the body. - // A type may not have a public private default constructor as per its spec so - // typically the default constructor will be private and in such cases we must - // manually override the access specifier from private to public just for the - // duration of this default initialization. + // A type may not have a public default constructor as per its spec so + // typically if this is the case the default constructor will be private and + // in such cases we must manually override the access specifier from private + // to public just for the duration of this default initialization. bool handleSpecialType(FieldDecl *FD, QualType Ty) { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); AccessSpecifier default_constructor_access; From 172d717e70b8fb07b70d72cb7c3c5e44d1eec548 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 8 Nov 2024 14:42:54 -0800 Subject: [PATCH 10/24] Modify handling of SYCL special types to account for default constructor access specifier --- .../include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index ed99115e0a391..006af453fe4fb 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -49,7 +49,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory using decoratedPtr = typename sycl::detail::DecoratedType< value_type, access::address_space::local_space>::type *; -public: // Frontend requires special types to have a default constructor in device // compilation mode in order to have a uniform way of initializing an object of // special type to then call the __init method on it. This is purely an @@ -57,6 +56,8 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory #ifdef __SYCL_DEVICE_ONLY__ work_group_memory() = default; #endif + +public: work_group_memory(const indeterminate_t &) {}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; From e531b054ba19d6262cff301b80e8b6dfc5f79d0b Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 8 Nov 2024 14:53:51 -0800 Subject: [PATCH 11/24] Make default constructor private --- .../ext/oneapi/experimental/work_group_memory.hpp | 12 +++++------- 1 file changed, 5 insertions(+), 7 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index 006af453fe4fb..5905209ca99f8 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -49,16 +49,14 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory using decoratedPtr = typename sycl::detail::DecoratedType< value_type, access::address_space::local_space>::type *; -// Frontend requires special types to have a default constructor in device -// compilation mode in order to have a uniform way of initializing an object of -// special type to then call the __init method on it. This is purely an -// implementation detail and not part of the spec. -#ifdef __SYCL_DEVICE_ONLY__ + // Frontend requires special types to have a default constructor in order to + // have a uniform way of initializing an object of special type to then call + // the __init method on it. This is purely an implementation detail and not + // part of the spec. work_group_memory() = default; -#endif public: - work_group_memory(const indeterminate_t &) {}; + work_group_memory(const indeterminate_t &){}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Fri, 8 Nov 2024 18:27:01 -0500 Subject: [PATCH 12/24] Formatting changes --- sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index 5905209ca99f8..bf0c24349c6ce 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -56,7 +56,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory work_group_memory() = default; public: - work_group_memory(const indeterminate_t &){}; + work_group_memory(const indeterminate_t &) {}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Fri, 8 Nov 2024 20:34:42 -0500 Subject: [PATCH 13/24] Change naming convention to match rest of the code --- clang/lib/Sema/SemaSYCL.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 86a2d8a7d8edd..01356085f1712 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3956,20 +3956,20 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // to public just for the duration of this default initialization. bool handleSpecialType(FieldDecl *FD, QualType Ty) { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); - AccessSpecifier default_constructor_access; - CXXConstructorDecl *default_constructor; + AccessSpecifier DefaultConstructorAccess; + CXXConstructorDecl *DefaultConstructor; std::for_each(RecordDecl->ctor_begin(), RecordDecl->ctor_end(), [&](auto elem) { if (elem->isDefaultConstructor()) { - default_constructor_access = elem->getAccess(); + DefaultConstructorAccess = elem->getAccess(); elem->setAccess(AS_public); - default_constructor = elem; + DefaultConstructor = elem; } }); addFieldInit(FD, Ty, std::nullopt, InitializationKind::CreateDefault(KernelCallerSrcLoc)); - default_constructor->setAccess(default_constructor_access); + DefaultConstructor->setAccess(DefaultConstructorAccess); addFieldMemberExpr(FD, Ty); createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); From eb31ca09abe296313f4f5e92f8df62b20e692bd2 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Mon, 11 Nov 2024 10:56:27 -0800 Subject: [PATCH 14/24] Use std::find_if instead of std::for_each to find default constructor --- clang/lib/Sema/SemaSYCL.cpp | 15 +++++---------- 1 file changed, 5 insertions(+), 10 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 01356085f1712..8a02e815a62a5 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3957,16 +3957,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { bool handleSpecialType(FieldDecl *FD, QualType Ty) { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); AccessSpecifier DefaultConstructorAccess; - CXXConstructorDecl *DefaultConstructor; - std::for_each(RecordDecl->ctor_begin(), RecordDecl->ctor_end(), - [&](auto elem) { - if (elem->isDefaultConstructor()) { - DefaultConstructorAccess = elem->getAccess(); - elem->setAccess(AS_public); - DefaultConstructor = elem; - } - }); - + auto DefaultConstructor = + std::find_if(RecordDecl->ctor_begin(), RecordDecl->ctor_end(), + [](auto it) { return it->isDefaultConstructor(); }); + DefaultConstructorAccess = DefaultConstructor->getAccess(); + DefaultConstructor->setAccess(AS_public); addFieldInit(FD, Ty, std::nullopt, InitializationKind::CreateDefault(KernelCallerSrcLoc)); DefaultConstructor->setAccess(DefaultConstructorAccess); From 54abcd03b51bbdb2b971748c1ba49a05ea8e9526 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 12 Nov 2024 14:09:31 -0800 Subject: [PATCH 15/24] Revert back changes to default constructor --- clang/lib/Sema/SemaSYCL.cpp | 14 ++------------ .../ext/oneapi/experimental/work_group_memory.hpp | 9 +++++---- 2 files changed, 7 insertions(+), 16 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8a02e815a62a5..e45b038273d77 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3950,23 +3950,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } // Default inits the type, then calls the init-method in the body. - // A type may not have a public default constructor as per its spec so - // typically if this is the case the default constructor will be private and - // in such cases we must manually override the access specifier from private - // to public just for the duration of this default initialization. bool handleSpecialType(FieldDecl *FD, QualType Ty) { - const auto *RecordDecl = Ty->getAsCXXRecordDecl(); - AccessSpecifier DefaultConstructorAccess; - auto DefaultConstructor = - std::find_if(RecordDecl->ctor_begin(), RecordDecl->ctor_end(), - [](auto it) { return it->isDefaultConstructor(); }); - DefaultConstructorAccess = DefaultConstructor->getAccess(); - DefaultConstructor->setAccess(AS_public); addFieldInit(FD, Ty, std::nullopt, InitializationKind::CreateDefault(KernelCallerSrcLoc)); - DefaultConstructor->setAccess(DefaultConstructorAccess); + addFieldMemberExpr(FD, Ty); + const auto *RecordDecl = Ty->getAsCXXRecordDecl(); createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); CXXMethodDecl *FinalizeMethod = getMethodByName(RecordDecl, FinalizeMethodName); diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index bf0c24349c6ce..a42ffc0210b76 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -49,14 +49,15 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory using decoratedPtr = typename sycl::detail::DecoratedType< value_type, access::address_space::local_space>::type *; +public: // Frontend requires special types to have a default constructor in order to // have a uniform way of initializing an object of special type to then call - // the __init method on it. This is purely an implementation detail and not - // part of the spec. + // the __init method on it. This is currently not part of the spec. + // TODO: Remove this once https://github.com/intel/llvm/issues/16061 is + // closed. work_group_memory() = default; -public: - work_group_memory(const indeterminate_t &) {}; + work_group_memory(const indeterminate_t &){}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Tue, 12 Nov 2024 17:18:47 -0500 Subject: [PATCH 16/24] Fix formatting errors --- sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index a42ffc0210b76..12a7e112a9bf3 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -57,7 +57,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory // closed. work_group_memory() = default; - work_group_memory(const indeterminate_t &){}; + work_group_memory(const indeterminate_t &) {}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Thu, 14 Nov 2024 06:45:26 -0800 Subject: [PATCH 17/24] Revert "Fix formatting errors" This reverts commit 21c27fae5ee584be990d728d5dcad41076f27de3. --- sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index 12a7e112a9bf3..a42ffc0210b76 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -57,7 +57,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory // closed. work_group_memory() = default; - work_group_memory(const indeterminate_t &) {}; + work_group_memory(const indeterminate_t &){}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Thu, 14 Nov 2024 06:46:08 -0800 Subject: [PATCH 18/24] Revert "Revert "Fix formatting errors"" This reverts commit 8c94bfc356b884dbddeed6431e0740957f77ea88. --- sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index a42ffc0210b76..12a7e112a9bf3 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -57,7 +57,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory // closed. work_group_memory() = default; - work_group_memory(const indeterminate_t &){}; + work_group_memory(const indeterminate_t &) {}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Thu, 14 Nov 2024 06:49:41 -0800 Subject: [PATCH 19/24] Revert "Use std::find_if instead of std::for_each to find default constructor" This reverts commit eb31ca09abe296313f4f5e92f8df62b20e692bd2. --- clang/lib/Sema/SemaSYCL.cpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e45b038273d77..4da7cb0c61763 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3951,6 +3951,18 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // Default inits the type, then calls the init-method in the body. bool handleSpecialType(FieldDecl *FD, QualType Ty) { + const auto *RecordDecl = Ty->getAsCXXRecordDecl(); + AccessSpecifier DefaultConstructorAccess; + CXXConstructorDecl *DefaultConstructor; + std::for_each(RecordDecl->ctor_begin(), RecordDecl->ctor_end(), + [&](auto elem) { + if (elem->isDefaultConstructor()) { + DefaultConstructorAccess = elem->getAccess(); + elem->setAccess(AS_public); + DefaultConstructor = elem; + } + }); + addFieldInit(FD, Ty, std::nullopt, InitializationKind::CreateDefault(KernelCallerSrcLoc)); From be386bcc7843674a46aaf216760d1ce3a0804b5f Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 14 Nov 2024 06:54:15 -0800 Subject: [PATCH 20/24] Revert "Revert back changes to default constructor" This reverts commit 54abcd03b51bbdb2b971748c1ba49a05ea8e9526. --- clang/lib/Sema/SemaSYCL.cpp | 23 +++++++++---------- .../oneapi/experimental/work_group_memory.hpp | 7 +++--- 2 files changed, 14 insertions(+), 16 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 4da7cb0c61763..c8f70e2ef88b7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3950,25 +3950,24 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } // Default inits the type, then calls the init-method in the body. + // A type may not have a public default constructor as per its spec so + // typically if this is the case the default constructor will be private and + // in such cases we must manually override the access specifier from private + // to public just for the duration of this default initialization. bool handleSpecialType(FieldDecl *FD, QualType Ty) { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); AccessSpecifier DefaultConstructorAccess; - CXXConstructorDecl *DefaultConstructor; - std::for_each(RecordDecl->ctor_begin(), RecordDecl->ctor_end(), - [&](auto elem) { - if (elem->isDefaultConstructor()) { - DefaultConstructorAccess = elem->getAccess(); - elem->setAccess(AS_public); - DefaultConstructor = elem; - } - }); - + auto DefaultConstructor = + std::find_if(RecordDecl->ctor_begin(), RecordDecl->ctor_end(), + [](auto it) { return it->isDefaultConstructor(); }); + DefaultConstructorAccess = DefaultConstructor->getAccess(); + DefaultConstructor->setAccess(AS_public); + addFieldInit(FD, Ty, std::nullopt, InitializationKind::CreateDefault(KernelCallerSrcLoc)); - + DefaultConstructor->setAccess(DefaultConstructorAccess); addFieldMemberExpr(FD, Ty); - const auto *RecordDecl = Ty->getAsCXXRecordDecl(); createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); CXXMethodDecl *FinalizeMethod = getMethodByName(RecordDecl, FinalizeMethodName); diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index 12a7e112a9bf3..bf0c24349c6ce 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -49,14 +49,13 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory using decoratedPtr = typename sycl::detail::DecoratedType< value_type, access::address_space::local_space>::type *; -public: // Frontend requires special types to have a default constructor in order to // have a uniform way of initializing an object of special type to then call - // the __init method on it. This is currently not part of the spec. - // TODO: Remove this once https://github.com/intel/llvm/issues/16061 is - // closed. + // the __init method on it. This is purely an implementation detail and not + // part of the spec. work_group_memory() = default; +public: work_group_memory(const indeterminate_t &) {}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; From 17639aea07970917c42908b0ce6bb11f5e2ff2f0 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 14 Nov 2024 09:05:52 -0800 Subject: [PATCH 21/24] Bypass default constructor access specifier in frontend and add tests for inheriting from work group memory --- clang/lib/Sema/SemaSYCL.cpp | 37 ++++++++++------ clang/test/SemaSYCL/Inputs/sycl.hpp | 17 ++++++++ .../work_group_memory_inheritance.cpp | 43 +++++++++++++++++++ .../oneapi/experimental/work_group_memory.hpp | 10 +++-- 4 files changed, 91 insertions(+), 16 deletions(-) create mode 100644 clang/test/SemaSYCL/work_group_memory_inheritance.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c8f70e2ef88b7..25aef7fec7781 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3981,9 +3981,17 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } bool handleSpecialType(const CXXBaseSpecifier &BS, QualType Ty) { - const auto *RecordDecl = Ty->getAsCXXRecordDecl(); + const auto *BaseRecordDecl = BS.getType()->getAsCXXRecordDecl(); + AccessSpecifier DefaultConstructorAccess; + auto DefaultConstructor = + std::find_if(BaseRecordDecl->ctor_begin(), BaseRecordDecl->ctor_end(), + [](auto it) { return it->isDefaultConstructor(); }); + DefaultConstructorAccess = DefaultConstructor->getAccess(); + DefaultConstructor->setAccess(AS_public); + addBaseInit(BS, Ty, InitializationKind::CreateDefault(KernelCallerSrcLoc)); - createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); + DefaultConstructor->setAccess(DefaultConstructorAccess); + createSpecialMethodCall(BaseRecordDecl, getInitMethodName(), BodyStmts); return true; } @@ -4680,16 +4688,21 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { bool handleSyclSpecialType(const CXXRecordDecl *RD, const CXXBaseSpecifier &BC, QualType FieldTy) final { - const auto *AccTy = - cast(FieldTy->getAsRecordDecl()); - assert(AccTy->getTemplateArgs().size() >= 2 && - "Incorrect template args for Accessor Type"); - int Dims = static_cast( - AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); - int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11); - Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, - CurOffset + - offsetOf(RD, BC.getType()->getAsCXXRecordDecl())); + if (isSyclAccessorType(FieldTy)) { + const auto *AccTy = + cast(FieldTy->getAsRecordDecl()); + assert(AccTy->getTemplateArgs().size() >= 2 && + "Incorrect template args for Accessor Type"); + int Dims = static_cast( + AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); + int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11); + Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, + CurOffset + + offsetOf(RD, BC.getType()->getAsCXXRecordDecl())); + } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) { + addParam(FieldTy, SYCLIntegrationHeader::kind_work_group_memory, + offsetOf(RD, BC.getType()->getAsCXXRecordDecl())); + } return true; } diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 337320764de2f..5df1550ed2dcb 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -448,6 +448,23 @@ class __SYCL_TYPE(multi_ptr) multi_ptr { pointer_t m_Pointer; }; +// Dummy implementation of work_group_memory for use in SemaSYCL tests. +template +class __attribute__((sycl_special_class)) +__SYCL_TYPE(work_group_memory) work_group_memory { + +// Default constructor for objects later initialized with __init member. + work_group_memory() = default; + +public: + work_group_memory(handler &CGH) {} + + void __init(__attribute((opencl_local)) DataT *Ptr) { this->Ptr = Ptr; } + void use() const {} +private: + __attribute((opencl_local)) DataT *Ptr; +}; + namespace ext { namespace oneapi { namespace experimental { diff --git a/clang/test/SemaSYCL/work_group_memory_inheritance.cpp b/clang/test/SemaSYCL/work_group_memory_inheritance.cpp new file mode 100644 index 0000000000000..47c7db1f3ca52 --- /dev/null +++ b/clang/test/SemaSYCL/work_group_memory_inheritance.cpp @@ -0,0 +1,43 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -sycl-std=2020 %s | FileCheck %s + +// Check that AST is correctly generated for kernel arguments that inherit from work group memory. + +#include "sycl.hpp" + +sycl::queue myQueue; + +struct WorkGroupMemoryDerived : + sycl::work_group_memory { +}; + +int main() { + myQueue.submit([&](sycl::handler &h) { + WorkGroupMemoryDerived DerivedObject{ h }; + h.parallel_for([=] { + DerivedObject.use(); + }); + }); + return 0; +} + +// CHECK: FunctionDecl {{.*}}kernel 'void (__local int *)' +// CHECK-NEXT: ParmVarDecl {{.*}}used _arg__base '__local int *' +// CHECK-NEXT: CompoundStmt {{.*}} +// CHECK-NEXT: DeclStmt {{.*}} +// CHECK-NEXT: VarDecl {{.*}} used __SYCLKernel {{.*}} cinit +// CHECK-NEXT: InitListExpr {{.*}} +// CHECK-NEXT: InitListExpr {{.*}} 'WorkGroupMemoryDerived' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::work_group_memory' 'void () noexcept' +// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (__local int *)' lvalue .__init {{.*}} +// CHECK-NEXT: MemberExpr {{.*}} 'WorkGroupMemoryDerived' lvalue .DerivedObject +// CHECK-NEXT: DeclRefExpr {{.*}} lvalue Var {{.*}} '__SYCLKernel' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__local int *' +// CHECK-NEXT: DeclRefExpr {{.*}} '__local int *' lvalue ParmVar {{.*}} '_arg__base' '__local int *' +// CHECK-NEXT: CompoundStmt {{.*}} +// CHECK-NEXT: CXXOperatorCallExpr {{.*}} 'void' '()' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'auto (*)() const -> void' +// CHECK-NEXT: DeclRefExpr {{.*}}'auto () const -> void' lvalue CXXMethod {{.*}} 'operator()' 'auto () const -> void' +// CHECK-NEXT: ImplicitCastExpr {{.*}} +// CHECK-NEXT: DeclRefExpr {{.*}}lvalue Var {{.*}} '__SYCLKernel' + diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index bf0c24349c6ce..6d0021f7cb3e8 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -55,8 +55,12 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory // part of the spec. work_group_memory() = default; +#ifdef __SYCL_DEVICE_ONLY__ + void __init(decoratedPtr ptr) { this->ptr = ptr; } +#endif + public: - work_group_memory(const indeterminate_t &) {}; + work_group_memory(const indeterminate_t &){}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template ptr = ptr; } -#endif + private: decoratedPtr ptr; }; From 6cc99d34cb4bcb31aa3b131fd71ed4000ef4fe7e Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 14 Nov 2024 09:09:22 -0800 Subject: [PATCH 22/24] Add TODO to revisit the handling of special types --- clang/lib/Sema/SemaSYCL.cpp | 2 ++ sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 2 ++ 2 files changed, 4 insertions(+) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 25aef7fec7781..9175fb84b2b7c 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3954,6 +3954,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // typically if this is the case the default constructor will be private and // in such cases we must manually override the access specifier from private // to public just for the duration of this default initialization. + // TODO: Revisit this approach once https://github.com/intel/llvm/issues/16061 + // is closed. bool handleSpecialType(FieldDecl *FD, QualType Ty) { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); AccessSpecifier DefaultConstructorAccess; diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index 6d0021f7cb3e8..c98696b81e99c 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -53,6 +53,8 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory // have a uniform way of initializing an object of special type to then call // the __init method on it. This is purely an implementation detail and not // part of the spec. + // TODO: Revisit this once https://github.com/intel/llvm/issues/16061 is + // closed. work_group_memory() = default; #ifdef __SYCL_DEVICE_ONLY__ From 90bceb08d1c8d7527c0f86793a65fcf4bab6cd5e Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Thu, 14 Nov 2024 09:22:53 -0800 Subject: [PATCH 23/24] Formatting changes --- clang/lib/Sema/SemaSYCL.cpp | 2 +- sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9175fb84b2b7c..7d53638c8eff3 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3964,7 +3964,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { [](auto it) { return it->isDefaultConstructor(); }); DefaultConstructorAccess = DefaultConstructor->getAccess(); DefaultConstructor->setAccess(AS_public); - + addFieldInit(FD, Ty, std::nullopt, InitializationKind::CreateDefault(KernelCallerSrcLoc)); DefaultConstructor->setAccess(DefaultConstructorAccess); diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index c98696b81e99c..d7c9138a2a23d 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -62,7 +62,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory #endif public: - work_group_memory(const indeterminate_t &){}; + work_group_memory(const indeterminate_t &) {}; work_group_memory(const work_group_memory &rhs) = default; work_group_memory &operator=(const work_group_memory &rhs) = default; template Date: Thu, 14 Nov 2024 16:51:18 -0500 Subject: [PATCH 24/24] Fix regex in filecheck pattern --- clang/test/SemaSYCL/work_group_memory_inheritance.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/work_group_memory_inheritance.cpp b/clang/test/SemaSYCL/work_group_memory_inheritance.cpp index 47c7db1f3ca52..6f1f6badbdc59 100644 --- a/clang/test/SemaSYCL/work_group_memory_inheritance.cpp +++ b/clang/test/SemaSYCL/work_group_memory_inheritance.cpp @@ -20,7 +20,7 @@ int main() { return 0; } -// CHECK: FunctionDecl {{.*}}kernel 'void (__local int *)' +// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (__local int *)' // CHECK-NEXT: ParmVarDecl {{.*}}used _arg__base '__local int *' // CHECK-NEXT: CompoundStmt {{.*}} // CHECK-NEXT: DeclStmt {{.*}}