forked from intel/llvm
-
Notifications
You must be signed in to change notification settings - Fork 4
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[SYCL][Docs] Allow copy-construction of device_global (intel#15075)
This commit makes it possible to copy-construct device_global variables if they do not have the device_image_scope property. The restriction on device_image_scope is due to static construction not being allowed in device code, which they would require, while other device_globals have USM storage which will be initialized by the host code, so the constructor on the device is a simple zero-initialization. --------- Signed-off-by: Larsen, Steffen <[email protected]> Co-authored-by: John Pennycook <[email protected]>
- Loading branch information
1 parent
5d5ec9e
commit e6e45d0
Showing
4 changed files
with
210 additions
and
6 deletions.
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,68 @@ | ||
| // RUN: %{build} -std=c++23 -o %t.out | ||
| // RUN: %{run} %t.out | ||
| // | ||
| // UNSUPPORTED: opencl && gpu | ||
| // UNSUPPORTED-TRACKER: GSD-4287 | ||
| // | ||
| // Tests the copy ctor on device_global without device_image_scope. | ||
|
|
||
| #include <sycl/detail/core.hpp> | ||
|
|
||
| namespace oneapiext = sycl::ext::oneapi::experimental; | ||
|
|
||
| oneapiext::device_global<const int> DGInit1{3}; | ||
| oneapiext::device_global<const int> DGCopy1{DGInit1}; | ||
|
|
||
| oneapiext::device_global<int> DGInit2{4}; | ||
| oneapiext::device_global<int> DGCopy2{DGInit2}; | ||
|
|
||
| oneapiext::device_global<float> DGInit3{5.0f}; | ||
| oneapiext::device_global<int> DGCopy3{DGInit3}; | ||
|
|
||
| oneapiext::device_global<const int, decltype(oneapiext::properties{ | ||
| oneapiext::device_image_scope})> | ||
| DGInit4{6}; | ||
| oneapiext::device_global<const int> DGCopy4{DGInit4}; | ||
|
|
||
| oneapiext::device_global<const int> DGInit5{7}; | ||
| oneapiext::device_global<const int, decltype(oneapiext::properties{ | ||
| oneapiext::host_access_read})> | ||
| DGCopy5{DGInit5}; | ||
|
|
||
| int main() { | ||
| sycl::queue Q; | ||
|
|
||
| int ReadVals[10] = {0, 0}; | ||
| { | ||
| sycl::buffer<int, 1> ReadValsBuff{ReadVals, 10}; | ||
|
|
||
| Q.submit([&](sycl::handler &CGH) { | ||
| sycl::accessor ReadValsAcc{ReadValsBuff, CGH, sycl::write_only}; | ||
| CGH.single_task([=]() { | ||
| ReadValsAcc[0] = DGInit1.get(); | ||
| ReadValsAcc[1] = DGCopy1.get(); | ||
| ReadValsAcc[2] = DGInit2.get(); | ||
| ReadValsAcc[3] = DGCopy2.get(); | ||
| ReadValsAcc[4] = DGInit3.get(); | ||
| ReadValsAcc[5] = DGCopy3.get(); | ||
| ReadValsAcc[6] = DGInit4.get(); | ||
| ReadValsAcc[7] = DGCopy4.get(); | ||
| ReadValsAcc[8] = DGInit5.get(); | ||
| ReadValsAcc[9] = DGCopy5.get(); | ||
| }); | ||
| }).wait_and_throw(); | ||
| } | ||
|
|
||
| assert(ReadVals[0] == 3); | ||
| assert(ReadVals[1] == 3); | ||
| assert(ReadVals[2] == 4); | ||
| assert(ReadVals[3] == 4); | ||
| assert(ReadVals[4] == 5); | ||
| assert(ReadVals[5] == 5); | ||
| assert(ReadVals[6] == 6); | ||
| assert(ReadVals[7] == 6); | ||
| assert(ReadVals[8] == 7); | ||
| assert(ReadVals[9] == 7); | ||
|
|
||
| return 0; | ||
| } |
29 changes: 29 additions & 0 deletions
29
sycl/test/extensions/device_global/device_global_copy_negative.cpp
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,29 @@ | ||
| // RUN: %clangxx -std=c++23 -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s | ||
| // | ||
| // Tests that the copy ctor on device_global with device_image_scope is | ||
| // unavailable. | ||
|
|
||
| #include <sycl/sycl.hpp> | ||
|
|
||
| namespace oneapiext = sycl::ext::oneapi::experimental; | ||
|
|
||
| using device_image_properties = | ||
| decltype(oneapiext::properties{oneapiext::device_image_scope}); | ||
|
|
||
| // expected-error@sycl/ext/oneapi/device_global/device_global.hpp:* {{call to deleted constructor}} | ||
| oneapiext::device_global<const int, device_image_properties> DGInit1{3}; | ||
| oneapiext::device_global<const int, device_image_properties> DGCopy1{DGInit1}; | ||
|
|
||
| // expected-error@sycl/ext/oneapi/device_global/device_global.hpp:* {{call to deleted constructor}} | ||
| oneapiext::device_global<int, device_image_properties> DGInit2{3}; | ||
| oneapiext::device_global<int, device_image_properties> DGCopy2{DGInit2}; | ||
|
|
||
| // expected-error@+2 {{call to deleted constructor}} | ||
| oneapiext::device_global<int, device_image_properties> DGInit3{3}; | ||
| oneapiext::device_global<float, device_image_properties> DGCopy3{DGInit3}; | ||
|
|
||
| // expected-error@+2 {{call to deleted constructor}} | ||
| oneapiext::device_global<const int> DGInit4{3}; | ||
| oneapiext::device_global<const int, device_image_properties> DGCopy4{DGInit4}; | ||
|
|
||
| int main() { return 0; } |