[SYCL] Adding support for 'sycl_special_class' attribute. #3892
Fixed build errors
Signed-off-by: Zahira Ammarguellat <zahira.ammarguellat@intel.com>
- Loading branch information
commit 425bb978c0674123d179e8be0d8b6316e247aa44
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,80 @@ | ||
| // RUN: %clang_cc1 -S -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s | ||
|
|
||
|
There was a problem hiding this comment. Do we need this test? Is this checking anything additional to what we already check in existing accessor/sampler/stream tests? If we're keeping test, please add a comment describing test There was a problem hiding this comment. Please address comment. I don't think this test is required since we already have other accessor/sampler/stream tests. Are you testing anything additional here? |
||
| #include "Inputs/sycl.hpp" | ||
|
|
||
| cl::sycl::queue myQueue; | ||
| cl::sycl::handler SH; | ||
|
|
||
| class AccessorBase { | ||
| int A; | ||
| public: | ||
| cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, | ||
| cl::sycl::access::target::local> | ||
| acc; | ||
| }; | ||
|
|
||
| class accessor { | ||
| public: | ||
| int field; | ||
| }; | ||
|
|
||
| class stream { | ||
| public: | ||
| int field; | ||
| }; | ||
|
|
||
| class sampler { | ||
| public: | ||
| int field; | ||
| }; | ||
|
|
||
| int main() { | ||
|
|
||
| AccessorBase Accessor1; | ||
| accessor Accessor2 = {1}; | ||
| stream Stream2; | ||
| cl::sycl::sampler Sampler1; | ||
| sampler Sampler2; | ||
|
|
||
| myQueue.submit([&](cl::sycl::handler &h) { | ||
| h.single_task<class kernel_function1>([=]() { | ||
| Accessor1.acc.use(); | ||
| }); | ||
| h.single_task<class kernel_function2>([=]() { | ||
| int a = Accessor2.field; | ||
| }); | ||
|
|
||
| cl::sycl::stream Stream1{0, 0, SH}; | ||
| h.single_task<class kernel_function3>([=]() { | ||
| int a = Stream2.field; | ||
| }); | ||
|
|
||
| h.single_task<class kernelfunction4>([=] { | ||
| Sampler1.use(); | ||
| }); | ||
|
|
||
| h.single_task<class kernelfunction5>([=] { | ||
| int a = Sampler2.field; | ||
| }); | ||
|
|
||
| }); | ||
|
|
||
| return 0; | ||
| } | ||
|
|
||
| // CHECK: %[[RANGE_TYPE:"struct.*cl::sycl::range"]] | ||
| // CHECK: %[[ID_TYPE:"struct.*cl::sycl::id"]] | ||
| // CHECK: define dso_local spir_kernel void @{{.*}}kernel_function1 | ||
| // CHECK-SAME: i32 [[ARG_A:%[a-zA-Z0-9_]+]], | ||
| // CHECK-SAME: i32 addrspace(1)* [[ACC1_DATA:%[a-zA-Z0-9_]+]], | ||
| // CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC1_DATA:%[a-zA-Z0-9_]+]], | ||
| // CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC2_DATA:%[a-zA-Z0-9_]+]], | ||
| // CHECK-SAME: %[[ID_TYPE]]* byval(%[[ID_TYPE]]) align 4 [[ACC3_DATA:%[a-zA-Z0-9_]+]]) | ||
|
|
||
| // CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.AccessorBase, %class.{{.*}}.AccessorBase addrspace(4)* %3, i32 0, i32 1 | ||
| // CHECK: call spir_func void @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2016ELNS2_11placeholderE0EEC1Ev(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* align 4 dereferenceable_or_null(12) [[ACC_FIELD]]) | ||
|
|
||
| // CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.AccessorBase, %class{{.*}}.AccessorBase addrspace(4)* %5, i32 0, i32 1 | ||
| // CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* %_arg_acc.addr.ascast, align 8 | ||
| // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{.*}} [[ACC1_FIELD]], i32 addrspace(1)* [[ACC1_DATA_LOAD]] | ||
|
There was a problem hiding this comment. This test missed the checks for the code produced by |
||
|
|
||
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -102,7 +102,7 @@ struct DeviceValueType<dataT, access::target::local> { | |
| template <typename dataT, int dimensions, access::mode accessmode, | ||
| access::target accessTarget = access::target::global_buffer, | ||
| access::placeholder isPlaceholder = access::placeholder::false_t> | ||
|
There was a problem hiding this comment. I think you accidentally removed There was a problem hiding this comment. @elizabethandrews That's the next thing I am trying to understand. Do all 3 classes (accessor, sampler and stream) have a property_list? What's the purpose of it? |
||
| class __attribute__((sycl_special_class(accessor))) accessor { | ||
|
|
||
| public: | ||
| void use(void) const {} | ||
|
|
@@ -162,7 +162,7 @@ struct _ImageImplT { | |
| }; | ||
|
|
||
| template <typename dataT, int dimensions, access::mode accessmode> | ||
| class __attribute__((sycl_special_class(accessor))) accessor<dataT, dimensions, accessmode, access::target::image, access::placeholder::false_t> { | ||
| public: | ||
| void use(void) const {} | ||
| template <typename... T> | ||
|
|
@@ -181,7 +181,7 @@ struct sampler_impl { | |
| #endif | ||
| }; | ||
|
|
||
| class __attribute__((sycl_special_class(sampler))) sampler { | ||
| struct sampler_impl impl; | ||
| #ifdef __SYCL_DEVICE_ONLY__ | ||
| void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; } | ||
|
|
@@ -285,7 +285,7 @@ class handler { | |
| } | ||
| }; | ||
|
|
||
| class __attribute__((sycl_special_class(stream))) stream { | ||
| accessor<int, 1, access::mode::read> acc; | ||
|
|
||
| public: | ||
|
|
||
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,87 @@ | ||
| // RUN: %clang_cc1 -S -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -ast-dump -sycl-std=2020 %s | FileCheck %s | ||
|
There was a problem hiding this comment. I'm not sure if following flags are necessary:
|
||
|
|
||
bader marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| #include "sycl.hpp" | ||
|
|
||
| sycl::queue myQueue; | ||
| sycl::handler H; | ||
|
There was a problem hiding this comment. Why these are declared as a global objects instead of |
||
|
|
||
| class AccessorBase { | ||
| int A; | ||
| public: | ||
| sycl::accessor<int, 1, sycl::access::mode::read_write, | ||
| sycl::access::target::local> | ||
| acc; | ||
| }; | ||
|
|
||
| class accessor { | ||
| public: | ||
| int field; | ||
| }; | ||
|
|
||
| class stream { | ||
| public: | ||
| int field; | ||
| }; | ||
|
|
||
| class sampler { | ||
| public: | ||
| int field; | ||
| }; | ||
|
|
||
| int main() { | ||
|
|
||
| AccessorBase Accessor1; | ||
| accessor Accessor2 = {1}; | ||
| sycl::stream Stream1{0, 0, H}; | ||
| stream Stream2; | ||
| sycl::sampler Sampler1; | ||
| sampler Sampler2; | ||
|
|
||
| myQueue.submit([&](sycl::handler &h) { | ||
| h.single_task<class kernel_function1>([=]() { | ||
| Accessor1.acc.use(); | ||
| }); | ||
| h.single_task<class kernel_function2>([=]() { | ||
| int a = Accessor2.field; | ||
| }); | ||
|
|
||
| h.single_task<class kernel_function3>([=]() { | ||
| Stream1.use(); | ||
| }); | ||
| h.single_task<class kernel_function4>([=]() { | ||
| int a = Stream2.field; | ||
| }); | ||
|
|
||
| h.single_task<class kernelfunction5>([=] { | ||
| Sampler1.use(); | ||
| }); | ||
|
|
||
| h.single_task<class kernelfunction6>([=] { | ||
| int a = Sampler2.field; | ||
| }); | ||
|
|
||
| }); | ||
|
|
||
| return 0; | ||
| } | ||
|
|
||
| // CHECK: ClassTemplateDecl {{.*}} accessor | ||
| // CHECK: CXXRecordDecl {{.*}} class accessor definition | ||
| // CHECK: SYCLSpecialClassAttr {{.*}} Accessor | ||
| // CHECK: CXXRecordDecl {{.*}} implicit class accessor | ||
|
|
||
| // CHECK: ClassTemplateSpecializationDecl {{.*}} class accessor definition | ||
| // CHECK: SYCLSpecialClassAttr{{.*}} Accessor | ||
| // CHECK: CXXRecordDecl {{.*}} prev {{.*}} implicit class accessor | ||
|
|
||
| // CHECK: ClassTemplateSpecializationDecl {{.*}} class accessor definition | ||
| // CHECK: SYCLSpecialClassAttr{{.*}} Accessor | ||
| // CHECK: CXXRecordDecl {{.*}} prev {{.*}} implicit class accessor | ||
|
|
||
| // CHECK: CXXRecordDecl {{.*}} referenced class sampler definition | ||
| // CHECK: SYCLSpecialClassAttr {{.*}} Sampler | ||
| // CHECK: CXXRecordDecl {{.*}} implicit class sampler | ||
|
|
||
| // CHECK: CXXRecordDecl {{.*}} prev {{.*}} referenced class stream definition | ||
| // CHECK: SYCLSpecialClassAttr {{.*}} Stream | ||
| // CHECK: CXXRecordDecl {{.*}} implicit referenced class stream | ||
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Same for the one above, this makes the test less fragile.