diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/Reductions.ipynb b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/Reductions.ipynb index 792ac181d6..f1060f99be 100644 --- a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/Reductions.ipynb +++ b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/Reductions.ipynb @@ -433,10 +433,10 @@ "\n", "It is common for parallel kernels to produce a single output resulting from some combination of all inputs (e.g. the sum). Writing efficient reductions is a complex task, depending on both device and runtime characteristics. Providing an abstraction for reductions in SYCL would greatly improve programmer productivity.\n", "\n", - "`sycl::ext::oneapi::reduction` object in parallel_for encapsulates the reduction variable, an optional operator identity and the reduction operator as shown below:\n", + "`sycl::reduction` object in parallel_for encapsulates the reduction variable, an optional operator identity and the reduction operator as shown below:\n", "\n", "```cpp\n", - " q.parallel_for(nd_range<1>{N, B}, sycl::ext::oneapi::reduction(sum, 0, sycl::plus<>()), [=](nd_item<1> it, auto& temp) {\n", + " q.parallel_for(nd_range<1>{N, B}, sycl::reduction(sum, sycl::plus<>()), [=](nd_item<1> it, auto& temp) {\n", " int i = it.get_global_id(0);\n", " temp.combine(data[i]);\n", " });\n", @@ -454,7 +454,7 @@ "cell_type": "markdown", "metadata": {}, "source": [ - "The code below uses __sycl::ext::oneapi::reduction__ object in _parallel_for_ to compute the reduction with just one kernel using Unified Shared Memory(USM) for memory management.\n", + "The code below uses __sycl::reduction__ object in _parallel_for_ to compute the reduction with just one kernel using Unified Shared Memory(USM) for memory management.\n", "\n", "The DPC++ code below demonstrates reduction in parallel_for with USM: Inspect code, there are no modifications necessary:\n", "\n", @@ -469,7 +469,7 @@ "metadata": {}, "outputs": [], "source": [ - "%%writefile lab/sum_oneapi_reduction_usm.cpp\n", + "%%writefile lab/sum_reduction_usm.cpp\n", "//==============================================================\n", "// Copyright © 2020 Intel Corporation\n", "//\n", @@ -496,7 +496,7 @@ " *sum = 0;\n", "\n", " //# nd-range kernel parallel_for with reduction parameter\n", - " q.parallel_for(nd_range<1>{N, B}, ext::oneapi::reduction(sum, 0, plus<>()), [=](nd_item<1> it, auto& temp) {\n", + " q.parallel_for(nd_range<1>{N, B}, reduction(sum, plus<>()), [=](nd_item<1> it, auto& temp) {\n", " auto i = it.get_global_id(0);\n", " temp.combine(data[i]);\n", " }).wait();\n", @@ -523,7 +523,7 @@ "metadata": {}, "outputs": [], "source": [ - "! chmod 755 q; chmod 755 run_sum_oneapi_reduction_usm.sh; if [ -x \"$(command -v qsub)\" ]; then ./q run_sum_oneapi_reduction_usm.sh; else ./run_sum_oneapi_reduction_usm.sh; fi" + "! chmod 755 q; chmod 755 run_sum_reduction_usm.sh; if [ -x \"$(command -v qsub)\" ]; then ./q run_sum_reduction_usm.sh; else ./run_sum_reduction_usm.sh; fi" ] }, { @@ -545,7 +545,7 @@ "cell_type": "markdown", "metadata": {}, "source": [ - "The code below uses __sycl::ext::oneapi::reduction__ object in _parallel_for_ to compute the reduction with just one kernel using SYCL buffers and accessors for memory management.\n", + "The code below uses __sycl::reduction__ object in _parallel_for_ to compute the reduction with just one kernel using SYCL buffers and accessors for memory management.\n", "\n", "The DPC++ code below demonstrates reduction in parallel_for with Buffers: Inspect code, there are no modifications necessary:\n", "\n", @@ -560,7 +560,7 @@ "metadata": {}, "outputs": [], "source": [ - "%%writefile lab/sum_oneapi_reduction_buffers.cpp\n", + "%%writefile lab/sum_reduction_buffers.cpp\n", "//==============================================================\n", "// Copyright © 2020 Intel Corporation\n", "//\n", @@ -586,12 +586,11 @@ " buffer buf_sum(&sum, range(1));\n", "\n", " q.submit([&](handler& h) {\n", - " //# create accessors for buffers\n", + " //# create accessors for buffer\n", " accessor acc_data(buf_data, h, read_only);\n", - " accessor acc_sum(buf_sum, h);\n", "\n", " //# nd-range kernel parallel_for with reduction parameter\n", - " h.parallel_for(nd_range<1>{N, B}, ext::oneapi::reduction(acc_sum, 0, plus<>()), [=](nd_item<1> it, auto& temp) {\n", + " h.parallel_for(nd_range<1>{N, B}, reduction(buf_sum, h, plus<>()), [=](nd_item<1> it, auto& temp) {\n", " auto i = it.get_global_id(0);\n", " temp.combine(acc_data[i]);\n", " });\n", @@ -617,7 +616,7 @@ "metadata": {}, "outputs": [], "source": [ - "! chmod 755 q; chmod 755 run_sum_oneapi_reduction_buffers.sh; if [ -x \"$(command -v qsub)\" ]; then ./q run_sum_oneapi_reduction_buffers.sh; else ./run_sum_oneapi_reduction_buffers.sh; fi" + "! chmod 755 q; chmod 755 run_sum_reduction_buffers.sh; if [ -x \"$(command -v qsub)\" ]; then ./q run_sum_reduction_buffers.sh; else ./run_sum_reduction_buffers.sh; fi" ] }, { @@ -639,7 +638,7 @@ "cell_type": "markdown", "metadata": {}, "source": [ - "The code below uses multiple __sycl::ext::oneapi::reduction__ objects in _parallel_for_ to compute the reductions with just one kernel using SYCL buffers and accessors for memory management.\n", + "The code below uses multiple __sycl::reduction__ objects in _parallel_for_ to compute the reductions with just one kernel using SYCL buffers and accessors for memory management.\n", "\n", "Multiple reductions are also supported with just one kernel, the code snippet below shows how to definne a kernel using parallel_for with multiple reduction objects:\n", "\n", @@ -693,14 +692,11 @@ " q.submit([&](handler& h) {\n", " //# create accessors for data and results\n", " accessor acc_data(buf_data, h, read_only);\n", - " accessor acc_sum(buf_sum, h);\n", - " accessor acc_min(buf_min, h);\n", - " accessor acc_max(buf_max, h);\n", " \n", " //# define reduction objects for sum, min, max reduction\n", - " auto reduction_sum = ext::oneapi::reduction(acc_sum, 0, plus<>());\n", - " auto reduction_min = ext::oneapi::reduction(acc_min, 0, minimum<>());\n", - " auto reduction_max = ext::oneapi::reduction(acc_max, 0, maximum<>());\n", + " auto reduction_sum = reduction(buf_sum, h, plus<>());\n", + " auto reduction_min = reduction(buf_min, h, minimum<>());\n", + " auto reduction_max = reduction(buf_max, h, maximum<>());\n", " \n", " //# parallel_for with multiple reduction objects\n", " h.parallel_for(nd_range<1>{N, B}, reduction_sum, reduction_min, reduction_max, [=](nd_item<1> it, auto& temp_sum, auto& temp_min, auto& temp_max) {\n", @@ -760,7 +756,7 @@ "cell_type": "markdown", "metadata": {}, "source": [ - "The code below uses __sycl::ext::oneapi::reduction__ object in _parallel_for_ to compute the reduction object that uses a custom operator to find minumum value and index.\n", + "The code below uses __sycl::reduction__ object in _parallel_for_ to compute the reduction object that uses a custom operator to find minumum value and index.\n", "\n", "The DPC++ code below demonstrates reduction in parallel_for with custom user defined operator to perform reduction: Inspect code, there are no modifications necessary:\n", "\n", @@ -775,7 +771,7 @@ "metadata": {}, "outputs": [], "source": [ - "%%writefile lab/oneapi_reduction_custom_operator.cpp\n", + "%%writefile lab/reduction_custom_operator.cpp\n", "//==============================================================\n", "// Copyright © 2020 Intel Corporation\n", "//\n", @@ -815,7 +811,7 @@ " //# custom operator for reduction to find minumum and index\n", " pair operator_identity = {std::numeric_limits::max(), std::numeric_limits::min()};\n", " *result = operator_identity;\n", - " auto reduction_object = ext::oneapi::reduction(result, operator_identity, minimum>());\n", + " auto reduction_object = reduction(result, operator_identity, minimum>());\n", "\n", " //# parallel_for with user defined reduction object\n", " q.parallel_for(nd_range<1>{N, B}, reduction_object, [=](nd_item<1> item, auto& temp) {\n", @@ -845,7 +841,7 @@ "metadata": {}, "outputs": [], "source": [ - "! chmod 755 q; chmod 755 run_oneapi_reduction_custom_operator.sh; if [ -x \"$(command -v qsub)\" ]; then ./q run_oneapi_reduction_custom_operator.sh; else ./run_oneapi_reduction_custom_operator.sh; fi" + "! chmod 755 q; chmod 755 run_reduction_custom_operator.sh; if [ -x \"$(command -v qsub)\" ]; then ./q run_reduction_custom_operator.sh; else ./run_reduction_custom_operator.sh; fi" ] }, { @@ -867,15 +863,8 @@ "cell_type": "markdown", "metadata": {}, "source": [ - "`sycl::reduce_over_group` function for sub_group/work_group and `sycl::ext::oneapi::reduction` in parallel_for helps to optimize and simplify reduction computation in DPC++" + "`sycl::reduce_over_group` function for sub_group/work_group and `sycl::reduction` in parallel_for helps to optimize and simplify reduction computation in DPC++" ] - }, - { - "cell_type": "code", - "execution_count": null, - "metadata": {}, - "outputs": [], - "source": [] } ], "metadata": { diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/assets/hwmapping.png b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/assets/hwmapping.png index dd6d6a50f0..11747050c7 100644 Binary files a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/assets/hwmapping.png and b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/assets/hwmapping.png differ diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/multiple_reductions_buffers.cpp b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/multiple_reductions_buffers.cpp index 3f7f939c54..b2ee47eb6d 100644 --- a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/multiple_reductions_buffers.cpp +++ b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/multiple_reductions_buffers.cpp @@ -28,14 +28,11 @@ int main() { q.submit([&](handler& h) { //# create accessors for data and results accessor acc_data(buf_data, h, read_only); - accessor acc_sum(buf_sum, h); - accessor acc_min(buf_min, h); - accessor acc_max(buf_max, h); //# define reduction objects for sum, min, max reduction - auto reduction_sum = ext::oneapi::reduction(acc_sum, 0, plus<>()); - auto reduction_min = ext::oneapi::reduction(acc_min, 0, minimum<>()); - auto reduction_max = ext::oneapi::reduction(acc_max, 0, maximum<>()); + auto reduction_sum = reduction(buf_sum, h, plus<>()); + auto reduction_min = reduction(buf_min, h, minimum<>()); + auto reduction_max = reduction(buf_max, h, maximum<>()); //# parallel_for with multiple reduction objects h.parallel_for(nd_range<1>{N, B}, reduction_sum, reduction_min, reduction_max, [=](nd_item<1> it, auto& temp_sum, auto& temp_min, auto& temp_max) { diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/oneapi_reduction_custom_operator.cpp b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/reduction_custom_operator.cpp similarity index 93% rename from DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/oneapi_reduction_custom_operator.cpp rename to DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/reduction_custom_operator.cpp index 23e7dfffa9..4f4d76e7e3 100644 --- a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/oneapi_reduction_custom_operator.cpp +++ b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/reduction_custom_operator.cpp @@ -37,7 +37,7 @@ int main() { //# custom operator for reduction to find minumum and index pair operator_identity = {std::numeric_limits::max(), std::numeric_limits::min()}; *result = operator_identity; - auto reduction_object = ext::oneapi::reduction(result, operator_identity, minimum>()); + auto reduction_object = reduction(result, operator_identity, minimum>()); //# parallel_for with user defined reduction object q.parallel_for(nd_range<1>{N, B}, reduction_object, [=](nd_item<1> item, auto& temp) { diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/sum_oneapi_reduction_buffers.cpp b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/sum_reduction_buffers.cpp similarity index 100% rename from DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/sum_oneapi_reduction_buffers.cpp rename to DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/sum_reduction_buffers.cpp diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/sum_oneapi_reduction_usm.cpp b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/sum_reduction_usm.cpp similarity index 100% rename from DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/sum_oneapi_reduction_usm.cpp rename to DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/lab/sum_reduction_usm.cpp diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_oneapi_reduction_custom_operator.sh b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_reduction_custom_operator.sh similarity index 100% rename from DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_oneapi_reduction_custom_operator.sh rename to DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_reduction_custom_operator.sh diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_sum_oneapi_reduction_buffers.sh b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_sum_reduction_buffers.sh similarity index 100% rename from DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_sum_oneapi_reduction_buffers.sh rename to DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_sum_reduction_buffers.sh diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_sum_oneapi_reduction_usm.sh b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_sum_reduction_usm.sh similarity index 100% rename from DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_sum_oneapi_reduction_usm.sh rename to DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/run_sum_reduction_usm.sh diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/multiple_reductions_buffers.cpp b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/multiple_reductions_buffers.cpp index 3f7f939c54..b2ee47eb6d 100644 --- a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/multiple_reductions_buffers.cpp +++ b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/multiple_reductions_buffers.cpp @@ -28,14 +28,11 @@ int main() { q.submit([&](handler& h) { //# create accessors for data and results accessor acc_data(buf_data, h, read_only); - accessor acc_sum(buf_sum, h); - accessor acc_min(buf_min, h); - accessor acc_max(buf_max, h); //# define reduction objects for sum, min, max reduction - auto reduction_sum = ext::oneapi::reduction(acc_sum, 0, plus<>()); - auto reduction_min = ext::oneapi::reduction(acc_min, 0, minimum<>()); - auto reduction_max = ext::oneapi::reduction(acc_max, 0, maximum<>()); + auto reduction_sum = reduction(buf_sum, h, plus<>()); + auto reduction_min = reduction(buf_min, h, minimum<>()); + auto reduction_max = reduction(buf_max, h, maximum<>()); //# parallel_for with multiple reduction objects h.parallel_for(nd_range<1>{N, B}, reduction_sum, reduction_min, reduction_max, [=](nd_item<1> it, auto& temp_sum, auto& temp_min, auto& temp_max) { diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/oneapi_reduction_custom_operator.cpp b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/reduction_custom_operator.cpp similarity index 93% rename from DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/oneapi_reduction_custom_operator.cpp rename to DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/reduction_custom_operator.cpp index 23e7dfffa9..4f4d76e7e3 100644 --- a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/oneapi_reduction_custom_operator.cpp +++ b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/reduction_custom_operator.cpp @@ -37,7 +37,7 @@ int main() { //# custom operator for reduction to find minumum and index pair operator_identity = {std::numeric_limits::max(), std::numeric_limits::min()}; *result = operator_identity; - auto reduction_object = ext::oneapi::reduction(result, operator_identity, minimum>()); + auto reduction_object = reduction(result, operator_identity, minimum>()); //# parallel_for with user defined reduction object q.parallel_for(nd_range<1>{N, B}, reduction_object, [=](nd_item<1> item, auto& temp) { diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/sum_oneapi_reduction_buffers.cpp b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/sum_reduction_buffers.cpp similarity index 100% rename from DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/sum_oneapi_reduction_buffers.cpp rename to DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/sum_reduction_buffers.cpp diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/sum_oneapi_reduction_usm.cpp b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/sum_reduction_usm.cpp similarity index 100% rename from DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/sum_oneapi_reduction_usm.cpp rename to DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/08_DPCPP_Reduction/src/sum_reduction_usm.cpp diff --git a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/Makefile b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/Makefile index 2c6226e6c3..964c35c2d0 100644 --- a/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/Makefile +++ b/DirectProgramming/DPC++/Jupyter/oneapi-essentials-training/Makefile @@ -96,14 +96,14 @@ prog31: 07_DPCPP_Library/lab/maximum_function.cpp prog32: 07_DPCPP_Library/lab/minimum_function.cpp dpcpp -o 07_DPCPP_Library/bin/prog32 07_DPCPP_Library/lab/minimum_function.cpp -prog33: 08_DPCPP_Reduction/lab/oneapi_reduction_custom_operator.cpp - dpcpp -o 08_DPCPP_Reduction/prog33 08_DPCPP_Reduction/lab/oneapi_reduction_custom_operator.cpp +prog33: 08_DPCPP_Reduction/lab/reduction_custom_operator.cpp + dpcpp -o 08_DPCPP_Reduction/prog33 08_DPCPP_Reduction/lab/reduction_custom_operator.cpp -prog34: 08_DPCPP_Reduction/lab/sum_oneapi_reduction_buffers.cpp - dpcpp -o 08_DPCPP_Reduction/prog34 08_DPCPP_Reduction/lab/sum_oneapi_reduction_buffers.cpp +prog34: 08_DPCPP_Reduction/lab/sum_reduction_buffers.cpp + dpcpp -o 08_DPCPP_Reduction/prog34 08_DPCPP_Reduction/lab/sum_reduction_buffers.cpp -prog35: 08_DPCPP_Reduction/lab/sum_oneapi_reduction_usm.cpp - dpcpp -o 08_DPCPP_Reduction/prog35 08_DPCPP_Reduction/lab/sum_oneapi_reduction_usm.cpp +prog35: 08_DPCPP_Reduction/lab/sum_reduction_usm.cpp + dpcpp -o 08_DPCPP_Reduction/prog35 08_DPCPP_Reduction/lab/sum_reduction_usm.cpp prog36: 08_DPCPP_Reduction/lab/sum_single_task.cpp dpcpp -o 08_DPCPP_Reduction/prog36 08_DPCPP_Reduction/lab/sum_single_task.cpp