From 71c71fdb0eff2aecaba355cd6b110a246b18cecf Mon Sep 17 00:00:00 2001 From: Gordon Date: Fri, 14 Sep 2018 14:47:02 +0100 Subject: [PATCH 01/10] CP016: Add initial proposal for sub groups. * Add initial proposal for adding sub groups to SYCL. --- sub-groups/index.md | 26 ++++++++++++ sub-groups/sycl-2.2/index.md | 78 ++++++++++++++++++++++++++++++++++++ 2 files changed, 104 insertions(+) create mode 100644 sub-groups/index.md create mode 100644 sub-groups/sycl-2.2/index.md diff --git a/sub-groups/index.md b/sub-groups/index.md new file mode 100644 index 0000000..317c1cf --- /dev/null +++ b/sub-groups/index.md @@ -0,0 +1,26 @@ +# Sub Groups + +| Proposal ID | CP016 | +|-------------|--------| +| Name | Sub Groups | +| Date of Creation | 14 September 2018 | +| Target | SYCL 2.2 | +| Current Status | _Work In Progress_ | +| Reply-to | Ruyman Reyes | +| Original author | Ruyman Reyes | +| Contributors | Ruyman Reyes , Gordon Brown , Victor Lomuller | + +## Overview + +This proposal aims to define an interface for using OpenCL 2.x sub groups in +SYCL he provisional SYCL 2.2 specification (revision date 2016/02/15) already +contains SVM, but this proposal aims to make SVM in SYCL 2.2 more generic, +easier to program, better defined, and not necessarily tied to OpenCL 2.2. + +## References + +[1] SYCL 1.2.1 specification +https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf + +[2] SYCL 2.2 provisional specification (revision date 2016/02/15) +https://www.khronos.org/registry/SYCL/specs/sycl-2.2.pdf \ No newline at end of file diff --git a/sub-groups/sycl-2.2/index.md b/sub-groups/sycl-2.2/index.md new file mode 100644 index 0000000..ffe60ef --- /dev/null +++ b/sub-groups/sycl-2.2/index.md @@ -0,0 +1,78 @@ +# Sub Groups + +This proposal aims to define an interface for using OpenCL 2.x sub groups in +SYCL he provisional SYCL 2.2 specification (revision date 2016/02/15) already +contains SVM, but this proposal aims to make SVM in SYCL 2.2 more generic, +easier to program, better defined, and not necessarily tied to OpenCL 2.2. + +## sub_group class + +New class template `sub_group` for identifying the sub group range and the current sub group id and also for providing sub group barriers. + +```cpp +template +class sub_group { + public: + + range nd_item::get_sub_group_range() const; + + size_t nd_item::get_sub_group_range(int dimension) const; + + size_t nd_item::get_sub_group_linear_range() const; + + id nd_item::get_sub_group_id() const; + + size_t nd_item::get_sub_group_id(int dimension) const; + + size_t nd_item::get_sub_group_linear_id() const; + + void nd_item::barrier(access::fence_space accessSpace + = access::fence_space::global_and_local) const; + + template + bool all_of(Predicate predicate) const; + + template + bool any_of(Predicate predicate) const; +}; +``` + +## Free functions + +```cpp +template +bool all_of(group group, Predicate predicate); + +template +bool all_of(sub_group subGroup, Predicate predicate); + +template +bool any_of(group group, Predicate predicate); + +template +bool any_of(sub_group subGroup, Predicate predicate); + +template +void barrier(group group, access::fence_space accessSpace + = access::fence_space::global_and_local) const; + +template +void barrier(sub_group subGroup, access::fence_space accessSpace + = access::fence_space::global_and_local) const; +``` + +## Extensions to the nd_item class + +New member function `get_sub_group` for identifying the current sub group and gaining access to sub group operations. + +```cpp +... + +group nd_item::get_sub_group() const; + +... +``` + +## Example + +TODO: Add example \ No newline at end of file From 23dc44af48454c770ca93a4e6214a50b473b6270 Mon Sep 17 00:00:00 2001 From: Ruyman Reyes Date: Fri, 14 Sep 2018 15:24:56 +0100 Subject: [PATCH 02/10] Working on the wording and versions --- sub-groups/index.md | 21 ++++++++----- sub-groups/{sycl-2.2 => sycl-1.2.1}/index.md | 33 +++++++++++++++----- 2 files changed, 39 insertions(+), 15 deletions(-) rename sub-groups/{sycl-2.2 => sycl-1.2.1}/index.md (63%) diff --git a/sub-groups/index.md b/sub-groups/index.md index 317c1cf..e39a1fd 100644 --- a/sub-groups/index.md +++ b/sub-groups/index.md @@ -1,10 +1,10 @@ -# Sub Groups +# Basic sub-group extension | Proposal ID | CP016 | |-------------|--------| -| Name | Sub Groups | +| Name | Basic sub group extension | | Date of Creation | 14 September 2018 | -| Target | SYCL 2.2 | +| Target | SYCL 1.2.1 | | Current Status | _Work In Progress_ | | Reply-to | Ruyman Reyes | | Original author | Ruyman Reyes | @@ -12,10 +12,13 @@ ## Overview -This proposal aims to define an interface for using OpenCL 2.x sub groups in -SYCL he provisional SYCL 2.2 specification (revision date 2016/02/15) already -contains SVM, but this proposal aims to make SVM in SYCL 2.2 more generic, -easier to program, better defined, and not necessarily tied to OpenCL 2.2. +This vendor extension aims to define an interface to expose sub-group functionality, +as defined in the SYCL 2.2 provisional and the OpenCL 2.2 provisional, +in SYCL 1.2.1. + +The extension is only targeting OpenCL devices that expose +`cl_codeplay_basic_sub_group` vendor extension. + ## References @@ -23,4 +26,6 @@ easier to program, better defined, and not necessarily tied to OpenCL 2.2. https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf [2] SYCL 2.2 provisional specification (revision date 2016/02/15) -https://www.khronos.org/registry/SYCL/specs/sycl-2.2.pdf \ No newline at end of file +https://www.khronos.org/registry/SYCL/specs/sycl-2.2.pdf + +[3] XXX diff --git a/sub-groups/sycl-2.2/index.md b/sub-groups/sycl-1.2.1/index.md similarity index 63% rename from sub-groups/sycl-2.2/index.md rename to sub-groups/sycl-1.2.1/index.md index ffe60ef..6f79b6a 100644 --- a/sub-groups/sycl-2.2/index.md +++ b/sub-groups/sycl-1.2.1/index.md @@ -1,13 +1,24 @@ -# Sub Groups +# Basic Sub group support This proposal aims to define an interface for using OpenCL 2.x sub groups in -SYCL he provisional SYCL 2.2 specification (revision date 2016/02/15) already -contains SVM, but this proposal aims to make SVM in SYCL 2.2 more generic, -easier to program, better defined, and not necessarily tied to OpenCL 2.2. +SYCL the provisional SYCL 1.2.1 specification, relying on the underlying +OpenCL implementation supporting the extension `cl_codeplay_basic_sub_groups`. -## sub_group class +The extension exposes to programmers the ability to identify sub-groups +on a work-group, count the number of sub-groups available. -New class template `sub_group` for identifying the sub group range and the current sub group id and also for providing sub group barriers. +## Namespace `basic_sub_group` + +All new functionality is exposed under the `basic_sub_group` namespace +in the codeplay vendor namespace. +When the vendor extension `basic_sub_group` is available, the macro +`SYCL_CODEPLAY_BASIC_SUB_GROUP` is defined in the header. + +### New sub\_group class + +The extension adds a new class template `sub_group` that identifies the +sub group range and the current sub group id. +It also for providing sub group barriers. ```cpp template @@ -29,6 +40,11 @@ class sub_group { void nd_item::barrier(access::fence_space accessSpace = access::fence_space::global_and_local) const; + /* T is permitted to be int, unsigned int, long, unsigned long, + float, half, double */ + template + T broadcast(size_t subGroupId, T value); + template bool all_of(Predicate predicate) const; @@ -40,6 +56,9 @@ class sub_group { ## Free functions ```cpp +template +T broadcast(sub_group subGroup, size_t subGroupId, T value); + template bool all_of(group group, Predicate predicate); @@ -75,4 +94,4 @@ group nd_item::get_sub_group() const; ## Example -TODO: Add example \ No newline at end of file +TODO: Add example From f08fbdcf2be4a4ee60db7d77dabdfcf626059710 Mon Sep 17 00:00:00 2001 From: Gordon Date: Fri, 14 Sep 2018 15:34:23 +0100 Subject: [PATCH 03/10] CP016: Add some minor changes to sub groups proposal. * Add links to OpenCL 2.2 and OpenCL C++ 1.0 specifications. * Add specification of the codeplay::nd_item class template. * Add proposal to readme. * Add minor formatting changes. --- README.md | 1 + sub-groups/index.md | 7 ++++++- sub-groups/sycl-1.2.1/index.md | 8 +++++--- 3 files changed, 12 insertions(+), 4 deletions(-) diff --git a/README.md b/README.md index 7338c0e..6e57db0 100644 --- a/README.md +++ b/README.md @@ -54,3 +54,4 @@ Each proposal in the table below will be tagged with one of the following states | CP013 | [Supporting Heterogeneous & Distributed Computing Through Affinity](affinity/index.md) | ISO C++ SG1, SG14 | 15 November 2017 | 12 August 2018 | _Work in Progress_ | | CP014 | [Shared Virtual Memory](svm/index.md) | SYCL 2.2 | 22 January 2018 | 22 January 2018 | _Work in Progress_ | | CP015 | [Specialization Constant](spec-constant/index.md) | SYCL 1.2.1 extension / SYCL 2.2 | 24 April 2018 | 24 April 2018 | _Work in Progress_ | +| CP016 | [Sub Groups](spec-constant/index.md) | SYCL 1.2.1 extension | 14 September 2018 | 14 September 2018 | _Work in Progress_ | diff --git a/sub-groups/index.md b/sub-groups/index.md index e39a1fd..4bee132 100644 --- a/sub-groups/index.md +++ b/sub-groups/index.md @@ -28,4 +28,9 @@ https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf [2] SYCL 2.2 provisional specification (revision date 2016/02/15) https://www.khronos.org/registry/SYCL/specs/sycl-2.2.pdf -[3] XXX +[3] OpenCL 2.2 API specification +https://www.khronos.org/registry/OpenCL/specs/2.2/pdf/OpenCL_API.pdf + +[4] OpenCL C++ 1.0 specification +https://www.khronos.org/registry/OpenCL/specs/2.2/pdf/OpenCL_Cxx.pdf + diff --git a/sub-groups/sycl-1.2.1/index.md b/sub-groups/sycl-1.2.1/index.md index 6f79b6a..1ad0cce 100644 --- a/sub-groups/sycl-1.2.1/index.md +++ b/sub-groups/sycl-1.2.1/index.md @@ -10,11 +10,11 @@ on a work-group, count the number of sub-groups available. ## Namespace `basic_sub_group` All new functionality is exposed under the `basic_sub_group` namespace -in the codeplay vendor namespace. +in the `codeplay` vendor extension namespace. When the vendor extension `basic_sub_group` is available, the macro `SYCL_CODEPLAY_BASIC_SUB_GROUP` is defined in the header. -### New sub\_group class +### Class `sub_group` The extension adds a new class template `sub_group` that identifies the sub group range and the current sub group id. @@ -41,7 +41,7 @@ class sub_group { = access::fence_space::global_and_local) const; /* T is permitted to be int, unsigned int, long, unsigned long, - float, half, double */ + float, half, double */ template T broadcast(size_t subGroupId, T value); @@ -82,6 +82,8 @@ void barrier(sub_group subGroup, access::fence_space accessSpace ## Extensions to the nd_item class +Extensions to the `nd_item` interface will be exposed via the a derived `nd_item` class template in the `codeplay` vendor extension namespace. + New member function `get_sub_group` for identifying the current sub group and gaining access to sub group operations. ```cpp From 4497bf85bda29fa94dbad7f5e29280a5219d9632 Mon Sep 17 00:00:00 2001 From: Gordon Date: Fri, 14 Sep 2018 15:42:12 +0100 Subject: [PATCH 04/10] CP016: Add explicit definition of codeplay::nd_item. * Move definition of codeplay extension nd_item to a class definition. * Add namespaces for class definitions. --- sub-groups/sycl-1.2.1/index.md | 44 ++++++++++++++++++++++++---------- 1 file changed, 31 insertions(+), 13 deletions(-) diff --git a/sub-groups/sycl-1.2.1/index.md b/sub-groups/sycl-1.2.1/index.md index 1ad0cce..e10e56b 100644 --- a/sub-groups/sycl-1.2.1/index.md +++ b/sub-groups/sycl-1.2.1/index.md @@ -21,6 +21,10 @@ sub group range and the current sub group id. It also for providing sub group barriers. ```cpp +namespace cl { +namespace sycl { +namespace codeplay { + template class sub_group { public: @@ -45,39 +49,43 @@ class sub_group { template T broadcast(size_t subGroupId, T value); + /* Predicate must be a callable type which returns bool */ template bool all_of(Predicate predicate) const; + /* Predicate must be a callable type which returns bool */ template bool any_of(Predicate predicate) const; }; + +} // namespace codeplay +} // namespace sycl +} // namespace cl ``` ## Free functions ```cpp +namespace cl { +namespace sycl { +namespace codeplay { + template T broadcast(sub_group subGroup, size_t subGroupId, T value); -template -bool all_of(group group, Predicate predicate); - template bool all_of(sub_group subGroup, Predicate predicate); -template -bool any_of(group group, Predicate predicate); - template bool any_of(sub_group subGroup, Predicate predicate); -template -void barrier(group group, access::fence_space accessSpace - = access::fence_space::global_and_local) const; - template void barrier(sub_group subGroup, access::fence_space accessSpace = access::fence_space::global_and_local) const; + +} // namespace codeplay +} // namespace sycl +} // namespace cl ``` ## Extensions to the nd_item class @@ -87,11 +95,21 @@ Extensions to the `nd_item` interface will be exposed via the a derived `nd_item New member function `get_sub_group` for identifying the current sub group and gaining access to sub group operations. ```cpp -... +namespace cl { +namespace sycl { +namespace codeplay { -group nd_item::get_sub_group() const; +template +class nd_item : public ::cl::sycl::nd_item { +public: + + group nd_item::get_sub_group() const; + +}; -... +} // namespace codeplay +} // namespace sycl +} // namespace cl ``` ## Example From 56d5f4c287cc9a4af776e45831f2f113a5cc4e9f Mon Sep 17 00:00:00 2001 From: Ruyman Reyes Date: Fri, 14 Sep 2018 15:43:07 +0100 Subject: [PATCH 05/10] Added wording for memory model --- sub-groups/sycl-1.2.1/index.md | 42 +++++++++++++++++++++++++++++++--- 1 file changed, 39 insertions(+), 3 deletions(-) diff --git a/sub-groups/sycl-1.2.1/index.md b/sub-groups/sycl-1.2.1/index.md index 6f79b6a..94a047c 100644 --- a/sub-groups/sycl-1.2.1/index.md +++ b/sub-groups/sycl-1.2.1/index.md @@ -1,12 +1,37 @@ # Basic Sub group support -This proposal aims to define an interface for using OpenCL 2.x sub groups in +This proposal aims to define an interface for using OpenCL 2.2 sub groups in SYCL the provisional SYCL 1.2.1 specification, relying on the underlying OpenCL implementation supporting the extension `cl_codeplay_basic_sub_groups`. The extension exposes to programmers the ability to identify sub-groups on a work-group, count the number of sub-groups available. +Details of the execution and memory model changes can be found in the +documentation for the Codeplay's OpenCL vendor extension `cl_codeplay_basic_sub_groups` +once available. + +## Execution model + +When this vendor extension is available, the execution model of SYCL 1.2.1 +is extended to also include support for sub-groups of threads inside of a +work-group. +Overall, these sub-groups work following the description of the OpenCL 2.2 +sub-groups, with some restrictions: + +* The number of sub-groups available for each work-group is determined +at compile-time and remains the same during the execution of the SYCL application. +* The number of threads per sub-group is known at compile-time, and remains the +same during execution of the SYCL application. +* Only those functions defined in this proposal are available. +In particular, there is no sub-group pipe communication. + +## Memory model + +Sub-groups can access global and local memory, but, given there is no +memory-scope to the atomic or barriers operations in SYCL 1.2.1, there is no +possibility to specify an equivalent of sub-group memory scope. + ## Namespace `basic_sub_group` All new functionality is exposed under the `basic_sub_group` namespace @@ -80,7 +105,7 @@ void barrier(sub_group subGroup, access::fence_space accessSpace = access::fence_space::global_and_local) const; ``` -## Extensions to the nd_item class +## Extensions to the nd\_item class New member function `get_sub_group` for identifying the current sub group and gaining access to sub group operations. @@ -94,4 +119,15 @@ group nd_item::get_sub_group() const; ## Example -TODO: Add example +```cpp +template +void my_subgroup_load(sub_group subG, global_ptr myArray) { + + float4 f; + if (subG.get_id() == 0) { + f.load(myArray); + } + barrier(subG, access::fence_space::global_and_local); + float4 res = broadcast(subG, 0, f); +} +``` From b9afd9ae045b64b2c1dd7b708d3ee7ab875f72ad Mon Sep 17 00:00:00 2001 From: Ruyman Reyes Date: Fri, 14 Sep 2018 15:45:07 +0100 Subject: [PATCH 06/10] Completed sentence --- sub-groups/sycl-1.2.1/index.md | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sub-groups/sycl-1.2.1/index.md b/sub-groups/sycl-1.2.1/index.md index 8836998..1c81dae 100644 --- a/sub-groups/sycl-1.2.1/index.md +++ b/sub-groups/sycl-1.2.1/index.md @@ -5,7 +5,8 @@ SYCL the provisional SYCL 1.2.1 specification, relying on the underlying OpenCL implementation supporting the extension `cl_codeplay_basic_sub_groups`. The extension exposes to programmers the ability to identify sub-groups -on a work-group, count the number of sub-groups available. +on a work-group, count the number of sub-groups available and perform +a broadcast from one work-item on a sub-group to the rest. Details of the execution and memory model changes can be found in the documentation for the Codeplay's OpenCL vendor extension `cl_codeplay_basic_sub_groups` From 88d23ea13aa247d4136977f16fdaa101a68bd099 Mon Sep 17 00:00:00 2001 From: Ruyman Reyes Date: Fri, 14 Sep 2018 15:50:19 +0100 Subject: [PATCH 07/10] Fixed interface --- sub-groups/sycl-1.2.1/index.md | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sub-groups/sycl-1.2.1/index.md b/sub-groups/sycl-1.2.1/index.md index 77f4d75..b683a22 100644 --- a/sub-groups/sycl-1.2.1/index.md +++ b/sub-groups/sycl-1.2.1/index.md @@ -55,11 +55,11 @@ template class sub_group { public: - range nd_item::get_sub_group_range() const; + constexpr range nd_item::get_sub_group_range() const; - size_t nd_item::get_sub_group_range(int dimension) const; + constexpr size_t nd_item::get_sub_group_range(int dimension) const; - size_t nd_item::get_sub_group_linear_range() const; + constexpr size_t nd_item::get_sub_group_linear_range() const; id nd_item::get_sub_group_id() const; @@ -129,7 +129,7 @@ template class nd_item : public ::cl::sycl::nd_item { public: - group nd_item::get_sub_group() const; + sub_group nd_item::get_sub_group() const; }; From 411d1a0b5f88425a67f4ccc4ac702c3c28a1056b Mon Sep 17 00:00:00 2001 From: Gordon Date: Fri, 14 Sep 2018 15:51:36 +0100 Subject: [PATCH 08/10] CP016: Add description of example for sub group proposal. * Add description of sub groups example. --- sub-groups/sycl-1.2.1/index.md | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/sub-groups/sycl-1.2.1/index.md b/sub-groups/sycl-1.2.1/index.md index 0799fb6..96271ae 100644 --- a/sub-groups/sycl-1.2.1/index.md +++ b/sub-groups/sycl-1.2.1/index.md @@ -139,12 +139,14 @@ public: ## Example +Below is trivial example showing how you would use `sub_group` to broadcast a value from one work-item within a sub-group to all other work-items in the sub-group. + ```cpp template void my_subgroup_load(sub_group subG, global_ptr myArray) { float4 f; - if (subG.get_id() == 0) { + if (subG.get_id() == 0) { f.load(myArray); } barrier(subG, access::fence_space::global_and_local); From 84390e65479ec221fda36332c6af93523cf8e909 Mon Sep 17 00:00:00 2001 From: Ruyman Reyes Date: Fri, 14 Sep 2018 15:58:58 +0100 Subject: [PATCH 09/10] Fixed the name of the OpenCL vendor extension --- sub-groups/index.md | 2 +- sub-groups/sycl-1.2.1/index.md | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/sub-groups/index.md b/sub-groups/index.md index 4bee132..22da9da 100644 --- a/sub-groups/index.md +++ b/sub-groups/index.md @@ -17,7 +17,7 @@ as defined in the SYCL 2.2 provisional and the OpenCL 2.2 provisional, in SYCL 1.2.1. The extension is only targeting OpenCL devices that expose -`cl_codeplay_basic_sub_group` vendor extension. +`cl_codeplay_basic_subgroups` vendor extension. ## References diff --git a/sub-groups/sycl-1.2.1/index.md b/sub-groups/sycl-1.2.1/index.md index b683a22..04726c6 100644 --- a/sub-groups/sycl-1.2.1/index.md +++ b/sub-groups/sycl-1.2.1/index.md @@ -2,14 +2,14 @@ This proposal aims to define an interface for using OpenCL 2.2 sub groups in SYCL the provisional SYCL 1.2.1 specification, relying on the underlying -OpenCL implementation supporting the extension `cl_codeplay_basic_sub_groups`. +OpenCL implementation supporting the extension `cl_codeplay_basic_subgroups`. The extension exposes to programmers the ability to identify sub-groups on a work-group, count the number of sub-groups available and perform a broadcast from one work-item on a sub-group to the rest. Details of the execution and memory model changes can be found in the -documentation for the Codeplay's OpenCL vendor extension `cl_codeplay_basic_sub_groups` +documentation for the Codeplay's OpenCL vendor extension `cl_codeplay_basic_subgroups` once available. ## Execution model From ea3889797c6ee81ef13501c06bebb295b20924e3 Mon Sep 17 00:00:00 2001 From: Ruyman Date: Sat, 15 Sep 2018 21:17:02 +0100 Subject: [PATCH 10/10] Update index.md * Removed unnecessary namespace --- sub-groups/sycl-1.2.1/index.md | 19 ++++++++++--------- 1 file changed, 10 insertions(+), 9 deletions(-) diff --git a/sub-groups/sycl-1.2.1/index.md b/sub-groups/sycl-1.2.1/index.md index 3bcc570..5c89922 100644 --- a/sub-groups/sycl-1.2.1/index.md +++ b/sub-groups/sycl-1.2.1/index.md @@ -55,20 +55,19 @@ template class sub_group { public: - constexpr range nd_item::get_sub_group_range() const; + constexpr range get_sub_group_range() const; - constexpr size_t nd_item::get_sub_group_range(int dimension) const; + constexpr size_t get_sub_group_range(int dimension) const; - constexpr size_t nd_item::get_sub_group_linear_range() const; + constexpr size_t get_sub_group_linear_range() const; - id nd_item::get_sub_group_id() const; + id get_sub_group_id() const; - size_t nd_item::get_sub_group_id(int dimension) const; + size_t get_sub_group_id(int dimension) const; - size_t nd_item::get_sub_group_linear_id() const; + size_t get_sub_group_linear_id() const; - void nd_item::barrier(access::fence_space accessSpace - = access::fence_space::global_and_local) const; + void barrier(access::fence_space accessSpace = access::fence_space::global_and_local) const; /* T is permitted to be int, unsigned int, long, unsigned long, float, half, double */ @@ -129,7 +128,7 @@ template class nd_item : public ::cl::sycl::nd_item { public: - sub_group nd_item::get_sub_group() const; + sub_group get_sub_group() const; }; @@ -143,6 +142,8 @@ public: Below is trivial example showing how you would use `sub_group` to broadcast a value from one work-item within a sub-group to all other work-items in the sub-group. ```cpp +using namespace cl::sycl; + template void my_subgroup_load(sub_group subG, global_ptr myArray) {