Add bin and lab subfolders to examples/SC20

Committing oneTBB tutorials for Supercomputing Conference 2020
Committing oneTBB tutorials CppCon2020
2020-11-06 10:45:56 +03:00 · 2020-11-05 15:30:59 +03:00 · 2020-09-18 17:57:53 +03:00
210 changed files with 734862 additions and 0 deletions
@@ -0,0 +1,838 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# FPGA Development Flow and optimizations"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "##### Sections\n",
+    "- [oneAPI with Intel® FPGAs](#oneAPI-with-Intel®-FPGAs)\n",
+    "- [Stage 1: Emulation](#Stage-1:-Emulation)\n",
+    "- [Stage 2: Optimization Report Generation](#Stage-2:-Optimization-Report-Generation)\n",
+    "- [Stage 3: Optimizing the previous code](#Stage-3:-Optimizing-the-previous-code)\n",
+    "- [References to Learn More](#References-to-Learn-More)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "***\n",
+    "# oneAPI with Intel® FPGAs\n",
+    "\n",
+    "The development flow for Intel FPGAs with oneAPI includes multiple stages. The purpose of these stages is to:\n",
+    "* Ensure functionality of your code (you get the correct answers from your computation)\n",
+    "* Ensure the custom hardware built to implement your code has optimal performance\n",
+    "\n",
+    "Without having to endure the lengthy compile to a full FPGA executable each time.\n",
+    "\n",
+    "The detailed flow is shown in the diagram below.\n",
+    "\n",
+    "In this lab, we will practice the first 2 stages of the flow - emulating your code to make sure your code is functional, and generating an optimization report to see how optimized the hardware image generated from your code is. (A subsequent lab will give you practice working with the optimization report.)\n",
+    "\n",
+    "<img src=\"assets/fpga_flow.png\">"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "***\n",
+    "# Stage 1: Emulation\n",
+    "\n",
+    "The first stage developing code for FPGAs with oneAPI is __emulation__. The purpose of emulation is to make sure that your code is __functional__, or in other words, that you __get the correct answers from your computations__.\n",
+    "\n",
+    "The compile time for this stage will be very quick, usually seconds.\n",
+    "\n",
+    "This quick compile time allows you to iterate through this stage many times, until your code is functionally correct.\n",
+    "\n",
+    "__Now, let's give it a try!__\n",
+    "\n",
+    "The code below just adds two vectors. Please have a quick look to get an overview of what it is doing.\n",
+    "\n",
+    "We will continue using this simple piece of code to learn about the development flow for FPGAs with oneAPI and howto apply optimizations.\n",
+    "\n",
+    "__After you are finished examining the code, click ▶ to save the code to a file.__"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": 1,
+   "metadata": {},
+   "outputs": [
+    {
+     "name": "stdout",
+     "output_type": "stream",
+     "text": [
+      "Writing lab/loop.cpp\n"
+     ]
+    }
+   ],
+   "source": [
+    "%%writefile lab/loop.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "#include <CL/sycl.hpp>\n",
+    "#include <iomanip>\n",
+    "#include <iostream>\n",
+    "#include <vector>\n",
+    "\n",
+    "// dpc_common.hpp can be found in the dev-utilities include folder.\n",
+    "// e.g., $ONEAPI_ROOT/dev-utilities/latest/include/dpc_common.hpp\n",
+    "#include \"dpc_common.hpp\"\n",
+    "\n",
+    "// Header locations and some DPC++ extensions changed between beta09 and beta10\n",
+    "// Temporarily modify the code sample to accept either version\n",
+    "#define BETA09 20200827\n",
+    "#if __SYCL_COMPILER_VERSION <= BETA09\n",
+    "  #include <CL/sycl/intel/fpga_extensions.hpp>\n",
+    "  namespace INTEL = sycl::intel;  // Namespace alias for backward compatibility\n",
+    "#else\n",
+    "  #include <CL/sycl/INTEL/fpga_extensions.hpp>\n",
+    "#endif\n",
+    "\n",
+    "using namespace sycl;\n",
+    "\n",
+    "// This function instantiates the vector add kernel, which contains\n",
+    "// a loop that adds up the two summand arrays and stores the result\n",
+    "// into sum. \n",
+    "void VecAdd(const std::vector<float> &summands1,\n",
+    "            const std::vector<float> &summands2, std::vector<float> &sum,\n",
+    "            size_t array_size) {\n",
+    "\n",
+    "\n",
+    "#if defined(FPGA_EMULATOR)\n",
+    "  INTEL::fpga_emulator_selector device_selector;\n",
+    "#else\n",
+    "  INTEL::fpga_selector device_selector;\n",
+    "#endif\n",
+    "\n",
+    "  try {\n",
+    "    queue q(device_selector, dpc_common::exception_handler,\n",
+    "            property::queue::enable_profiling{});\n",
+    "\n",
+    "    buffer buffer_summands1(summands1);\n",
+    "    buffer buffer_summands2(summands2);\n",
+    "    // Use verbose SYCL 1.2 syntax for the output buffer.\n",
+    "    // (This will become unnecessary in a future compiler version.)\n",
+    "    buffer<float, 1> buffer_sum(sum.data(), array_size);\n",
+    "\n",
+    "    event e = q.submit([&](handler &h) {\n",
+    "      auto acc_summands1 = buffer_summands1.get_access<access::mode::read>(h);\n",
+    "      auto acc_summands2 = buffer_summands2.get_access<access::mode::read>(h);\n",
+    "      auto acc_sum = buffer_sum.get_access<access::mode::discard_write>(h);\n",
+    "\n",
+    "      h.single_task([=]() {\n",
+    "        for (size_t i = 0; i < array_size; i++) {\n",
+    "          acc_sum[i] = acc_summands1[i] + acc_summands2[i];\n",
+    "        }\n",
+    "      });\n",
+    "    });\n",
+    "\n",
+    "    double start = e.get_profiling_info<info::event_profiling::command_start>();\n",
+    "    double end = e.get_profiling_info<info::event_profiling::command_end>();\n",
+    "    // convert from nanoseconds to ms\n",
+    "    double kernel_time = (double)(end - start) * 1e-6;\n",
+    "\n",
+    "    std::cout << \" kernel time : \" << kernel_time << \" ms\\n\";\n",
+    "    std::cout << \"Throughput for kernel \" << \": \";\n",
+    "    std::cout << std::fixed << std::setprecision(3)\n",
+    "              << ((double)array_size / kernel_time) / 1e6f << \" GFlops\\n\";\n",
+    "\n",
+    "  } catch (sycl::exception const &e) {\n",
+    "    // Catches exceptions in the host code\n",
+    "    std::cout << \"Caught a SYCL host exception:\\n\" << e.what() << \"\\n\";\n",
+    "\n",
+    "    // Most likely the runtime couldn't find FPGA hardware!\n",
+    "    if (e.get_cl_code() == CL_DEVICE_NOT_FOUND) {\n",
+    "      std::cout << \"If you are targeting an FPGA, please ensure that your \"\n",
+    "                   \"system has a correctly configured FPGA board.\\n\";\n",
+    "      std::cout << \"If you are targeting the FPGA emulator, compile with \"\n",
+    "                   \"-DFPGA_EMULATOR.\\n\";\n",
+    "    }\n",
+    "    std::terminate();\n",
+    "  }\n",
+    "}\n",
+    "\n",
+    "int main(int argc, char *argv[]) {\n",
+    "  size_t array_size = 1 << 26;\n",
+    "\n",
+    "  if (argc > 1) {\n",
+    "    std::string option(argv[1]);\n",
+    "    if (option == \"-h\" || option == \"--help\") {\n",
+    "      std::cout << \"Usage: \\n<executable> <data size>\\n\\nFAILED\\n\";\n",
+    "      return 1;\n",
+    "    } else {\n",
+    "      array_size = std::stoi(option);\n",
+    "    }\n",
+    "  }\n",
+    "\n",
+    "  std::vector<float> summands1(array_size);\n",
+    "  std::vector<float> summands2(array_size);\n",
+    "\n",
+    "  std::vector<float> sum(array_size);\n",
+    "\n",
+    "  // Initialize the two summand arrays (arrays to be added to each other) to\n",
+    "  // 1:N and N:1, so that the sum of all elements is N + 1\n",
+    "  for (size_t i = 0; i < array_size; i++) {\n",
+    "    summands1[i] = static_cast<float>(i + 1);\n",
+    "    summands2[i] = static_cast<float>(array_size - i);\n",
+    "  }\n",
+    "\n",
+    "  std::cout << \"Input Array Size:  \" << array_size << \"\\n\";\n",
+    "\n",
+    "  // Instantiate VecAdd kernel that contains a loop that adds up the two summand arrays.\n",
+    "  VecAdd(summands1, summands2, sum, array_size);\n",
+    "\n",
+    "  // Verify that the output data is the same for every unroll factor\n",
+    "  for (size_t i = 0; i < array_size; i++) {\n",
+    "    if (sum[i] != summands1[i] + summands2[i]) {\n",
+    "      std::cout << \"FAILED: The results are incorrect\\n\";\n",
+    "      return 1;\n",
+    "    }\n",
+    "  }\n",
+    "  std::cout << \"PASSED: The results are correct\\n\";\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "__Now, you will compile the code to target the FPGA emulator.__\n"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! echo \"##\" $(whoami) is compiling a simple FPGA code with a loop\n",
+    "! dpcpp -fintelfpga lab/loop.cpp -DFPGA_EMULATOR -o bin/loop.emu\n",
+    "! bin/loop.emu"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "***\n",
+    "## Stage 2: Optimization Report Generation\n",
+    "\n",
+    "In this next section of the lab, you will compile the kernel using different command line options with Intel's oneAPI DPC++ compiler in order to create an optimization report. You will also be using the Jupyter Lab interface to browse and open a file, so that will be explained to you.\n"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "__Let's compile the code and generate an optimization report.__\n",
+    "\n",
+    "The commands you need to do this are shown below. (We are using the two-step method since there is a current issue showing the source code in the report with the one-step method.)\n",
+    "\n",
+    "```\n",
+    "dpcpp -fintelfpga lab/loop.cpp -c -o bin/loop.o\n",
+    "dpcpp -fintelfpga bin/loop.o -fsycl-link -Xshardware -o bin/loop.a\n",
+    "```\n",
+    "\n",
+    "__This compilation may take approximately 2 minutes.__"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! rm -rf bin/loop.a\n",
+    "! echo \"##\" $(whoami) is working on generating the report for loop.cpp lab step 2\n",
+    "! dpcpp -fintelfpga lab/loop.cpp -c -o bin/loop.o\n",
+    "! dpcpp -fintelfpga bin/loop.o -fsycl-link -Xshardware -o bin/loop.a\n",
+    "! echo \"Compilation done.\""
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "You may wonder if these two previous dpcpp calls can be summarized into a single one. The answer is yes, and it would be:\n",
+    "```\n",
+    "dpcpp -fintelfpga lab/loop.cpp -fsycl-link -Xshardware -o bin/loop.a\n",
+    "```\n",
+    "But then you may get a warning saying:\n",
+    "```\n",
+    "aoc: Warning: Cannot find dependency file \"/home/u32284/tmp/loop.d\" for source file \"/home/u32284/tmp/loop-172fb0.spv\". Source code will not be available in the HLD Reports. Ensure you ran dpcpp with the -fintelfpga flag.\n",
+    "```\n",
+    "Which means that you won't see the `source code` in the resulting reports. Try at your own risk."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "! echo \"##\" $(whoami) is working on generating the report for loop.cpp lab step 2\n",
+    "! dpcpp -fintelfpga lab/loop.cpp -fsycl-link -Xshardware -o bin/loop.a\n",
+    "! echo \"The compile is finished.\""
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "__When you see the \"Compilation done.\" message above, an optimization report file will have been generated for our example. You will see a warning if you compile more than once. The warning can be ignored.__\n",
+    "\n",
+    "Now, let's examine the generated report file.\n",
+    "\n",
+    "Within the Jupyter Lab environment, it’s easy to navigate to the generated optimization report through the file browser on the leftmost panel. Browse to the directory `bin/loop.prj/reports/` (Double click on directories to open.)\n",
+    "\n",
+    "__Double-click on report.html.__ The report will open up as another tab beside the notebook tab in Jupyter Lab, as shown below.\n",
+    "\n",
+    "__You will need to click on \"Trust HTML\" for the report to display correctly.__\n",
+    "\n",
+    "<img src=\"assets/FPGA-report.png\">\n",
+    "\n",
+    "You've now learned how to use the first two stages of the FPGA development flow with oneAPI! You will spend most of your development time in these two stages. In the next sections of the lab, you will learn different optimization techniques."
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "***\n",
+    "## Stage 3: Optimizing the code for FPGA\n",
+    "\n",
+    "As you may have noticed, the `Initiation Interval` is too high (`II=286`) and as a result the FPGA device would be running at 240Mz which is not ideal. To overcome those kind of issues, in this exercise we will exemplarily apply three optimizations:\n",
+    "\n",
+    "1. Give a hint to the compiler promising that the vectors `acc_sum`, `acc_summands1` and `acc_summands2` do not share elements (they are non-overlapping regions of memory) and therefore they do not result in loop-carried dependencies.\n",
+    "2. Unroll the loop to have more pipelines processing elements of the arrays per clock cycle\n",
+    "3. Replicate the kernel to better utilize the FPGA HW resources and also enabling more operations per sencod. For the sake of expediency, we will replicate the kernel but changing the unroll factor for each replica. Replication is implemented in modern C++ using templated functions and classes.\n",
+    "\n",
+    "In the following code, some changes are needed:\n",
+    "\n",
+    "1. Inspect the code cell below and make the following modifications.\n",
+    "  1. STEP A: Insert in the q.single_task the intel::kernel_args_restrict C++ attribute (as in Slide 16 in 10-FPGA-Optimizations)\n",
+    "  2. STEP B: Insert in the q.single_task the pragma unroll using the template argument unroll_factor (as in Slide 13) \n",
+    "  3. STEP C: Add in the main() function a fith invocation to VecAdd now using unroll_factor=16 and sum_unrollx16\n",
+    "2. When the modifications are complete, click run ▶ to save the code to a file.\n",
+    "3. Run ▶ the cell in the __Build and check the report of the modified code__ section below the code snippet to compile and execute the code in the saved file.\n",
+    "\n"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": 2,
+   "metadata": {},
+   "outputs": [
+    {
+     "name": "stdout",
+     "output_type": "stream",
+     "text": [
+      "Writing lab/loop_unroll.cpp\n"
+     ]
+    }
+   ],
+   "source": [
+    "%%writefile lab/loop_unroll.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "#include <CL/sycl.hpp>\n",
+    "#include <iomanip>\n",
+    "#include <iostream>\n",
+    "#include <vector>\n",
+    "\n",
+    "// dpc_common.hpp can be found in the dev-utilities include folder.\n",
+    "// e.g., $ONEAPI_ROOT/dev-utilities//include/dpc_common.hpp\n",
+    "#include \"dpc_common.hpp\"\n",
+    "\n",
+    "// Header locations and some DPC++ extensions changed between beta09 and beta10\n",
+    "// Temporarily modify the code sample to accept either version\n",
+    "#define BETA09 20200827\n",
+    "#if __SYCL_COMPILER_VERSION <= BETA09\n",
+    "  #include <CL/sycl/intel/fpga_extensions.hpp>\n",
+    "  namespace INTEL = sycl::intel;  // Namespace alias for backward compatibility\n",
+    "#else\n",
+    "  #include <CL/sycl/INTEL/fpga_extensions.hpp>\n",
+    "#endif\n",
+    "\n",
+    "using namespace sycl;\n",
+    "\n",
+    "template <int unroll_factor> class VAdd;\n",
+    "\n",
+    "// This function instantiates the vector add kernel, which contains\n",
+    "// a loop that adds up the two summand arrays and stores the result\n",
+    "// into sum. This loop will be unrolled by the specified unroll_factor.\n",
+    "template <int unroll_factor>\n",
+    "void VecAdd(const std::vector<float> &summands1,\n",
+    "            const std::vector<float> &summands2, std::vector<float> &sum,\n",
+    "            size_t array_size) {\n",
+    "\n",
+    "\n",
+    "#if defined(FPGA_EMULATOR)\n",
+    "  INTEL::fpga_emulator_selector device_selector;\n",
+    "#else\n",
+    "  INTEL::fpga_selector device_selector;\n",
+    "#endif\n",
+    "\n",
+    "  try {\n",
+    "    queue q(device_selector, dpc_common::exception_handler,\n",
+    "            property::queue::enable_profiling{});\n",
+    "\n",
+    "    buffer buffer_summands1(summands1);\n",
+    "    buffer buffer_summands2(summands2);\n",
+    "    // Use verbose SYCL 1.2 syntax for the output buffer.\n",
+    "    // (This will become unnecessary in a future compiler version.)\n",
+    "    buffer<float, 1> buffer_sum(sum.data(), array_size);\n",
+    "\n",
+    "    event e = q.submit([&](handler &h) {\n",
+    "      auto acc_summands1 = buffer_summands1.get_access<access::mode::read>(h);\n",
+    "      auto acc_summands2 = buffer_summands2.get_access<access::mode::read>(h);\n",
+    "      auto acc_sum = buffer_sum.get_access<access::mode::discard_write>(h);\n",
+    "\n",
+    "//This template argument to single_task allows for the correct name_mangling for all the kernels\n",
+    "      h.single_task<VAdd<unroll_factor>>([=]()\n",
+    "//STEP A: Insert here the intel::kernel_args_restrict C++ attribute (as in Slide 16 in 10-FPGA-Optimizations)\n",
+    "                                         {\n",
+    "        // Unroll the loop fully or partially, depending on unroll_factor\n",
+    "//STEP B: Insert here the pragma unroll using the template argument unroll_factor (as in Slide 13)\n",
+    "        for (size_t i = 0; i < array_size; i++) {\n",
+    "          acc_sum[i] = acc_summands1[i] + acc_summands2[i];\n",
+    "        }\n",
+    "      });\n",
+    "    });\n",
+    "\n",
+    "    double start = e.get_profiling_info<info::event_profiling::command_start>();\n",
+    "    double end = e.get_profiling_info<info::event_profiling::command_end>();\n",
+    "    // convert from nanoseconds to ms\n",
+    "    double kernel_time = (double)(end - start) * 1e-6;\n",
+    "\n",
+    "    std::cout << \"unroll_factor \" << unroll_factor\n",
+    "              << \" kernel time : \" << kernel_time << \" ms\\n\";\n",
+    "    std::cout << \"Throughput for kernel with unroll_factor \" << unroll_factor\n",
+    "              << \": \";\n",
+    "    std::cout << std::fixed << std::setprecision(3)\n",
+    "              << ((double)array_size / kernel_time) / 1e6f << \" GFlops\\n\";\n",
+    "\n",
+    "  } catch (sycl::exception const &e) {\n",
+    "    // Catches exceptions in the host code\n",
+    "    std::cout << \"Caught a SYCL host exception:\\n\" << e.what() << \"\\n\";\n",
+    "\n",
+    "    // Most likely the runtime couldn't find FPGA hardware!\n",
+    "    if (e.get_cl_code() == CL_DEVICE_NOT_FOUND) {\n",
+    "      std::cout << \"If you are targeting an FPGA, please ensure that your \"\n",
+    "                   \"system has a correctly configured FPGA board.\\n\";\n",
+    "      std::cout << \"If you are targeting the FPGA emulator, compile with \"\n",
+    "                   \"-DFPGA_EMULATOR.\\n\";\n",
+    "    }\n",
+    "    std::terminate();\n",
+    "  }\n",
+    "}\n",
+    "\n",
+    "int main(int argc, char *argv[]) {\n",
+    "  size_t array_size = 1 << 26;\n",
+    "\n",
+    "  if (argc > 1) {\n",
+    "    std::string option(argv[1]);\n",
+    "    if (option == \"-h\" || option == \"--help\") {\n",
+    "      std::cout << \"Usage: \\n<executable> <data size>\\n\\nFAILED\\n\";\n",
+    "      return 1;\n",
+    "    } else {\n",
+    "      array_size = std::stoi(option);\n",
+    "    }\n",
+    "  }\n",
+    "\n",
+    "  std::vector<float> summands1(array_size);\n",
+    "  std::vector<float> summands2(array_size);\n",
+    "\n",
+    "  std::vector<float> sum_unrollx1(array_size);\n",
+    "  std::vector<float> sum_unrollx2(array_size);\n",
+    "  std::vector<float> sum_unrollx4(array_size);\n",
+    "  std::vector<float> sum_unrollx8(array_size);\n",
+    "  std::vector<float> sum_unrollx16(array_size);\n",
+    "\n",
+    "  // Initialize the two summand arrays (arrays to be added to each other) to\n",
+    "  // 1:N and N:1, so that the sum of all elements is N + 1\n",
+    "  for (size_t i = 0; i < array_size; i++) {\n",
+    "    summands1[i] = static_cast<float>(i + 1);\n",
+    "    summands2[i] = static_cast<float>(array_size - i);\n",
+    "  }\n",
+    "\n",
+    "  std::cout << \"Input Array Size:  \" << array_size << \"\\n\";\n",
+    "\n",
+    "  // Instantiate VecAdd kernel with different unroll factors: 1, 2, 4, 8, 16\n",
+    "  // The VecAdd kernel contains a loop that adds up the two summand arrays.\n",
+    "  // This loop will be unrolled by the specified unroll factor.\n",
+    "  // The sum array is expected to be identical, regardless of the unroll factor.\n",
+    "  VecAdd<1>(summands1, summands2, sum_unrollx1, array_size);\n",
+    "  VecAdd<2>(summands1, summands2, sum_unrollx2, array_size);\n",
+    "  VecAdd<4>(summands1, summands2, sum_unrollx4, array_size);\n",
+    "  VecAdd<8>(summands1, summands2, sum_unrollx8, array_size);\n",
+    "\n",
+    "//STEP C: Add a fith invocation to VecAdd now using unroll_factor=16 and sum_unrollx16\n",
+    "\n",
+    "  // Verify that the output data is the same for every unroll factor\n",
+    "  for (size_t i = 0; i < array_size; i++) {\n",
+    "    if (sum_unrollx1[i] != summands1[i] + summands2[i] ||\n",
+    "        sum_unrollx1[i] != sum_unrollx2[i] ||\n",
+    "        sum_unrollx1[i] != sum_unrollx4[i] ||\n",
+    "        sum_unrollx1[i] != sum_unrollx8[i] ||\n",
+    "        sum_unrollx1[i] != sum_unrollx16[i]) {\n",
+    "      std::cout << \"FAILED: The results are incorrect\\n\";\n",
+    "      return 1;\n",
+    "    }\n",
+    "  }\n",
+    "  std::cout << \"PASSED: The results are correct\\n\";\n",
+    "  return 0;\n",
+    "}\n"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and check the report of the modified code\n",
+    "\n",
+    "Select the cell below and click Run ▶ to compile and generate the report for the code that you modified above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! rm -rf lab/loop_unroll.a\n",
+    "! echo \"##\" $(whoami) is working on optimizing with loop unrolling\n",
+    "! dpcpp -fintelfpga lab/loop_unroll.cpp -c -o lab/loop_unroll.o\n",
+    "! dpcpp -fintelfpga lab/loop_unroll.o -fsycl-link -Xshardware -o lab/loop_unroll.a\n",
+    "! echo \"The compile is finished.\""
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "Browse to the directory `lab`/loop_unroll.prj/reports/ (Double click on directories to open them.)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Solution (Don't peak unless you have to)"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": 3,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [
+    {
+     "name": "stdout",
+     "output_type": "stream",
+     "text": [
+      "Writing solutions/loop_unroll.cpp\n"
+     ]
+    }
+   ],
+   "source": [
+    "%%writefile solutions/loop_unroll.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "#include <CL/sycl.hpp>\n",
+    "#include <iomanip>\n",
+    "#include <iostream>\n",
+    "#include <vector>\n",
+    "\n",
+    "// dpc_common.hpp can be found in the dev-utilities include folder.\n",
+    "// e.g., $ONEAPI_ROOT/dev-utilities//include/dpc_common.hpp\n",
+    "#include \"dpc_common.hpp\"\n",
+    "\n",
+    "// Header locations and some DPC++ extensions changed between beta09 and beta10\n",
+    "// Temporarily modify the code sample to accept either version\n",
+    "#define BETA09 20200827\n",
+    "#if __SYCL_COMPILER_VERSION <= BETA09\n",
+    "  #include <CL/sycl/intel/fpga_extensions.hpp>\n",
+    "  namespace INTEL = sycl::intel;  // Namespace alias for backward compatibility\n",
+    "#else\n",
+    "  #include <CL/sycl/INTEL/fpga_extensions.hpp>\n",
+    "#endif\n",
+    "\n",
+    "using namespace sycl;\n",
+    "\n",
+    "template <int unroll_factor> class VAdd;\n",
+    "\n",
+    "// This function instantiates the vector add kernel, which contains\n",
+    "// a loop that adds up the two summand arrays and stores the result\n",
+    "// into sum. This loop will be unrolled by the specified unroll_factor.\n",
+    "template <int unroll_factor>\n",
+    "void VecAdd(const std::vector<float> &summands1,\n",
+    "            const std::vector<float> &summands2, std::vector<float> &sum,\n",
+    "            size_t array_size) {\n",
+    "\n",
+    "\n",
+    "#if defined(FPGA_EMULATOR)\n",
+    "  INTEL::fpga_emulator_selector device_selector;\n",
+    "#else\n",
+    "  INTEL::fpga_selector device_selector;\n",
+    "#endif\n",
+    "\n",
+    "  try {\n",
+    "    queue q(device_selector, dpc_common::exception_handler,\n",
+    "            property::queue::enable_profiling{});\n",
+    "\n",
+    "    buffer buffer_summands1(summands1);\n",
+    "    buffer buffer_summands2(summands2);\n",
+    "    // Use verbose SYCL 1.2 syntax for the output buffer.\n",
+    "    // (This will become unnecessary in a future compiler version.)\n",
+    "    buffer<float, 1> buffer_sum(sum.data(), array_size);\n",
+    "\n",
+    "    event e = q.submit([&](handler &h) {\n",
+    "      auto acc_summands1 = buffer_summands1.get_access<access::mode::read>(h);\n",
+    "      auto acc_summands2 = buffer_summands2.get_access<access::mode::read>(h);\n",
+    "      auto acc_sum = buffer_sum.get_access<access::mode::discard_write>(h);\n",
+    "\n",
+    "      h.single_task<VAdd<unroll_factor>>([=]()\n",
+    "                                         [[intel::kernel_args_restrict]] {\n",
+    "        // Unroll the loop fully or partially, depending on unroll_factor\n",
+    "        #pragma unroll unroll_factor\n",
+    "        for (size_t i = 0; i < array_size; i++) {\n",
+    "          acc_sum[i] = acc_summands1[i] + acc_summands2[i];\n",
+    "        }\n",
+    "      });\n",
+    "    });\n",
+    "\n",
+    "    double start = e.get_profiling_info<info::event_profiling::command_start>();\n",
+    "    double end = e.get_profiling_info<info::event_profiling::command_end>();\n",
+    "    // convert from nanoseconds to ms\n",
+    "    double kernel_time = (double)(end - start) * 1e-6;\n",
+    "\n",
+    "    std::cout << \"unroll_factor \" << unroll_factor\n",
+    "              << \" kernel time : \" << kernel_time << \" ms\\n\";\n",
+    "    std::cout << \"Throughput for kernel with unroll_factor \" << unroll_factor\n",
+    "              << \": \";\n",
+    "    std::cout << std::fixed << std::setprecision(3)\n",
+    "              << ((double)array_size / kernel_time) / 1e6f << \" GFlops\\n\";\n",
+    "\n",
+    "  } catch (sycl::exception const &e) {\n",
+    "    // Catches exceptions in the host code\n",
+    "    std::cout << \"Caught a SYCL host exception:\\n\" << e.what() << \"\\n\";\n",
+    "\n",
+    "    // Most likely the runtime couldn't find FPGA hardware!\n",
+    "    if (e.get_cl_code() == CL_DEVICE_NOT_FOUND) {\n",
+    "      std::cout << \"If you are targeting an FPGA, please ensure that your \"\n",
+    "                   \"system has a correctly configured FPGA board.\\n\";\n",
+    "      std::cout << \"If you are targeting the FPGA emulator, compile with \"\n",
+    "                   \"-DFPGA_EMULATOR.\\n\";\n",
+    "    }\n",
+    "    std::terminate();\n",
+    "  }\n",
+    "}\n",
+    "\n",
+    "int main(int argc, char *argv[]) {\n",
+    "  size_t array_size = 1 << 26;\n",
+    "\n",
+    "  if (argc > 1) {\n",
+    "    std::string option(argv[1]);\n",
+    "    if (option == \"-h\" || option == \"--help\") {\n",
+    "      std::cout << \"Usage: \\n<executable> <data size>\\n\\nFAILED\\n\";\n",
+    "      return 1;\n",
+    "    } else {\n",
+    "      array_size = std::stoi(option);\n",
+    "    }\n",
+    "  }\n",
+    "\n",
+    "  std::vector<float> summands1(array_size);\n",
+    "  std::vector<float> summands2(array_size);\n",
+    "\n",
+    "  std::vector<float> sum_unrollx1(array_size);\n",
+    "  std::vector<float> sum_unrollx2(array_size);\n",
+    "  std::vector<float> sum_unrollx4(array_size);\n",
+    "  std::vector<float> sum_unrollx8(array_size);\n",
+    "  std::vector<float> sum_unrollx16(array_size);\n",
+    "\n",
+    "  // Initialize the two summand arrays (arrays to be added to each other) to\n",
+    "  // 1:N and N:1, so that the sum of all elements is N + 1\n",
+    "  for (size_t i = 0; i < array_size; i++) {\n",
+    "    summands1[i] = static_cast<float>(i + 1);\n",
+    "    summands2[i] = static_cast<float>(array_size - i);\n",
+    "  }\n",
+    "\n",
+    "  std::cout << \"Input Array Size:  \" << array_size << \"\\n\";\n",
+    "\n",
+    "  // Instantiate VecAdd kernel with different unroll factors: 1, 2, 4, 8, 16\n",
+    "  // The VecAdd kernel contains a loop that adds up the two summand arrays.\n",
+    "  // This loop will be unrolled by the specified unroll factor.\n",
+    "  // The sum array is expected to be identical, regardless of the unroll factor.\n",
+    "  VecAdd<1>(summands1, summands2, sum_unrollx1, array_size);\n",
+    "  VecAdd<2>(summands1, summands2, sum_unrollx2, array_size);\n",
+    "  VecAdd<4>(summands1, summands2, sum_unrollx4, array_size);\n",
+    "  VecAdd<8>(summands1, summands2, sum_unrollx8, array_size);\n",
+    "  VecAdd<16>(summands1, summands2, sum_unrollx16, array_size);\n",
+    "\n",
+    "  // Verify that the output data is the same for every unroll factor\n",
+    "  for (size_t i = 0; i < array_size; i++) {\n",
+    "    if (sum_unrollx1[i] != summands1[i] + summands2[i] ||\n",
+    "        sum_unrollx1[i] != sum_unrollx2[i] ||\n",
+    "        sum_unrollx1[i] != sum_unrollx4[i] ||\n",
+    "        sum_unrollx1[i] != sum_unrollx8[i] ||\n",
+    "        sum_unrollx1[i] != sum_unrollx16[i]) {\n",
+    "      std::cout << \"FAILED: The results are incorrect\\n\";\n",
+    "      return 1;\n",
+    "    }\n",
+    "  }\n",
+    "  std::cout << \"PASSED: The results are correct\\n\";\n",
+    "  return 0;\n",
+    "}\n"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "Generate the report:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! rm -rf solutions/loop_unroll.a\n",
+    "! echo \"##\" $(whoami) is working on optimizing with loop unrolling\n",
+    "! dpcpp -fintelfpga solutions/loop_unroll.cpp -c -o solutions/loop_unroll.o\n",
+    "! dpcpp -fintelfpga solutions/loop_unroll.o -fsycl-link -Xshardware -o solutions/loop_unroll.a\n",
+    "! echo \"The compile is finished.\""
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "Browse to the directory `solutions`/loop_unroll.prj/reports/ (Double click on diretories to push down into them.)\n",
+    "\n",
+    "The report should now look like:\n",
+    "\n",
+    "<img src=\"assets/FPGA-unroll-report.png\">\n"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "Feel free to compile it as we did in the first hands-on session and you can expect a similar output running on an Intel Arria10 FPGA that is available in DevCloud:\n",
+    "```\n",
+    "Input Array Size:  67108864\n",
+    "unroll_factor 1 kernel time : 242.849 ms\n",
+    "Throughput for kernel with unroll_factor 1: 0.276 GFlops\n",
+    "unroll_factor 2 kernel time : 122.319 ms\n",
+    "Throughput for kernel with unroll_factor 2: 0.549 GFlops\n",
+    "unroll_factor 4 kernel time : 63.950 ms\n",
+    "Throughput for kernel with unroll_factor 4: 1.049 GFlops\n",
+    "unroll_factor 8 kernel time : 39.567 ms\n",
+    "Throughput for kernel with unroll_factor 8: 1.696 GFlops\n",
+    "unroll_factor 16 kernel time : 37.500 ms\n",
+    "Throughput for kernel with unroll_factor 16: 1.790 GFlops\n",
+    "PASSED: The results are correct\n",
+    "```\n"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "***\n",
+    "## References to Learn More\n",
+    "\n",
+    "Please refer to the following resources to learn more. This is a great thing to do if you have extra time during the lab!\n",
+    "\n",
+    "#### FPGA Specific Documentation\n",
+    "\n",
+    "* [Website hub for using FPGAs with oneAPI](https://software.intel.com/content/www/us/en/develop/tools/oneapi/components/fpga.html)\n",
+    "* [Intel® oneAPI Programming Guide](https://software.intel.com/content/www/us/en/develop/download/intel-oneapi-programming-guide.html)\n",
+    "* [Intel® oneAPI DPC++ FPGA Optimization Guide](https://software.intel.com/content/www/us/en/develop/download/oneapi-fpga-optimization-guide.html)\n",
+    "* [oneAPI Samples on GitHub](https://github.com/oneapi-src/oneAPI-samples)\n",
+    "\n",
+    "#### Intel® oneAPI Toolkit documentation\n",
+    "* [Intel® oneAPI main page](https://software.intel.com/oneapi \"oneAPI main page\")\n",
+    "* [Intel® DevCloud Signup](https://software.intel.com/en-us/devcloud/oneapi \"Intel DevCloud\")  Sign up here if you do not have an account.\n",
+    "* [Intel® DevCloud Connect](https://devcloud.intel.com/datacenter/connect)  Login to the DevCloud here.\n",
+    "* [oneAPI Specification elements](https://www.oneapi.com/spec/)\n",
+    "* [DPC++ reference](https://docs.oneapi.com/versions/latest/dpcpp/)\n",
+    "\n",
+    "#### SYCL \n",
+    "* [SYCL* Specification (for version 1.2.1)](https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf)\n",
+    "\n",
+    "#### Modern C++\n",
+    "* [CPPReference](https://en.cppreference.com/w/)\n",
+    "* [CPlusPlus](http://www.cplusplus.com/)"
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3.7 (Intel® oneAPI)",
+   "language": "python",
+   "name": "c009-intel_distribution_of_python_3_oneapi-beta05-python"
+  },
+  "language_info": {
+   "codemirror_mode": {
+    "name": "ipython",
+    "version": 3
+   },
+   "file_extension": ".py",
+   "mimetype": "text/x-python",
+   "name": "python",
+   "nbconvert_exporter": "python",
+   "pygments_lexer": "ipython3",
+   "version": "3.7.7"
+  },
+  "toc": {
+   "base_numbering": 1,
+   "nav_menu": {},
+   "number_sections": true,
+   "sideBar": true,
+   "skip_h1_title": false,
+   "title_cell": "Table of Contents",
+   "title_sidebar": "Contents",
+   "toc_cell": false,
+   "toc_position": {
+    "height": "525.6px",
+    "left": "28px",
+    "top": "137.8px",
+    "width": "301.109px"
+   },
+   "toc_section_display": true,
+   "toc_window_display": true
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 4
+}
@@ -0,0 +1,442 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# Introduction to oneDPL"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "##### Sections\n",
+    "- _Code_: [Hello, PSTL!](#Hello,-PSTL!)\n",
+    "- _Code_: [Fancy Iterators](#Fancy-iterators)\n",
+    "- _Code_: [Monte Carlo Example](#Monte-Carlo-Pi)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "***\n",
+    "# Hello, PSTL!\n",
+    "\n",
+    "To begin, we'll build and run a very simple C++ application that invokes the classic STL algorithm ``sort`` on an input sequence of random integers.\n",
+    "\n",
+    "To make this exercise more interesting we parallelize our sort at a high-level. To do that, we simply add a parallel execution policy and take the algorithm for a second run. Adding a parallel execution policy is as easy as inserting `std::execution::par` as a first parameter. This tells the compiler and runtime that it's safe to execute iterations in parallel on multiple threads of execution. It's a hint by the user but not a mandatory requirement to run any optimizations for the runtime. A high-quality implementation like Parallel STL with TBB as a backend could benefit by splitting the sort into finer grained tasks. Those can be executed by TBB's scheduler in a multithreaded fashion. Of course that adds some overhead, such as maintaining a thread pool or synchronize the tasks. So it's not advised to add a parallel execution policy to each and every algorithm. It's a user's (Yes that's you) responsibility to ensure a parallel optimization hint is legal and beneficial.\n",
+    "\n",
+    "Consequently, in our simple example we time both runs so that we can compare them. Please note that intentionally any warm-up for the parallel execution is skipped.\n",
+    "\n",
+    "Inspect the code below - there are no modifications necessary. Run the first cell to create the file, then run the cell below it to compile and execute the code.\n",
+    "1. Inspect the code cell below, then click run ▶ to save the code to a file\n",
+    "2. Run ▶ the cell in the __Build and Run__ section below the code snippet to compile and execute the code in the saved file\n",
+    "\n",
+    "What are your expectations regarding the speedup? We're likely running on an system with two Intel Xeon processor with 6 cores each (12 hardware threads) in the current DevCloud configuration, as of Sept'20. \n",
+    "\n",
+    "What happens if we decrease the size of the input sequence to let's say 50? "
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/sort.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <iostream>\n",
+    "#include <vector>\n",
+    "#include <random>\n",
+    "\n",
+    "#include <oneapi/dpl/execution>\n",
+    "#include <oneapi/dpl/algorithm>\n",
+    "\n",
+    "int main() {\n",
+    "  std::default_random_engine e1(0);\n",
+    "  std::uniform_int_distribution<int> d(0, 5000);\n",
+    "\n",
+    "  // initialize input vector with random numbers\n",
+    "  std::vector<int> data(5000000);\n",
+    "  for(auto& i:data) {\n",
+    "    i = d(e1);\n",
+    "  }\n",
+    "\n",
+    "  {\n",
+    "    std::vector<int> input(data);\n",
+    "    std::cout << \"Running sequentially:\\n\";\n",
+    "    auto st0 = std::chrono::high_resolution_clock::now();\n",
+    "    std::sort(input.begin(), input.end());\n",
+    "    auto st1 = std::chrono::high_resolution_clock::now();\n",
+    "    std::cout << \"Serial time   = \" << 1e-9 * (st1-st0).count() << \" seconds\\n\";\n",
+    "  }\n",
+    "\n",
+    "  {\n",
+    "    std::vector<int> input(data);\n",
+    "    std::cout << \"\\nRunning in parallel:\\n\";\n",
+    "    auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "    std::sort(std::execution::par, input.begin(), input.end());\n",
+    "    auto pt1 = std::chrono::high_resolution_clock::now();\n",
+    "    std::cout << \"Parallel time = \" << 1e-9 * (pt1-pt0).count() << \" seconds\\n\";\n",
+    "  }\n",
+    "    \n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile scripts/run_sort.sh\n",
+    "#!/bin/bash\n",
+    "#==========================================\n",
+    "# Copyright (c) 2020 Intel Corporation\n",
+    "#\n",
+    "# SPDX-License-Identifier: Apache-2.0\n",
+    "#==========================================\n",
+    "\n",
+    "source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1\n",
+    "/bin/echo \"##\" $(whoami) is compiling oneDPL example\n",
+    "rm -rf bin/sort\n",
+    "dpcpp lab/sort.cpp -o bin/sort -tbb\n",
+    "bin/sort"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "Select the cell below and click Run ▶ to compile and execute the code above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 scripts/run_sort.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_sort.sh; else ./scripts/run_sort.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "***\n",
+    "# Fancy iterators\n",
+    "\n",
+    "Next we want to demonstrate how the applicability of classic STL algorithms can be extended by special iterators, so called fancy iterators. Imagine the following problem: A logic sequence of pairs (key and data elements) has to be sorted according to its key element. The elements are stored in two containers, one vector that holds the keys and another vector that holds the data elements.\n",
+    "\n",
+    "oneDPL as part of Intel oneAPI supports fancy iterators that can be used with the C++17 parallel algorithms. For example a `counting_iterator` is provided that represents a linear increasing sequence. With one big advantage compared to a classic vector that holds the data. This can actually save memory bandwidth. Because the sequence does not need a representation in memory.\n",
+    "\n",
+    "In our example we use a `zip_iterator` that can be used to tie two or even multiple sequences. The resulting iterator can now be used as an input or output of STL algorithm. Like in the following example.\n",
+    "\n",
+    "Complete the instructions in the following code template to implement a sort by key using a zip iterator."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/fancy_sort.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <iostream>\n",
+    "#include <vector>\n",
+    "\n",
+    "#include <oneapi/dpl/iterator>\n",
+    "#include <oneapi/dpl/execution>\n",
+    "#include <oneapi/dpl/algorithm>\n",
+    "\n",
+    "int main() {\n",
+    "\n",
+    "  {\n",
+    "    using oneapi::dpl::make_zip_iterator;\n",
+    "    std::vector<int> keys = { 0, 1, 0, 1, 0, 1, 0, 1, 0, 1};\n",
+    "    std::vector<int> data = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9};\n",
+    "    // *** Step 1: create a zip iterator with iterator to begin of keys and data as input.\n",
+    "    auto zip_in = data.begin();\n",
+    "    const size_t n = std::distance(keys.begin(), keys.end());  \n",
+    "\n",
+    "    std::cout << \"\\nRunning in parallel:\\n\";\n",
+    "    auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "    // *** Step 2: Replace left and right side of comparison with get<>() to extract and compare key\n",
+    "    auto custom_func = [](auto l, auto r){ using std::get; return l < r;};\n",
+    "    // *** Step 3: Replace all instances of iterator data.begin() with your zip iterator\n",
+    "    std::sort(std::execution::par, zip_in, zip_in + n, custom_func);\n",
+    "    auto pt1 = std::chrono::high_resolution_clock::now();\n",
+    "    std::cout << \"Parallel time = \" << 1e-9 * (pt1-pt0).count() << \" seconds\\n\";\n",
+    "      \n",
+    "    std::cout << \"Sorted keys:\" ;\n",
+    "    for(auto i:keys)\n",
+    "      std::cout << i << \" \";\n",
+    "    std::cout << \"\\n\" << \"Sorted data:\";\n",
+    "    for(auto i:data)\n",
+    "      std::cout << i << \" \";\n",
+    "    std::cout << \"\\n\";\n",
+    "  }\n",
+    "    \n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Solution (Don't peak unless you have to)"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "%%writefile lab/fancy_sort.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <iostream>\n",
+    "#include <vector>\n",
+    "\n",
+    "#include <oneapi/dpl/iterator>\n",
+    "#include <oneapi/dpl/execution>\n",
+    "#include <oneapi/dpl/algorithm>\n",
+    "\n",
+    "int main() {\n",
+    "\n",
+    "  {\n",
+    "    using oneapi::dpl::make_zip_iterator;\n",
+    "    std::vector<int> keys = { 0, 1, 0, 1, 0, 1, 0, 1, 0, 1};\n",
+    "    std::vector<int> data = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9};\n",
+    "    auto zip_in = make_zip_iterator(keys.begin(), data.begin());\n",
+    "    const size_t n = std::distance(keys.begin(), keys.end());  \n",
+    "    auto custom_func = [](auto l, auto r){ using std::get; return get<0>(l) < get<0>(r);};\n",
+    "      \n",
+    "    std::cout << \"\\nRunning in parallel:\\n\";\n",
+    "    auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "    std::sort(std::execution::par, zip_in, zip_in + n, custom_func);\n",
+    "    auto pt1 = std::chrono::high_resolution_clock::now();\n",
+    "    std::cout << \"Parallel time = \" << 1e-9 * (pt1-pt0).count() << \" seconds\\n\";\n",
+    "      \n",
+    "    std::cout << \"Sorted keys:\" ;\n",
+    "    for(auto i:keys)\n",
+    "      std::cout << i << \" \";\n",
+    "    std::cout << \"\\n\" << \"Sorted data:\";\n",
+    "    for(auto i:data)\n",
+    "      std::cout << i << \" \";\n",
+    "    std::cout << \"\\n\";\n",
+    "  }\n",
+    "    \n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile scripts/run_fancy_sort.sh\n",
+    "#!/bin/bash\n",
+    "#==========================================\n",
+    "# Copyright (c) 2020 Intel Corporation\n",
+    "#\n",
+    "# SPDX-License-Identifier: Apache-2.0\n",
+    "#==========================================\n",
+    "\n",
+    "source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1\n",
+    "/bin/echo \"##\" $(whoami) is compiling oneDPL example\n",
+    "rm -rf bin/fancy_sort\n",
+    "dpcpp lab/fancy_sort.cpp -o bin/fancy_sort -tbb\n",
+    "bin/fancy_sort"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "Select the cell below and click Run ▶ to compile and execute the code above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 scripts/run_fancy_sort.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_fancy_sort.sh; else ./scripts/run_fancy_sort.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "***\n",
+    "# Monte Carlo Pi\n",
+    "\n",
+    "Let's recreate the Direct programming DPC++ example using the DPC++ library (oneDPL). We'll use the algorithm ``transform_reduce``, which can for example be used to parallelize the computation of a dot product, by applying a parallel execution policy. Well or the monte carlo simulation computing PI.  "
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/monte_carlo_pi.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <iostream>\n",
+    "#include <vector>\n",
+    "\n",
+    "#include <CL/sycl.hpp>\n",
+    "\n",
+    "#include <oneapi/dpl/random>\n",
+    "\n",
+    "#include <oneapi/dpl/iterator>\n",
+    "#include <oneapi/dpl/execution>\n",
+    "#include <oneapi/dpl/algorithm>\n",
+    "\n",
+    "#define N 500\n",
+    "#define LOCAL_N 1000\n",
+    "\n",
+    "#define SEED 667\n",
+    "#define PI 3.1415926535897932384626433832795\n",
+    "\n",
+    "int main() {\n",
+    "    \n",
+    "    using oneapi::dpl::counting_iterator;\n",
+    "    double estimated_pi = 3.14;\n",
+    "    {\n",
+    "        sycl::queue q; //(sycl::gpu_selector{});\n",
+    "        auto policy = oneapi::dpl::execution::make_device_policy(q);\n",
+    " \n",
+    "        auto sum = transform_reduce( policy, counting_iterator<int>(0), counting_iterator<int>(N), 0, std::plus<float>{}\n",
+    "                                    , [=](auto n){\n",
+    "                                        float local_sum = 0.0f;\n",
+    "                                        // Get random coords\n",
+    "                                        oneapi::std::minstd_rand engine(SEED, n);\n",
+    "                                        oneapi::std::uniform_real_distribution<float> distr(-1.0f,1.0f);\n",
+    "                                        for(int i = 0; i < LOCAL_N; ++i) {\n",
+    "                                            float x = distr(engine);\n",
+    "                                            float y = distr(engine);\n",
+    "                                            auto hypotenuse_sqr = x * x + y * y;\n",
+    "                                            if (hypotenuse_sqr <= 1.0)\n",
+    "                                                local_sum += 1.0;\n",
+    "                                        }\n",
+    "                                        return local_sum / (float)LOCAL_N;\n",
+    "                                    });\n",
+    "\n",
+    "        estimated_pi = 4.0 * (float)sum / N;\n",
+    "    }\n",
+    "        \n",
+    "    // Printing Results\n",
+    "    std::cout << \"Estimated value of Pi = \" << estimated_pi << std::endl;\n",
+    "    std::cout << \"Exact value of Pi = \" << PI << std::endl;\n",
+    "    std::cout << \"Absolute error = \" << fabs( PI-estimated_pi ) << std::endl;\n",
+    "\n",
+    "    return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile scripts/run_monte_carlo_pi.sh\n",
+    "#!/bin/bash\n",
+    "#==========================================\n",
+    "# Copyright (c) 2020 Intel Corporation\n",
+    "#\n",
+    "# SPDX-License-Identifier: Apache-2.0\n",
+    "#==========================================\n",
+    "\n",
+    "source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1\n",
+    "/bin/echo \"##\" $(whoami) is compiling oneDPL example\n",
+    "rm -rf bin/monte_carlo_pi\n",
+    "dpcpp lab/monte_carlo_pi.cpp -o bin/monte_carlo_pi -tbb\n",
+    "bin/monte_carlo_pi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "Select the cell below and click Run ▶ to compile and execute the code above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 scripts/run_monte_carlo_pi.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_monte_carlo_pi.sh; else ./scripts/run_monte_carlo_pi.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### "
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3.7 (Intel® oneAPI)",
+   "language": "python",
+   "name": "c009-intel_distribution_of_python_3_oneapi-beta05-python"
+  },
+  "language_info": {
+   "codemirror_mode": {
+    "name": "ipython",
+    "version": 3
+   },
+   "file_extension": ".py",
+   "mimetype": "text/x-python",
+   "name": "python",
+   "nbconvert_exporter": "python",
+   "pygments_lexer": "ipython3",
+   "version": "3.7.7"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 4
+}
@@ -0,0 +1 @@
+bin directory needed for notebook
@@ -0,0 +1,149 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#pragma once
+
+#include <exception>
+#include <iostream>
+#include <numeric>
+#include <stdlib.h>
+
+#include <CL/sycl.hpp>
+
+namespace common {
+
+// this exception handler with catch async exceptions
+static auto exception_handler = [](cl::sycl::exception_list eList) {
+  for (std::exception_ptr const &e : eList) {
+    try {
+      std::rethrow_exception(e);
+    } catch (std::exception const &e) {
+#if _DEBUG
+      std::cout << "Failure" << std::endl;
+#endif
+      std::terminate();
+    }
+  }
+};
+
+template<typename ArrayType>
+void print_array(const char* text, const ArrayType& array) {
+    std::cout << text;
+  for (const auto& s : array) std::cout << s << ' ';
+  std::cout << "\n";
+}
+
+void print_array(const char* text, const float * array, int array_size) {
+    std::cout << text;
+  for (int i=0; i<array_size; ++i) std::cout << array[i] << ' ';
+  std::cout << "\n";
+}
+
+template<typename ArrayType>
+void init_input_arrays(ArrayType &a_array, ArrayType &b_array) {
+  std::iota(a_array.begin(), a_array.end(), 0);
+  std::iota(b_array.begin(), b_array.end(), 0);
+}
+
+template<typename ArrayType>
+void validate_results(float alpha, const ArrayType& a_array, const ArrayType& b_array,
+		      const ArrayType& c_sycl, const ArrayType& c_tbb) {
+  // Create gold standard / serial result
+  ArrayType c_gold;
+
+  for (size_t i = 0; i < a_array.size(); ++i)
+    c_gold[i] = a_array[i] + alpha * b_array[i];
+
+  // Compare golden triad with heterogeneous triad
+  if (!std::equal(std::begin(c_sycl), std::end(c_sycl), std::begin(c_gold)))
+    std::cout << "SYCL triad error.\n";
+  else
+    std::cout << "SYCL triad correct.\n";
+
+  // Compare golden triad with TBB triad
+  if (!std::equal(std::begin(c_tbb), end(c_tbb),
+                  std::begin(c_gold)))
+    std::cout << "TBB triad error.\n";
+  else
+    std::cout << "TBB triad correct.\n";
+}
+    
+template<typename ArrayType>
+void validate_hetero_results(float ratio, float alpha_sycl, float alpha_tbb, const ArrayType& a_array, 
+                             const ArrayType& b_array, const ArrayType& c_array) {
+
+  std::size_t ratio_switch_pt = static_cast<std::size_t>(std::ceil(c_array.size() * ratio));
+    
+  // Create gold standard / serial result
+  ArrayType c_gold;
+  for (size_t i = 0; i < ratio_switch_pt; ++i)
+    c_gold[i] = a_array[i] + alpha_sycl * b_array[i];
+  for (size_t i = ratio_switch_pt; i < c_array.size(); ++i)
+    c_gold[i] = a_array[i] + alpha_tbb * b_array[i];
+    
+  // Compare golden triad with heterogeneous triad
+  if (!std::equal(std::begin(c_array), std::end(c_array), std::begin(c_gold)))
+    std::cout << "Hetero triad error.\n";
+  else
+    std::cout << "Hetero triad correct.\n";
+}
+
+void validate_usm_results(float ratio, float alpha_sycl, float alpha_tbb, const float* a_array, 
+                             const float* b_array, const float* c_array, int array_size) {
+
+  std::size_t ratio_switch_pt = static_cast<std::size_t>(std::ceil(array_size * ratio));
+    
+  // Create gold standard / serial result
+  float* c_gold=new float[array_size];
+  for (size_t i = 0; i < ratio_switch_pt; ++i)
+    c_gold[i] = a_array[i] + alpha_sycl * b_array[i];
+  for (size_t i = ratio_switch_pt; i < array_size; ++i)
+    c_gold[i] = a_array[i] + alpha_tbb * b_array[i];
+    
+  // Compare golden triad with heterogeneous triad
+  if (!std::equal(c_array, c_array+array_size, c_gold))
+    std::cout << "Hetero triad error.\n";
+  else
+    std::cout << "Hetero triad correct.\n";
+  delete[] c_gold;
+}
+    
+
+template<typename ArrayType>
+void print_results(float alpha, const ArrayType& a_array, const ArrayType& b_array,
+                   const ArrayType& c_sycl, const ArrayType& c_tbb) {
+  print_array("input array a_array: ", a_array);
+  print_array("input array b_array: ", b_array);
+  std::cout << "alpha: " << alpha << "\n";
+  print_array("output array c using SYCL on GPU: ", c_sycl);
+  print_array("output array c using TBB on CPU: ", c_tbb);
+}
+    
+template<typename ArrayType>
+void print_hetero_results(float ratio, float alpha_sycl, float alpha_tbb, const ArrayType& a_array, 
+                          const ArrayType& b_array, const ArrayType& c_array) {
+  print_array("input array a_array: ", a_array);
+  print_array("input array b_array: ", b_array);
+  std::cout << "ratio: " << ratio << "\n";
+  std::cout << "switch_index: " << static_cast<std::size_t>(std::ceil(c_array.size() * ratio)) << "\n";
+  std::cout << "alpha_sycl: " << alpha_sycl << "\n";
+  std::cout << "alpha_tbb: " << alpha_tbb << "\n";
+  print_array("output array c_array: ", c_array);
+}
+
+void print_usm_results(float ratio, float alpha_sycl, float alpha_tbb, const float *a_array, 
+                          const float *b_array, const float *c_array, int array_size) {
+  print_array("input array a_array: ", a_array, array_size);
+  print_array("input array b_array: ", b_array, array_size);
+  std::cout << "ratio: " << ratio << "\n";
+  std::cout << "switch_index: " << static_cast<std::size_t>(std::ceil(array_size * ratio)) << "\n";
+  std::cout << "alpha_sycl: " << alpha_sycl << "\n";
+  std::cout << "alpha_tbb: " << alpha_tbb << "\n";
+  print_array("output array c_tbb: ", c_array, array_size);
+}
+    
+};  // namespace common
+
@@ -0,0 +1,17 @@
+//==============================================================
+// Copyright (c) 2019-2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#ifndef _GAMMA_UTILS_HPP
+#define _GAMMA_UTILS_HPP
+
+#include "utils/Img.hpp"
+#include "utils/ImgAlgorithm.hpp"
+#include "utils/ImgFormat.hpp"
+#include "utils/ImgPixel.hpp"
+
+#include "utils/Other.hpp"
+
+#endif  // _GAMMA_UTILS_HPP
@@ -0,0 +1,204 @@
+//==============================================================
+// Copyright (c) 2019-2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#ifndef _GAMMA_UTILS_IMG_HPP
+#define _GAMMA_UTILS_IMG_HPP
+
+#include "ImgPixel.hpp"
+
+#include <algorithm>
+#include <fstream>
+#include <iostream>
+#include <vector>
+
+// Image class definition
+template <typename Format>
+class Img {
+ private:
+  Format _format;
+  std::int32_t _width;
+  std::int32_t _height;
+  std::vector<ImgPixel> _pixels;
+
+  using Iterator = std::vector<ImgPixel>::iterator;
+  using ConstIterator = std::vector<ImgPixel>::const_iterator;
+
+ public:
+  /////////////////////
+  // SPECIAL METHODS //
+  /////////////////////
+
+  Img(std::int32_t width, std::int32_t height);
+
+  void reset(std::int32_t width, std::int32_t height);
+
+  ///////////////
+  // ITERATORS //
+  ///////////////
+
+  Iterator begin() noexcept;
+  Iterator end() noexcept;
+  ConstIterator begin() const noexcept;
+  ConstIterator end() const noexcept;
+  ConstIterator cbegin() const noexcept;
+  ConstIterator cend() const noexcept;
+
+  /////////////
+  // GETTERS //
+  /////////////
+
+  std::int32_t width() const noexcept;
+  std::int32_t height() const noexcept;
+
+  ImgPixel const* data() const noexcept;
+  ImgPixel* data() noexcept;
+
+  ///////////////////
+  // FUNCTIONALITY //
+  ///////////////////
+
+  void write(std::string const& filename) const;
+
+  template <typename Functor>
+  void fill(Functor f);
+  void fill(ImgPixel pixel);
+  void fill(ImgPixel pixel, std::int32_t row, std::int32_t col);
+};
+
+///////////////////////////////////////////////
+// IMG CLASS IMPLEMENTATION: SPECIAL METHODS //
+///////////////////////////////////////////////
+
+template <typename Format>
+Img<Format>::Img(std::int32_t width, std::int32_t height)
+    : _format(width, height) {
+  _pixels.resize(width * height);
+
+  _width = width;
+  _height = height;
+}
+
+template <typename Format>
+void Img<Format>::reset(std::int32_t width, std::int32_t height) {
+  _pixels.resize(width * height);
+
+  _width = width;
+  _height = height;
+
+  _format.reset(width, height);
+}
+
+/////////////////////////////////////////
+// IMG CLASS IMPLEMENTATION: ITERATORS //
+/////////////////////////////////////////
+
+template <typename Format>
+typename Img<Format>::Iterator Img<Format>::begin() noexcept {
+  return _pixels.begin();
+}
+
+template <typename Format>
+typename Img<Format>::Iterator Img<Format>::end() noexcept {
+  return _pixels.end();
+}
+
+template <typename Format>
+typename Img<Format>::ConstIterator Img<Format>::begin() const noexcept {
+  return _pixels.begin();
+}
+
+template <typename Format>
+typename Img<Format>::ConstIterator Img<Format>::end() const noexcept {
+  return _pixels.end();
+}
+
+template <typename Format>
+typename Img<Format>::ConstIterator Img<Format>::cbegin() const noexcept {
+  return _pixels.begin();
+}
+
+template <typename Format>
+typename Img<Format>::ConstIterator Img<Format>::cend() const noexcept {
+  return _pixels.end();
+}
+
+///////////////////////////////////////
+// IMG CLASS IMPLEMENTATION: GETTERS //
+///////////////////////////////////////
+
+template <typename Format>
+std::int32_t Img<Format>::width() const noexcept {
+  return _width;
+}
+
+template <typename Format>
+std::int32_t Img<Format>::height() const noexcept {
+  return _height;
+}
+
+template <typename Format>
+ImgPixel const* Img<Format>::data() const noexcept {
+  return _pixels.data();
+}
+
+template <typename Format>
+ImgPixel* Img<Format>::data() noexcept {
+  return _pixels.data();
+}
+
+/////////////////////////////////////////////
+// IMG CLASS IMPLEMENTATION: FUNCTIONALITY //
+/////////////////////////////////////////////
+
+template <typename Format>
+void Img<Format>::write(std::string const& filename) const {
+  if (_pixels.empty()) {
+    std::cerr << "Img::write:: image is empty" << std::endl;
+    return;
+  }
+
+  std::ofstream filestream(filename, std::ios::binary);
+
+  _format.write(filestream, *this);
+}
+
+template <typename Format>
+template <typename Functor>
+void Img<Format>::fill(Functor f) {
+  if (_pixels.empty()) {
+    std::cerr << "Img::fill(Functor): image is empty" << std::endl;
+    return;
+  }
+
+  for (auto& pixel : _pixels) f(pixel);
+}
+
+template <typename Format>
+void Img<Format>::fill(ImgPixel pixel) {
+  if (_pixels.empty()) {
+    std::cerr << "Img::fill(ImgPixel): image is empty" << std::endl;
+    return;
+  }
+
+  std::fill(_pixels.begin(), _pixels.end(), pixel);
+}
+
+template <typename Format>
+void Img<Format>::fill(ImgPixel pixel, int row, int col) {
+  if (_pixels.empty()) {
+    std::cerr << "Img::fill(ImgPixel): image is empty" << std::endl;
+    return;
+  }
+
+  if (row >= _height || row < 0 || col >= _width || col < 0) {
+    std::cerr << "Img::fill(ImgPixel, int, int): out of range" << std::endl;
+    return;
+  }
+
+  _pixels.at(row * _width + col) = pixel;
+}
+
+#endif  // _GAMMA_UTILS_IMG_HPP
@@ -0,0 +1,50 @@
+//==============================================================
+// Copyright (c) 2019-2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#ifndef _GAMMA_UTILS_IMGALGORITHM_HPP
+#define _GAMMA_UTILS_IMGALGORITHM_HPP
+
+#include <cmath>
+#include <cstdint>
+
+// struct to store fractal that image will fill from
+class ImgFractal {
+ private:
+  const std::int32_t _width;
+  const std::int32_t _height;
+
+  double _cx = -0.7436;
+  double _cy = 0.1319;
+
+  double _magn = 2000000.0;
+  int _maxIterations = 1000;
+
+ public:
+  ImgFractal(std::int32_t width, std::int32_t height)
+      : _width(width), _height(height) {}
+
+  double operator()(std::int32_t x, std::int32_t y) const {
+    double fx = (double(x) - double(_width) / 2) * (1 / _magn) + _cx;
+    double fy = (double(y) - double(_height) / 2) * (1 / _magn) + _cy;
+
+    double res = 0;
+    double nx = 0;
+    double ny = 0;
+    double val = 0;
+
+    for (int i = 0; nx * nx + ny * ny <= 4 && i < _maxIterations; ++i) {
+      val = nx * nx - ny * ny + fx;
+      ny = 2 * nx * ny + fy;
+      nx = val;
+
+      res += std::exp(-std::sqrt(nx * nx + ny * ny));
+    }
+
+    return res;
+  }
+};
+
+#endif  // _GAMMA_UTILS_IMGALGORITHM_HPP
@@ -0,0 +1,98 @@
+//==============================================================
+// Copyright (c) 2019-2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#ifndef _GAMMA_UTILS_IMGFORMAT_HPP
+#define _GAMMA_UTILS_IMGFORMAT_HPP
+
+#include "ImgPixel.hpp"
+
+#include <fstream>
+
+namespace ImgFormat {
+
+// struct to store an image in BMP format
+struct BMP {
+ private:
+  using FileHeader = struct {
+    // not from specification
+    // was added for alignment
+    // store size of rest of the fields
+    std::uint16_t sizeRest;  // file header size in bytes
+
+    std::uint16_t type;
+    std::uint32_t size;  // file size in bytes
+    std::uint32_t reserved;
+    std::uint32_t offBits;  // cumulative header size in bytes
+  };
+
+  using InfoHeader = struct {
+    // from specification
+    // store size of rest of the fields
+    std::uint32_t size;  // info header size in bytes
+
+    std::int32_t width;   // image width in pixels
+    std::int32_t height;  // image height in pixels
+    std::uint16_t planes;
+    std::uint16_t bitCount;      // color depth
+    std::uint32_t compression;   // compression
+    std::uint32_t sizeImage;     // image map size in bytes
+    std::int32_t xPelsPerMeter;  // pixel per metre (y axis)
+    std::int32_t yPelsPerMeter;  // pixel per metre (y axis)
+    std::uint32_t clrUsed;       // color palette (0 is default)
+    std::uint32_t clrImportant;
+  };
+
+  FileHeader _fileHeader;
+  InfoHeader _infoHeader;
+
+ public:
+  BMP(std::int32_t width, std::int32_t height) noexcept {
+    reset(width, height);
+  }
+
+  void reset(std::int32_t width, std::int32_t height) noexcept {
+    uint32_t padSize = (4 - (width * sizeof(ImgPixel)) % 4) % 4;
+    uint32_t mapSize = width * height * sizeof(ImgPixel) + height * padSize;
+    uint32_t allSize = mapSize + _fileHeader.sizeRest + _infoHeader.size;
+
+    _fileHeader.sizeRest = 14;  // file header size in bytes
+    _fileHeader.type = 0x4d42;
+    _fileHeader.size = allSize;  // file size in bytes
+    _fileHeader.reserved = 0;
+    _fileHeader.offBits = 54;  // sizeRest + size -> 14 + 40 -> 54
+
+    _infoHeader.size = 40;        // info header size in bytes
+    _infoHeader.width = width;    // image width in pixels
+    _infoHeader.height = height;  // image height in pixels
+    _infoHeader.planes = 1;
+    _infoHeader.bitCount = 32;        // color depth
+    _infoHeader.compression = 0;      // compression
+    _infoHeader.sizeImage = mapSize;  // image map size in bytes
+    _infoHeader.xPelsPerMeter = 0;    // pixel per metre (x axis)
+    _infoHeader.yPelsPerMeter = 0;    // pixel per metre (y axis)
+    _infoHeader.clrUsed = 0;          // color palette (0 is default)
+    _infoHeader.clrImportant = 0;
+  }
+
+  template <template <class> class Image, typename Format>
+  void write(std::ofstream& ostream, Image<Format> const& image) const {
+    ostream.write(reinterpret_cast<char const*>(&_fileHeader.type),
+                  _fileHeader.sizeRest);
+
+    ostream.write(reinterpret_cast<char const*>(&_infoHeader),
+                  _infoHeader.size);
+
+    ostream.write(reinterpret_cast<char const*>(image.data()),
+                  image.width() * image.height() * sizeof(image.data()[0]));
+  }
+
+  FileHeader const& fileHeader() const noexcept { return _fileHeader; }
+  InfoHeader const& infoHeader() const noexcept { return _infoHeader; }
+};
+
+}  // namespace ImgFormat
+
+#endif  // _GAMMA_UTILS_IMGFORMAT_HPP
@@ -0,0 +1,40 @@
+//==============================================================
+// Copyright (c) 2019-2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#ifndef _GAMMA_UTILS_IMGPIXEL_HPP
+#define _GAMMA_UTILS_IMGPIXEL_HPP
+
+#include <cstdint>
+#include <ostream>
+
+// struct to store a pixel of image
+struct ImgPixel {
+  std::uint8_t b;
+  std::uint8_t g;
+  std::uint8_t r;
+  std::uint8_t a;
+
+  bool operator==(ImgPixel const& other) const {
+    return (b == other.b) && (g == other.g) && (r == other.r) && (a == other.a);
+  }
+
+  bool operator!=(ImgPixel const& other) const { return !(*this == other); }
+
+  void set(std::uint8_t blue, std::uint8_t green, std::uint8_t red,
+           std::uint8_t alpha) {
+    b = blue;
+    g = green;
+    r = red;
+    a = alpha;
+  }
+};
+
+std::ostream& operator<<(std::ostream& output, ImgPixel const& pixel) {
+  return output << "(" << unsigned(pixel.r) << ", " << unsigned(pixel.g) << ", "
+                << unsigned(pixel.b) << ", " << unsigned(pixel.a) << ")";
+}
+
+#endif  // _GAMMA_UTILS_IMGPIXEL_HPP
@@ -0,0 +1,31 @@
+//==============================================================
+// Copyright (c) 2019-2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#ifndef _GAMMA_UTILS_OTHER_HPP
+#define _GAMMA_UTILS_OTHER_HPP
+
+#include <chrono>
+
+// function for time measuring
+inline double get_time_in_sec() {
+  namespace ch = std::chrono;
+  return ch::time_point_cast<ch::milliseconds>(ch::steady_clock::now())
+             .time_since_epoch()
+             .count() *
+         1.e-3;
+}
+
+// function to check correctness
+template <typename It>
+bool check(It begin1, It end1, It begin2) {
+  for (; begin1 != end1; ++begin1, ++begin2) {
+    if (*begin1 != *begin2) return false;
+  }
+
+  return true;
+}
+
+#endif  // _GAMMA_UTILS_OTHER_HPP
@@ -0,0 +1 @@
+lab directory needed for notebook
@@ -0,0 +1,55 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+# Script to submit job in Intel(R) DevCloud
+# Version: 0.7
+#==========================================
+if [ -z "$1" ]; then
+    echo "Missing script argument, Usage: ./q run.sh"
+elif [ ! -f "$1" ]; then
+    echo "File $1 does not exist"
+else
+    echo "Job has been submitted to Intel(R) DevCloud and will execute soon."
+    echo ""
+    echo " If you do not see result in 60 seconds, please restart the Jupyter kernel:"
+    echo " Kernel -> 'Restart Kernel and Clear All Outputs...' and then try again"
+    echo ""
+    script=$1
+    shift
+    args=$@
+    # Remove old output files
+    rm *.sh.* > /dev/null 2>&1
+    # Submit job using qsub
+    qsub_id=`qsub -l nodes=1:gpu:ppn=2 -d . $script $@`
+    job_id="$(cut -d'.' -f1 <<<"$qsub_id")"
+    # Print qstat output
+    qstat 
+    # Wait for output file to be generated and display
+    echo ""
+    echo -ne "Waiting for Output "
+    script_name=${script/scripts\//}
+    until [ -f $script_name.o$job_id ]; do
+        sleep 1
+        echo -ne "█"
+        ((timeout++))
+        # Timeout if no output file generated within 60 seconds
+        if [ $timeout == 60 ]; then
+            echo ""
+            echo ""
+            echo "TimeOut 60 seconds: Job is still queued for execution, check for output file later ($script_name.o$job_id)"
+            echo ""
+            break
+        fi
+    done
+    # Print output and error file content if exist
+    if [ -n "$(find -name '*.sh.o'$job_id)" ]; then
+        echo " Done⬇"
+        cat $script_name.o$job_id
+        cat $script_name.e$job_id
+        echo "Job Completed in $timeout seconds."
+        rm *.sh.*$job_id > /dev/null 2>&1
+    fi
+fi
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB with SYCL hetero exercise using flow::async_node
+dpcpp -std=c++20 -O2 solutions/triad-hetero-async_node-solved.cpp -o bin/triad-hetero-async_node-solved -tbb -lpthread
+if [ $? -eq 0 ]; then bin/triad-hetero-async_node-solved; fi
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB with SYCL hetero exercise using flow::async_node
+dpcpp -std=c++20 -O2 lab/triad-hetero-async_node.cpp -o bin/triad-hetero-async_node -tbb -lpthread
+if [ $? -eq 0 ]; then bin/triad-hetero-async_node; fi
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB with SYCL exercise triad consecutive
+dpcpp -std=c++20 -O2 lab/triad-consecutive.cpp -o bin/triad-consecutive -tbb -lpthread
+if [ $? -eq 0 ]; then bin/triad-consecutive; fi
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction exercise pi-parallel 
+dpcpp -std=c++17 -O2 lab/pi-parallel.cpp -o bin/pi-parallel -tbb
+bin/pi-parallel
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction exercise pi-serial 
+dpcpp -std=c++17 -O2 lab/pi-serial.cpp -o bin/pi-serial
+bin/pi-serial
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction sample pi-parallel 
+dpcpp -std=c++17 -O2 solutions/pi-parallel.cpp -o bin/pi-parallel-solution -tbb
+bin/pi-parallel-solution
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB with SYCL hetero solution using a resumable task
+dpcpp -std=c++20 -O2 solutions/triad-hetero-suspend-resume-solved.cpp -o bin/triad-hetero-suspend-resume-solved -tbb -lpthread
+if [ $? -eq 0 ]; then bin/triad-hetero-suspend-resume-solved; fi
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB with SYCL hetero exercise using a resumable task
+dpcpp -std=c++20 -O2 lab/triad-hetero-suspend-resume.cpp -o bin/triad-hetero-suspend-resume -tbb -lpthread
+if [ $? -eq 0 ]; then bin/triad-hetero-suspend-resume; fi
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB with SYCL solution using a task_group
+dpcpp -std=c++20 -O2 solutions/triad-task_group-solved.cpp -o bin/triad-task_group-solved -tbb -lpthread
+if [ $? -eq 0 ]; then bin/triad-task_group-solved; fi
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB with SYCL exercise using a task_group
+dpcpp -std=c++20 -O2 lab/triad-task_group.cpp -o bin/triad-task_group -ltbb -lpthread
+if [ $? -eq 0 ]; then bin/triad-task_group; fi
@@ -0,0 +1,62 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <CL/sycl.hpp>
+#define SIZE 1024
+
+using namespace sycl;
+
+int main() {
+    
+    std::array<int, SIZE> c;
+    // Reference vector for sanity check
+    std::array<int, SIZE> c_res;
+    std::fill_n(c_res.begin(), SIZE, 1);
+    
+    {
+        range<1> r{SIZE};
+        queue q(gpu_selector{});
+        buffer<int, 1> a_buf{r};
+        buffer<int, 1> b_buf{r};
+        buffer<int, 1> c_buf{c.data(), r};
+    
+    // Create a simple data dependency graph,
+    // Following structure:
+    // (1) ----> (3) ---> (4)
+    // (2) __/   
+    
+        q.submit([&](auto& h) {
+            accessor a(a_buf, h, write_only);
+            h.parallel_for(r, [=](auto idx) {
+                a[idx] = idx; }); 
+        });
+        q.submit([&](auto& h) {
+            accessor b(b_buf, h, write_only);
+            h.parallel_for(r, [=](auto idx) {
+                b[idx] = -idx; }); 
+        });
+        q.submit([&](auto& h) {
+            accessor a(a_buf, h, read_only);
+            accessor b(b_buf, h, read_only);
+            accessor c(c_buf, h, write_only);
+            h.parallel_for(r, [=](auto idx) {
+                c[idx] = a[idx] + b[idx]; }); 
+        });
+        q.submit([&](auto& h) {
+            accessor c(c_buf, h, read_write);
+            h.parallel_for(r, [=](auto idx) {
+                c[idx] += 1; }); 
+        }).wait();
+    }
+    
+    if (std::equal(c.begin(), c.end(), c_res.begin())) {
+        std::cout << "DONE.\n";
+    } else {
+        std::cout << "ERROR. Mismatch in resulting data!\n";
+    }
+    
+    return 0;
+}
@@ -0,0 +1,127 @@
+//==============================================================
+// Copyright (c) Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+#include <CL/sycl.hpp>
+#include <iomanip>
+#include <iostream>
+#include <vector>
+
+// dpc_common.hpp can be found in the dev-utilities include folder.
+// e.g., $ONEAPI_ROOT/dev-utilities/latest/include/dpc_common.hpp
+#include "dpc_common.hpp"
+
+// Header locations and some DPC++ extensions changed between beta09 and beta10
+// Temporarily modify the code sample to accept either version
+#define BETA09 20200827
+#if __SYCL_COMPILER_VERSION <= BETA09
+  #include <CL/sycl/intel/fpga_extensions.hpp>
+  namespace INTEL = sycl::intel;  // Namespace alias for backward compatibility
+#else
+  #include <CL/sycl/INTEL/fpga_extensions.hpp>
+#endif
+
+using namespace sycl;
+
+// This function instantiates the vector add kernel, which contains
+// a loop that adds up the two summand arrays and stores the result
+// into sum. 
+void VecAdd(const std::vector<float> &summands1,
+            const std::vector<float> &summands2, std::vector<float> &sum,
+            size_t array_size) {
+
+
+#if defined(FPGA_EMULATOR)
+  INTEL::fpga_emulator_selector device_selector;
+#else
+  INTEL::fpga_selector device_selector;
+#endif
+
+  try {
+    queue q(device_selector, dpc_common::exception_handler,
+            property::queue::enable_profiling{});
+
+    buffer buffer_summands1(summands1);
+    buffer buffer_summands2(summands2);
+    // Use verbose SYCL 1.2 syntax for the output buffer.
+    // (This will become unnecessary in a future compiler version.)
+    buffer<float, 1> buffer_sum(sum.data(), array_size);
+
+    event e = q.submit([&](handler &h) {
+      auto acc_summands1 = buffer_summands1.get_access<access::mode::read>(h);
+      auto acc_summands2 = buffer_summands2.get_access<access::mode::read>(h);
+      auto acc_sum = buffer_sum.get_access<access::mode::discard_write>(h);
+
+      h.single_task([=]() {
+        for (size_t i = 0; i < array_size; i++) {
+          acc_sum[i] = acc_summands1[i] + acc_summands2[i];
+        }
+      });
+    });
+
+    double start = e.get_profiling_info<info::event_profiling::command_start>();
+    double end = e.get_profiling_info<info::event_profiling::command_end>();
+    // convert from nanoseconds to ms
+    double kernel_time = (double)(end - start) * 1e-6;
+
+    std::cout << " kernel time : " << kernel_time << " ms\n";
+    std::cout << "Throughput for kernel " << ": ";
+    std::cout << std::fixed << std::setprecision(3)
+              << ((double)array_size / kernel_time) / 1e6f << " GFlops\n";
+
+  } catch (sycl::exception const &e) {
+    // Catches exceptions in the host code
+    std::cout << "Caught a SYCL host exception:\n" << e.what() << "\n";
+
+    // Most likely the runtime couldn't find FPGA hardware!
+    if (e.get_cl_code() == CL_DEVICE_NOT_FOUND) {
+      std::cout << "If you are targeting an FPGA, please ensure that your "
+                   "system has a correctly configured FPGA board.\n";
+      std::cout << "If you are targeting the FPGA emulator, compile with "
+                   "-DFPGA_EMULATOR.\n";
+    }
+    std::terminate();
+  }
+}
+
+int main(int argc, char *argv[]) {
+  size_t array_size = 1 << 26;
+
+  if (argc > 1) {
+    std::string option(argv[1]);
+    if (option == "-h" || option == "--help") {
+      std::cout << "Usage: \n<executable> <data size>\n\nFAILED\n";
+      return 1;
+    } else {
+      array_size = std::stoi(option);
+    }
+  }
+
+  std::vector<float> summands1(array_size);
+  std::vector<float> summands2(array_size);
+
+  std::vector<float> sum(array_size);
+
+  // Initialize the two summand arrays (arrays to be added to each other) to
+  // 1:N and N:1, so that the sum of all elements is N + 1
+  for (size_t i = 0; i < array_size; i++) {
+    summands1[i] = static_cast<float>(i + 1);
+    summands2[i] = static_cast<float>(array_size - i);
+  }
+
+  std::cout << "Input Array Size:  " << array_size << "\n";
+
+  // Instantiate VecAdd kernel that contains a loop that adds up the two summand arrays.
+  VecAdd(summands1, summands2, sum, array_size);
+
+  // Verify that the output data is the same for every unroll factor
+  for (size_t i = 0; i < array_size; i++) {
+    if (sum[i] != summands1[i] + summands2[i]) {
+      std::cout << "FAILED: The results are incorrect\n";
+      return 1;
+    }
+  }
+  std::cout << "PASSED: The results are correct\n";
+  return 0;
+}
@@ -0,0 +1,150 @@
+//==============================================================
+// Copyright (c) Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+#include <CL/sycl.hpp>
+#include <iomanip>
+#include <iostream>
+#include <vector>
+
+// dpc_common.hpp can be found in the dev-utilities include folder.
+// e.g., $ONEAPI_ROOT/dev-utilities//include/dpc_common.hpp
+#include "dpc_common.hpp"
+
+// Header locations and some DPC++ extensions changed between beta09 and beta10
+// Temporarily modify the code sample to accept either version
+#define BETA09 20200827
+#if __SYCL_COMPILER_VERSION <= BETA09
+  #include <CL/sycl/intel/fpga_extensions.hpp>
+  namespace INTEL = sycl::intel;  // Namespace alias for backward compatibility
+#else
+  #include <CL/sycl/INTEL/fpga_extensions.hpp>
+#endif
+
+using namespace sycl;
+
+template <int unroll_factor> class VAdd;
+
+// This function instantiates the vector add kernel, which contains
+// a loop that adds up the two summand arrays and stores the result
+// into sum. This loop will be unrolled by the specified unroll_factor.
+template <int unroll_factor>
+void VecAdd(const std::vector<float> &summands1,
+            const std::vector<float> &summands2, std::vector<float> &sum,
+            size_t array_size) {
+
+
+#if defined(FPGA_EMULATOR)
+  INTEL::fpga_emulator_selector device_selector;
+#else
+  INTEL::fpga_selector device_selector;
+#endif
+
+  try {
+    queue q(device_selector, dpc_common::exception_handler,
+            property::queue::enable_profiling{});
+
+    buffer buffer_summands1(summands1);
+    buffer buffer_summands2(summands2);
+    // Use verbose SYCL 1.2 syntax for the output buffer.
+    // (This will become unnecessary in a future compiler version.)
+    buffer<float, 1> buffer_sum(sum.data(), array_size);
+
+    event e = q.submit([&](handler &h) {
+      auto acc_summands1 = buffer_summands1.get_access<access::mode::read>(h);
+      auto acc_summands2 = buffer_summands2.get_access<access::mode::read>(h);
+      auto acc_sum = buffer_sum.get_access<access::mode::discard_write>(h);
+
+      h.single_task<VAdd<unroll_factor>>([=]()
+                                         [[intel::kernel_args_restrict]] {
+        // Unroll the loop fully or partially, depending on unroll_factor
+        #pragma unroll unroll_factor
+        for (size_t i = 0; i < array_size; i++) {
+          acc_sum[i] = acc_summands1[i] + acc_summands2[i];
+        }
+      });
+    });
+
+    double start = e.get_profiling_info<info::event_profiling::command_start>();
+    double end = e.get_profiling_info<info::event_profiling::command_end>();
+    // convert from nanoseconds to ms
+    double kernel_time = (double)(end - start) * 1e-6;
+
+    std::cout << "unroll_factor " << unroll_factor
+              << " kernel time : " << kernel_time << " ms\n";
+    std::cout << "Throughput for kernel with unroll_factor " << unroll_factor
+              << ": ";
+    std::cout << std::fixed << std::setprecision(3)
+              << ((double)array_size / kernel_time) / 1e6f << " GFlops\n";
+
+  } catch (sycl::exception const &e) {
+    // Catches exceptions in the host code
+    std::cout << "Caught a SYCL host exception:\n" << e.what() << "\n";
+
+    // Most likely the runtime couldn't find FPGA hardware!
+    if (e.get_cl_code() == CL_DEVICE_NOT_FOUND) {
+      std::cout << "If you are targeting an FPGA, please ensure that your "
+                   "system has a correctly configured FPGA board.\n";
+      std::cout << "If you are targeting the FPGA emulator, compile with "
+                   "-DFPGA_EMULATOR.\n";
+    }
+    std::terminate();
+  }
+}
+
+int main(int argc, char *argv[]) {
+  size_t array_size = 1 << 26;
+
+  if (argc > 1) {
+    std::string option(argv[1]);
+    if (option == "-h" || option == "--help") {
+      std::cout << "Usage: \n<executable> <data size>\n\nFAILED\n";
+      return 1;
+    } else {
+      array_size = std::stoi(option);
+    }
+  }
+
+  std::vector<float> summands1(array_size);
+  std::vector<float> summands2(array_size);
+
+  std::vector<float> sum_unrollx1(array_size);
+  std::vector<float> sum_unrollx2(array_size);
+  std::vector<float> sum_unrollx4(array_size);
+  std::vector<float> sum_unrollx8(array_size);
+  std::vector<float> sum_unrollx16(array_size);
+
+  // Initialize the two summand arrays (arrays to be added to each other) to
+  // 1:N and N:1, so that the sum of all elements is N + 1
+  for (size_t i = 0; i < array_size; i++) {
+    summands1[i] = static_cast<float>(i + 1);
+    summands2[i] = static_cast<float>(array_size - i);
+  }
+
+  std::cout << "Input Array Size:  " << array_size << "\n";
+
+  // Instantiate VecAdd kernel with different unroll factors: 1, 2, 4, 8, 16
+  // The VecAdd kernel contains a loop that adds up the two summand arrays.
+  // This loop will be unrolled by the specified unroll factor.
+  // The sum array is expected to be identical, regardless of the unroll factor.
+  VecAdd<1>(summands1, summands2, sum_unrollx1, array_size);
+  VecAdd<2>(summands1, summands2, sum_unrollx2, array_size);
+  VecAdd<4>(summands1, summands2, sum_unrollx4, array_size);
+  VecAdd<8>(summands1, summands2, sum_unrollx8, array_size);
+  VecAdd<16>(summands1, summands2, sum_unrollx16, array_size);
+
+  // Verify that the output data is the same for every unroll factor
+  for (size_t i = 0; i < array_size; i++) {
+    if (sum_unrollx1[i] != summands1[i] + summands2[i] ||
+        sum_unrollx1[i] != sum_unrollx2[i] ||
+        sum_unrollx1[i] != sum_unrollx4[i] ||
+        sum_unrollx1[i] != sum_unrollx8[i] ||
+        sum_unrollx1[i] != sum_unrollx16[i]) {
+      std::cout << "FAILED: The results are incorrect\n";
+      return 1;
+    }
+  }
+  std::cout << "PASSED: The results are correct\n";
+  return 0;
+}
@@ -0,0 +1,55 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <chrono>
+#include <cmath>
+#include <iostream>
+#include <limits>
+#include <thread>
+#include <tbb/tbb.h>
+
+double calc_pi(int num_intervals) {
+  double dx = 1.0 / num_intervals;
+  double sum = tbb::parallel_reduce(
+    /* range = */ tbb::blocked_range<int>(0, num_intervals ), 
+    /* identity = */ 0.0,
+    /* func */ 
+    [=](const tbb::blocked_range<int>& r, double init) -> double {
+      for (int i = r.begin(); i != r.end(); ++i) {
+        double x = (i+0.5)*dx;
+        double h = std::sqrt(1-x*x);
+        init += h*dx;
+      }
+      return init;
+    },
+    std::plus<double>{}
+  );
+  double pi = 4 * sum;
+  return pi;
+}
+
+static void warmupTBB() {
+  int num_threads = std::thread::hardware_concurrency();
+  tbb::parallel_for(0, num_threads,
+    [](unsigned int) { 
+      std::this_thread::sleep_for(std::chrono::milliseconds(10)); 
+  });
+}
+
+int main() {
+  const int num_intervals = std::numeric_limits<int>::max();
+  double parallel_time = 0.0;
+  warmupTBB();
+  {
+    auto pt0 = std::chrono::high_resolution_clock::now();
+    double pi = calc_pi(num_intervals);
+    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();
+    std::cout << "parallel pi == " << pi << std::endl;
+  }
+
+  std::cout << "parallel_time == " << parallel_time << " seconds" << std::endl;
+  return 0;
+}
@@ -0,0 +1,67 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <CL/sycl.hpp>
+// For random number generation in device code:
+#include <oneapi/dpl/random>
+// Utility header for output Image
+#include "utils.hpp"
+
+using namespace sycl;
+
+const int iter = 10000;
+const int img_dimensions = 512;
+
+const int size_wg = 32;
+const int num_wg = (int)ceil((double)iter / (double)size_wg);
+
+// Return index of pixel corresponding to a set of simulation coordinates:
+int GetPixelIndex(double x, double y) {
+    const int radius = img_dimensions / 2;
+    int img_x = x * radius + radius;
+    int img_y = y * radius + radius;
+    return img_y * img_dimensions + img_x;
+}
+
+int main() {
+    Img<ImgFormat::BMP> image{img_dimensions, img_dimensions};
+    int sum = 0;
+    
+    queue q(default_selector{});
+    
+    std::uint32_t seed = 667;
+    {
+        buffer sum_buf(&sum, range(1));
+        buffer<ImgPixel,1> imgplot_buf(image.data(), image.width() * image.height());
+
+        queue q;
+        // Compute a dot-product by reducing all computed values using standard plus functor
+        q.submit([&](handler& cgh) {
+            accessor img_plot(imgplot_buf, cgh, read_write);
+            accessor sum(sum_buf, cgh, read_write);
+            cgh.parallel_for<class dot_product>(nd_range<1>{num_wg * size_wg, size_wg}, sycl::ONEAPI::reduction(sum, 0, std::plus<int>()), [=](nd_item<1> it, auto& sum) {
+                int i = it.get_global_id(0);
+                
+                if (i < iter) {
+                    // Get random coords
+                    oneapi::std::minstd_rand engine(seed, i * 2);
+                    oneapi::std::uniform_real_distribution<double> distr(-1.0,1.0);
+                    double x = distr(engine);
+                    double y = distr(engine);
+                    
+                    auto hypotenuse_sqr = x * x + y * y;
+                    if (hypotenuse_sqr <= 1.0) {
+                        sum += 1;
+                        img_plot[GetPixelIndex(x,y)].set(255,255,255,255);
+                    }
+                }
+            });
+        }).wait();   
+    }
+    std::cout << "DONE. \nPi: " << 4.0 * (double)sum / iter << "\n";
+    image.write("img/pi.png");
+    return 0;
+}
@@ -0,0 +1,153 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache 2.0
+// =============================================================
+
+#include "../common/common_utils.hpp"
+
+#include <cmath>  //for std::ceil
+#include <array>
+#include <atomic>
+#include <iostream>
+#include <thread>
+
+#include <CL/sycl.hpp>
+
+#include <tbb/blocked_range.h>
+#include <tbb/flow_graph.h>
+#include <tbb/global_control.h>
+#include <tbb/parallel_for.h>
+
+constexpr size_t array_size = 16;
+
+template<size_t ARRAY_SIZE>
+struct msg_t {
+  static constexpr size_t array_size = ARRAY_SIZE;
+  const float offload_ratio = 0.5;
+  const float alpha_0 = 0.5;
+  const float alpha_1 = 1.0;
+  std::array<float, array_size> a_array;  // input
+  std::array<float, array_size> b_array;  // input
+  std::array<float, array_size> c_sycl;   // GPU output
+  std::array<float, array_size> c_tbb;    // CPU output
+};
+
+using msg_ptr = std::shared_ptr<msg_t<array_size>>;
+
+using async_node_t = tbb::flow::async_node<msg_ptr, msg_ptr>;
+using gateway_t = async_node_t::gateway_type;
+
+class AsyncActivity {
+  msg_ptr msg;
+  gateway_t* gateway_ptr;
+  std::atomic<bool> submit_flag;
+  std::thread service_thread;
+
+ public:
+  AsyncActivity() : msg{nullptr}, gateway_ptr{nullptr}, submit_flag{false},
+    service_thread( [this] {
+      //Wait until other thread sets submit_flag=true
+      while( !submit_flag ) std::this_thread::yield();
+      // Here we go! Dispatch code to the GPU
+      // Execute the kernel over a portion of the array range
+      size_t array_size_sycl = std::ceil(msg->a_array.size() * msg->offload_ratio);
+      {  
+        sycl::buffer a_buffer{msg->a_array}, b_buffer{msg->b_array}, c_buffer{msg->c_sycl};
+        sycl::queue q{sycl::gpu_selector{}};
+        float alpha = msg->alpha_0;
+        q.submit([&, alpha](sycl::handler& h) {            
+          sycl::accessor a_accessor{a_buffer, h, sycl::read_only};
+          sycl::accessor b_accessor{b_buffer, h, sycl::read_only};
+          sycl::accessor c_accessor{c_buffer, h, sycl::write_only};
+          h.parallel_for(sycl::range<1>{array_size_sycl}, [=](sycl::id<1> index) {
+            c_accessor[index] = a_accessor[index] + b_accessor[index] * alpha;
+          });  
+        }).wait();
+      }
+      gateway_ptr->try_put(msg);
+      gateway_ptr->release_wait();
+    } ) {}
+
+  ~AsyncActivity() {
+    service_thread.join();
+  }
+
+  void submit(msg_ptr m, gateway_t& gateway) {
+    gateway.reserve_wait();
+    msg = m;
+    gateway_ptr = &gateway;
+    submit_flag = true;
+  }
+};
+
+int main() {
+  tbb::flow::graph g;
+
+  // Input node:
+  tbb::flow::input_node<msg_ptr> in_node{g, 
+    [&](tbb::flow_control& fc) -> msg_ptr {
+      static bool has_run = false;
+      if (has_run) fc.stop();
+      has_run = true; // This example only creates a message to feed the Flow Graph
+      msg_ptr msg = std::make_shared<msg_t<array_size>>();
+      common::init_input_arrays(msg->a_array, msg->b_array);
+      return msg;
+    }
+  };
+
+  // CPU node
+  tbb::flow::function_node<msg_ptr, msg_ptr> cpu_node{
+      g, tbb::flow::unlimited, [&](msg_ptr msg) -> msg_ptr {
+        size_t i_start = static_cast<size_t>(std::ceil(msg->array_size * msg->offload_ratio));
+        size_t i_end = static_cast<size_t>(msg->array_size);
+        auto &a_array = msg->a_array, &b_array = msg->b_array, &c_tbb = msg->c_tbb;
+        float alpha = msg->alpha_1;
+        tbb::parallel_for(tbb::blocked_range<size_t>{i_start, i_end},
+          [&, alpha](const tbb::blocked_range<size_t>& r) {
+            for (size_t i = r.begin(); i < r.end(); ++i)
+              c_tbb[i] = a_array[i] + alpha * b_array[i];
+            }
+        );
+        return msg;
+      }};
+
+  // async node -- GPU
+  AsyncActivity async_act;
+  async_node_t a_node{g, tbb::flow::unlimited,
+    [&async_act](msg_ptr msg, gateway_t& gateway) {
+      async_act.submit(msg, gateway);
+    }
+  };
+
+  // join node
+  using join_t = tbb::flow::join_node<std::tuple<msg_ptr, msg_ptr>>;
+  join_t node_join{g};
+
+  // out node
+  tbb::flow::function_node<join_t::output_type> out_node{g, tbb::flow::unlimited, 
+    [&](const join_t::output_type& two_msgs) {
+      msg_ptr msg = std::get<0>(two_msgs); //Both msg's point to the same data
+      //Merge GPU result into CPU array
+      std::size_t gpu_end = static_cast<std::size_t>(std::ceil(msg->array_size * msg->offload_ratio));
+      std::copy(msg->c_sycl.begin(), msg->c_sycl.begin()+gpu_end, msg->c_tbb.begin());
+      common::validate_hetero_results(msg->offload_ratio, msg->alpha_0, msg->alpha_1, 
+                                      msg->a_array, msg->b_array, msg->c_tbb);
+      if(msg->array_size<=64)
+        common::print_hetero_results(msg->offload_ratio, msg->alpha_0, msg->alpha_1, 
+                                     msg->a_array, msg->b_array, msg->c_tbb);
+    }
+  };  // end of out node
+
+  // construct graph
+  tbb::flow::make_edge(in_node, a_node);
+  tbb::flow::make_edge(in_node, cpu_node);
+  tbb::flow::make_edge(a_node, tbb::flow::input_port<0>(node_join));
+  tbb::flow::make_edge(cpu_node, tbb::flow::input_port<1>(node_join));
+  tbb::flow::make_edge(node_join, out_node);
+
+  in_node.activate();
+  g.wait_for_all();
+
+  return 0;
+}
@@ -0,0 +1,119 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache 2.0
+// =============================================================
+
+#include "../common/common_utils.hpp"
+
+#include <array>
+#include <atomic>
+#include <cmath>
+#include <iostream>
+#include <thread>
+#include <algorithm>
+
+#include <CL/sycl.hpp>
+
+#include <tbb/blocked_range.h>
+#include <tbb/task.h>
+#include <tbb/task_group.h>
+#include <tbb/parallel_for.h>
+
+template<size_t array_size>
+class AsyncActivity {
+    float alpha;
+    const float *a_array, *b_array; 
+    float *c_sycl;
+    sycl::queue& q;
+    float offload_ratio;
+    std::atomic<bool> submit_flag;
+    tbb::task::suspend_point suspend_point;
+    std::thread service_thread;
+
+public:
+    AsyncActivity(float alpha_sycl, const float *a, const float *b, float *c, sycl::queue &queue) : 
+      alpha{alpha_sycl}, a_array{a}, b_array{b}, c_sycl{c}, q{queue}, 
+      offload_ratio{0}, submit_flag{false},
+      service_thread([this] {
+        // Wait until the job will be submitted into the async activity
+        while(!submit_flag) std::this_thread::yield();
+        // Here submit_flag==true --> DISPATCH GPU computation
+        std::size_t array_size_sycl = std::ceil(array_size * offload_ratio);
+        float l_alpha=alpha;
+        const float *la=a_array, *lb=b_array;
+        float *lc=c_sycl;
+        q.submit([&](sycl::handler& h) {            
+            h.parallel_for(sycl::range<1>{array_size_sycl}, [=](sycl::id<1> index) {
+              lc[index] = la[index] + lb[index] * l_alpha;
+            });  
+        }).wait(); //The thread may spin or block here.
+  
+        // Pass a signal into the main thread that the GPU work is completed
+        tbb::task::resume(suspend_point);
+      }) {}
+
+    ~AsyncActivity() {
+        service_thread.join();
+    }
+
+    void submit( float ratio, tbb::task::suspend_point sus_point ) {
+        offload_ratio = ratio;
+        suspend_point = sus_point;
+        submit_flag = true;
+    }
+}; // class AsyncActivity
+
+int main() {
+  
+  constexpr float ratio = 0.5; // CPU or GPU offload ratio
+  // We use different alpha coefficients so that 
+  //we can identify the GPU and CPU part if we print c_array result
+  const float alpha_sycl = 0.5, alpha_tbb = 1.0;  
+  constexpr size_t array_size = 16;
+
+  sycl::queue q{sycl::gpu_selector{}};
+  std::cout << "Using device: " << q.get_device().get_info<sycl::info::device::name>() << '\n'; 
+
+  //This host allocation of c comes handy specially for integrated GPUs (CPU and GPU share mem)
+  float *a_array = malloc_host<float>(array_size, q); 
+  float *b_array = malloc_host<float>(array_size, q); 
+  float *c_sycl = malloc_host<float>(array_size, q);
+  float *c_tbb = new float[array_size];
+
+  // sets array values to 0..N
+  std::iota(a_array, a_array+array_size,0); 
+  std::iota(b_array, b_array+array_size,0);
+  
+  tbb::task_group tg;
+  AsyncActivity<array_size> activity{alpha_sycl, a_array, b_array, c_sycl, q};
+
+  //Spawn a task that runs a parallel_for on the CPU
+  tg.run([&, alpha_tbb]{
+   std::size_t i_start = static_cast<std::size_t>(std::ceil(array_size * ratio));
+   std::size_t i_end = array_size;
+   tbb::parallel_for(i_start, i_end, [=]( std::size_t index ) {
+     c_tbb[index] = a_array[index] + alpha_tbb * b_array[index];
+   });
+  });
+
+  //Spawn another task that asyncrhonously offloads computation to the GPU  
+    tbb::task::suspend([&]( tbb::task::suspend_point suspend_point ) {
+     activity.submit(ratio, suspend_point);
+    });
+
+  tg.wait();
+
+  //Merge GPU result into CPU array
+  std::size_t gpu_end = static_cast<std::size_t>(std::ceil(array_size * ratio));
+  std::copy(c_sycl, c_sycl+gpu_end, c_tbb);
+
+  common::validate_usm_results(ratio, alpha_sycl, alpha_tbb, a_array, b_array, c_tbb, array_size);
+  if(array_size<64)
+    common::print_usm_results(ratio, alpha_sycl, alpha_tbb, a_array, b_array, c_tbb, array_size);
+
+  free(a_array,q);
+  free(b_array,q);
+  free(c_sycl,q);
+  delete[] c_tbb;
+}
@@ -0,0 +1,59 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache 2.0
+// =============================================================
+
+#include "../common/common_utils.hpp"
+#include <array>
+#include <tbb/blocked_range.h>
+#include <tbb/parallel_for.h>
+#include "tbb/task_group.h"
+
+int main() {
+  const float alpha = 0.5;  // coeff for triad calculation
+  const size_t array_size = 16;
+
+  std::array<float, array_size> a_array, b_array, c_sycl, c_tbb;
+
+  // sets array values to 0..N
+  common::init_input_arrays(a_array, b_array); 
+
+  // create task_group
+  tbb::task_group tg;
+
+  // Run a TBB task that uses SYCL to offload to GPU, function run does not block
+  tg.run([&, alpha]() {
+    std::cout << "executing on the GPU using SYCL\n";
+    {  
+      sycl::buffer a_buffer{a_array}, b_buffer{b_array}, c_buffer{c_sycl};
+      sycl::queue q{sycl::default_selector{}};
+      q.submit([&](sycl::handler& h) {            
+        sycl::accessor a_accessor{a_buffer, h, sycl::read_only};
+        sycl::accessor b_accessor{b_buffer, h, sycl::read_only};
+        sycl::accessor c_accessor{c_buffer, h, sycl::write_only};
+        h.parallel_for(sycl::range<1>{array_size}, [=](sycl::id<1> index) {
+           c_accessor[index] = a_accessor[index] + b_accessor[index] * alpha;
+        });  
+      }).wait();
+    }
+  });
+
+  // Run a TBB task that uses SYCL to offload to CPU
+  tg.run([&, alpha]() {
+    std::cout << "executing on the CPU using TBB\n";
+
+    tbb::parallel_for(tbb::blocked_range<int>(0, a_array.size()),
+      [&](tbb::blocked_range<int> r) {
+        for (int index = r.begin(); index < r.end(); ++index) {
+          c_tbb[index] = a_array[index] + b_array[index] * alpha;
+        }
+    });
+  });
+
+  // wait for both TBB tasks to complete
+  tg.wait();
+
+  common::validate_results(alpha, a_array, b_array, c_sycl, c_tbb);
+  common::print_results(alpha, a_array, b_array, c_sycl, c_tbb);
+} 
@@ -0,0 +1,50 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <CL/sycl.hpp>
+
+#define N 1024
+
+using namespace sycl;
+
+int main() {
+    
+    std::vector<int> c_host(N);
+    std::vector<int> c_res(N);
+    std::fill_n(c_res.begin(), N, 1);
+    
+    {
+        queue q(gpu_selector{});
+        
+        int *a = malloc_shared<int>(N, q);
+        int *b = malloc_shared<int>(N, q);
+        int *c_dev = malloc_device<int>(N, q);
+        
+        for (int i = 0; i < N; ++i) {
+            a[i] = i + 1;
+            b[i] = -i;
+        }
+        
+        auto e1 = q.parallel_for(range<1>(N), [=](id<1> i) { c_dev[i] = a[i] + b[i]; });
+                 
+        q.submit([&](auto &h) {
+              h.depends_on(e1);
+              h.memcpy(c_host.data(), c_dev, sizeof(int) * N);
+        }).wait();
+                 
+        free(a,q);
+        free(b,q);
+        free(c_dev,q);
+    }
+
+    if (std::equal(c_host.begin(), c_host.end(), c_res.begin())) {
+        std::cout << "DONE.\n";
+    } else {
+        std::cout << "ERROR. Mismatch in resulting data!\n";
+    }
+    
+    return 0;
+}
@@ -0,0 +1,58 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <CL/sycl.hpp>
+#include <CL/sycl/INTEL/fpga_extensions.hpp>
+
+#define SIZE 1024
+
+using namespace sycl;
+
+int main() {
+    
+    {
+        range<1> r{SIZE};
+        #ifdef FPGA_EMULATOR
+        INTEL::fpga_emulator_selector device_selector;
+        #else
+        INTEL::fpga_selector device_selector;
+        #endif
+        queue q(device_selector);
+        
+        buffer<int, 1> a_buf{r};
+        buffer<int, 1> b_buf{r};
+        buffer<int, 1> c_buf{r};
+    
+    // a ---- c --- d
+    // b __/   
+    
+        q.submit([&](handler& h) {
+            accessor a(a_buf, h, write_only);
+            h.parallel_for(r, [=](auto idx) {
+                a[idx] = idx; }); 
+        });
+        q.submit([&](handler& h) {
+            accessor b(b_buf, h, write_only);
+            h.parallel_for(r, [=](auto idx) {
+                b[idx] = -idx; }); 
+        });
+        q.submit([&](handler& h) {
+            accessor a(a_buf, h, read_only);
+            accessor b(b_buf, h, read_only);
+            accessor c(c_buf, h, write_only);
+            h.parallel_for(r, [=](auto idx) {
+                c[idx] = a[idx] + b[idx]; }); 
+        });
+        q.submit([&](handler& h) {
+            accessor c(c_buf, h, read_write);
+            h.parallel_for(r, [=](auto idx) {
+                c[idx] += 1; }); 
+        }).wait();
+    }
+    std::cout << "DONE.\n";
+    
+    return 0;
+}
@@ -0,0 +1,58 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <CL/sycl.hpp>
+#define SIZE 1024
+
+using namespace cl::sycl;
+
+int main() {
+  // Initialize container.
+  std::array<int, SIZE> a, b, c;
+  for (int i = 0; i<SIZE; ++i) {
+    a[i] = i + 1;
+    b[i] = -i;
+    c[i] = 0;
+  }
+  {
+    std::cout << "\nAvailable devices:\n";
+    // Report available devices.
+    auto platforms = platform::get_platforms();
+    for (auto &platform : platforms) {
+      std::cout << "Platform: "
+        << platform.get_info<info::platform::name>() << std::endl;
+      auto devices = platform.get_devices();
+      for (auto &device : devices ) {
+        std::cout << " Device: "
+          << device.get_info<info::device::name>()
+          << std::endl;
+      }
+    }
+    
+    range<1> r{SIZE};
+    gpu_selector selector;
+      
+    queue q(selector);
+    
+    // Print out the device information.
+    std::cout << "\nRunning on device: "
+              << q.get_device().get_info<info::device::name>() << "\n";
+      
+    buffer<int, 1> a_buf(a.data(), r);
+    buffer<int, 1> b_buf(b.data(), r);
+    buffer<int, 1> c_buf(c.data(), r);
+    q.submit([&](auto &h) {
+      accessor a(a_buf, h, read_only);
+      accessor b(b_buf, h, read_only);
+      accessor c(c_buf, h, write_only);
+      h.parallel_for(r,[=](id<1> idx) {
+        c[idx] = a[idx] + b[idx];
+      });
+    }).wait();
+  }
+  std::cout << "DONE.\n";
+  return 0;
+}
@@ -0,0 +1,36 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <CL/sycl.hpp>
+#define SIZE 1024
+
+using namespace cl::sycl;
+
+int main() {
+  std::array<int, SIZE> a, b, c;
+  for (int i = 0; i<SIZE; ++i) {
+    a[i] = i + 1;
+    b[i] = -i;
+    c[i] = 0;
+  }
+  {
+    range<1> r{SIZE};
+    queue q;
+    buffer<int, 1> a_buf(a.data(), r);
+    buffer<int, 1> b_buf(b.data(), r);
+    buffer<int, 1> c_buf(c.data(), r);
+    q.submit([&](handler &h) {
+      accessor a(a_buf, h, read_only);
+      accessor b(b_buf, h, read_only);
+      accessor c(c_buf, h, write_only);
+      h.parallel_for(r,[=](id<1> idx) {
+        c[idx] = a[idx] + b[idx];
+      });
+    }).wait();
+  }
+  std::cout << "DONE.\n";
+  return 0;
+}
@@ -0,0 +1,300 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# Hello, TBB!"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "##### Sections\n",
+    "- _Code_: [Hello, TBB!](#Hello,-TBB!)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "This set of Jupyter Notebooks contains a series of exercises that provide a taste of the features of the oneTBB library. \n",
+    "To learn more about oneTBB, you can visit [oneTBB on Github](https://github.com/oneapi-src/oneTBB) or explore the Apress book \n",
+    "[Pro TBB: C++ Parallel Programming with Threading Building Blocks (TBB)](https://www.apress.com/gp/book/9781484243978).\n",
+    "\n",
+    "This module will familiarize you with the basic structure and flow of the tutorials, and allow you do a practice\n",
+    "build and run of a simple oneTBB application.\n",
+    "\n",
+    "## Learning Objectives\n",
+    "\n",
+    "* To become familiar with the structure and general flow of the tutorials\n",
+    "* Run a first example that uses oneTBB to execute two tasks in parallel.\n",
+    "\n",
+    "## General Flow of the Tutorials\n",
+    "\n",
+    "Typically each exercise in the oneTBB modules provides three sections: (1) a problem statement with\n",
+    "a working serial implementation, (2) the skeleton of a parallel implementation that describes steps\n",
+    "you need to perform to complete the exercise and (3) a completed solution. To get the most out of each\n",
+    "exercise, you should not look at the completed solution unless you really get stuck.\n",
+    "\n",
+    "The remainder of this module provides a simple example that will follows the flow described\n",
+    "above and results in a simple TBB parallel application.\n",
+    "\n",
+    "## Hello, TBB!\n",
+    "\n",
+    "Let's start by creating an application that invokes a function, `sleep_then_print` twice. One \n",
+    "invocation is ``sleep_then_print(2, \"Hello\")`` and the other is ``sleep_then_print(1, \", TBB!\")``.  \n",
+    "When executed sequentially, one after the other, the first function call will sleep for 2 seconds \n",
+    "then print \"Hello\" and the second function call will sleep for 1 more second and then print \", TBB!\".\n",
+    "\n",
+    "Inspect the code below - there are no modifications necessary. This code is our serial base case. Run the first \n",
+    "cell to create the file, then run the cell below it to compile and execute the code.\n",
+    "\n",
+    "1. Inspect the code cell below, then click run ▶ to save the code to a file\n",
+    "2. Run ▶ the cell in the __Build and Run__ section below the code snippet to compile and execute the code in the saved file"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/hello-tbb.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <iostream>\n",
+    "#include <thread>\n",
+    "\n",
+    "void sleep_then_print(unsigned sec, const char *str) {\n",
+    "  std::this_thread::sleep_for(std::chrono::seconds(sec));\n",
+    "  std::cout << str << std::flush;\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  std::cout << \"Running sequentially:\\n\";\n",
+    "  auto st0 = std::chrono::high_resolution_clock::now();\n",
+    "  sleep_then_print(2, \"Hello\");\n",
+    "  sleep_then_print(1, \", TBB!\");\n",
+    "  auto st1 = std::chrono::high_resolution_clock::now();\n",
+    "    \n",
+    "  std::cout << \"\\nSerial time   = \" << 1e-9 * (st1-st0).count() << \" seconds\\n\";\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and Run\n",
+    "Select the cell below and click Run ▶ to compile and execute the above example:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_hello-tbb.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_hello-tbb.sh; else ./scripts/run_hello-tbb.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Use tbb::parallel_invoke to add concurrency\n",
+    "\n",
+    "Now we will add some concurrency to the code.\n",
+    "\n",
+    "In the code cell below, the features of the TBB library are made available by including `tbb.h`. All oneTBB classes \n",
+    "and functions are in namespace `tbb`.\n",
+    "\n",
+    "An almost complete skeleton for the code is provide in the code cell. In the `main` function of our first example, \n",
+    "the same two function calls are made again but we will use `tbb::parallel_invoke` to execute them concurrently. \n",
+    "\n",
+    "The interface of parallel_invoke needed for this example is shown below:\n",
+    "\n",
+    "```cpp\n",
+    "template<typename... Functions>\n",
+    "void parallel_invoke(Functions&&... fs);\n",
+    "```\n",
+    "\n",
+    "This function executes two or more function objects in parallel. When executed concurrently, the two calls to \n",
+    "``sleep_then_print`` will take approximately 2 seconds in wallclock time to complete. And, as a consequence \n",
+    "of their concurrent execution, the output from the second call is likely to be printed first (it only slept \n",
+    "for 1 second before printing) and, as a result, the output might look jumbled, \", TBB!Hello\".\n",
+    "\n",
+    "For this exercise, complete the following steps:\n",
+    "\n",
+    "1. Inspect the code cell below and make the following modifications.\n",
+    "  1. Remove the `//` before all the lines, except the one with \"STEP 1\" in it.\n",
+    "2. When the modifications are complete, click run ▶ to save the code to a file.\n",
+    "3. Run ▶ the cell in the __Build and Run the modified code__ section below the code snippet to compile and execute the code in the saved file."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/hello-tbb-parallel.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <iostream>\n",
+    "#include <thread>\n",
+    "#include <tbb/tbb.h>\n",
+    "\n",
+    "void sleep_then_print(unsigned sec, const char *str) {\n",
+    "  std::this_thread::sleep_for(std::chrono::seconds(sec));\n",
+    "  std::cout << str << std::flush;\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  std::cout << \"\\nRunning in parallel:\\n\";\n",
+    "  auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "  // STEP 1: remove the '//' at the start of the lines to use parallel_invoke\n",
+    "  //  tbb::parallel_invoke(\n",
+    "  //  []() { \n",
+    "    sleep_then_print(2, \"Hello\"); \n",
+    "  // },\n",
+    "  //  []() { \n",
+    "    sleep_then_print(1, \", TBB!\"); \n",
+    "  // });\n",
+    "  auto pt1 = std::chrono::high_resolution_clock::now();\n",
+    "    \n",
+    "  std::cout << \"Parallel time = \" << 1e-9 * (pt1-pt0).count() << \" seconds\\n\";\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and Run\n",
+    "Select the cell below and click Run ▶ to compile and execute the above example:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_hello-tbb-parallel.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_hello-tbb-parallel.sh; else ./scripts/run_hello-tbb-parallel.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Solution (only peak if you need to)\n",
+    "\n",
+    "In general, a fully complete solution will be provided, but will be hidden in the notebook.  You can expand the cells to see,\n",
+    "save, build and run the solution if needed."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "%%writefile solutions/hello-tbb-solved.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <iostream>\n",
+    "#include <thread>\n",
+    "#include <tbb/tbb.h>\n",
+    "\n",
+    "void sleep_then_print(unsigned sec, const char *str) {\n",
+    "  std::this_thread::sleep_for(std::chrono::seconds(sec));\n",
+    "  std::cout << str << std::flush;\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  std::cout << \"\\nRunning in parallel:\\n\";\n",
+    "  auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "  // STEP 1: remove the '//' at the start of the lines to use parallel_invoke\n",
+    "  tbb::parallel_invoke(\n",
+    "    []() { \n",
+    "      sleep_then_print(2, \"Hello\"); \n",
+    "    },\n",
+    "    []() { \n",
+    "      sleep_then_print(1, \", TBB!\"); \n",
+    "    });\n",
+    "  auto pt1 = std::chrono::high_resolution_clock::now();\n",
+    "    \n",
+    "  std::cout << \"Parallel time = \" << 1e-9 * (pt1-pt0).count() << \" seconds\\n\";\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_hello-tbb-solved.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_hello-tbb-solved.sh; else ./scripts/run_hello-tbb-solved.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Next steps\n",
+    "\n",
+    "You have now completed a first exercise! The remaining modules have similar structure but include exercises\n",
+    "that demand more than just uncommenting lines.\n",
+    "\n",
+    "If you are ready, go to [the next module](../02_oneTBB_algorithms/oneTBB_algorithms.ipynb)."
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3.7 (Intel® oneAPI)",
+   "language": "python",
+   "name": "c009-intel_distribution_of_python_3_oneapi-beta05-python"
+  },
+  "language_info": {
+   "codemirror_mode": {
+    "name": "ipython",
+    "version": 3
+   },
+   "file_extension": ".py",
+   "mimetype": "text/x-python",
+   "name": "python",
+   "nbconvert_exporter": "python",
+   "pygments_lexer": "ipython3",
+   "version": "3.7.7"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 4
+}
@@ -0,0 +1,31 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <chrono>
+#include <iostream>
+#include <thread>
+#include <tbb/tbb.h>
+
+void sleep_then_print(unsigned sec, const char *str) {
+  std::this_thread::sleep_for(std::chrono::seconds(sec));
+  std::cout << str << std::flush;
+}
+
+int main() {
+  std::cout << "\nRunning in parallel:\n";
+  auto pt0 = std::chrono::high_resolution_clock::now();
+  // STEP 1: remove the '//' at the start of the lines to use parallel_invoke
+  //  tbb::parallel_invoke(
+  //  []() { 
+    sleep_then_print(2, "Hello"); 
+  // },
+  //  []() { 
+    sleep_then_print(1, ", TBB!"); 
+  // });
+  auto pt1 = std::chrono::high_resolution_clock::now();
+    
+  std::cout << "Parallel time = " << 1e-9 * (pt1-pt0).count() << " seconds\n";
+}
@@ -0,0 +1,25 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <chrono>
+#include <iostream>
+#include <thread>
+#include <tbb/tbb.h>
+
+void sleep_then_print(unsigned sec, const char *str) {
+  std::this_thread::sleep_for(std::chrono::seconds(sec));
+  std::cout << str << std::flush;
+}
+
+int main() {
+  std::cout << "Running sequentially:\n";
+  auto st0 = std::chrono::high_resolution_clock::now();
+  sleep_then_print(2, "Hello");
+  sleep_then_print(1, ", TBB!");
+  auto st1 = std::chrono::high_resolution_clock::now();
+    
+  std::cout << "\nSerial time   = " << 1e-9 * (st1-st0).count() << " seconds\n";
+}
@@ -0,0 +1 @@
+bin directory needed for notebook
@@ -0,0 +1 @@
+lab directory needed for notebook
@@ -0,0 +1,56 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+# Script to submit job in Intel(R) DevCloud
+# Version: 0.7
+#==========================================
+if [ -z "$1" ]; then
+    echo "Missing script argument, Usage: ./q run.sh"
+elif [ ! -f "$1" ]; then
+    echo "File $1 does not exist"
+else
+    echo "Job has been submitted to Intel(R) DevCloud and will execute soon."
+    echo ""
+    echo " If you do not see result in 60 seconds, please restart the Jupyter kernel:"
+    echo " Kernel -> 'Restart Kernel and Clear All Outputs...' and then try again"
+    echo ""
+    script=$1
+    shift
+    args=$@
+    # Remove old output files
+    rm *.sh.* > /dev/null 2>&1
+    # Submit job using qsub
+    qsub_id=`qsub -l nodes=1:ppn=2 -d . $script $@`
+    job_id="$(cut -d'.' -f1 <<<"$qsub_id")"
+    # Print qstat output
+    qstat 
+    # Wait for output file to be generated and display
+    echo ""
+    echo -ne "Waiting for Output "
+    script_name=${script/scripts\//}
+    echo $script_name.o$job_id
+    until [ -f $script_name.o$job_id ]; do
+        sleep 1
+        echo -ne "█"
+        ((timeout++))
+        # Timeout if no output file generated within 60 seconds
+        if [ $timeout == 60 ]; then
+            echo ""
+            echo ""
+            echo "TimeOut 60 seconds: Job is still queued for execution, check for output file later ($script.o$job_id)"
+            echo ""
+            break
+        fi
+    done
+    # Print output and error file content if exist
+    if [ -n "$(find -name '*.sh.o'$job_id)" ]; then
+        echo " Done⬇"
+        cat $script_name.o$job_id
+        cat $script_name.e$job_id
+        echo "Job Completed in $timeout seconds."
+        rm *.sh.*$job_id > /dev/null 2>&1
+    fi
+fi
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction sample hello-tbb-parallel 
+dpcpp -std=c++17 -O2 lab/hello-tbb-parallel.cpp -o bin/hello-tbb-parallel -tbb
+bin/hello-tbb-parallel
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction sample hello-tbb-solved 
+dpcpp -std=c++17 -O2 solutions/hello-tbb-solved.cpp -o bin/hello-tbb-solved -tbb
+bin/hello-tbb-solved
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction sample hello-tbb 
+dpcpp -std=c++17 -O2 lab/hello-tbb.cpp -o bin/hello-tbb -tbb
+bin/hello-tbb
@@ -0,0 +1 @@
+solutions directory needed for notebook
@@ -0,0 +1,31 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <chrono>
+#include <iostream>
+#include <thread>
+#include <tbb/tbb.h>
+
+void sleep_then_print(unsigned sec, const char *str) {
+  std::this_thread::sleep_for(std::chrono::seconds(sec));
+  std::cout << str << std::flush;
+}
+
+int main() {
+  std::cout << "\nRunning in parallel:\n";
+  auto pt0 = std::chrono::high_resolution_clock::now();
+  // STEP 1: remove the '//' at the start of the lines to use parallel_invoke
+  tbb::parallel_invoke(
+    []() { 
+      sleep_then_print(2, "Hello"); 
+    },
+    []() { 
+      sleep_then_print(1, ", TBB!"); 
+    });
+  auto pt1 = std::chrono::high_resolution_clock::now();
+    
+  std::cout << "Parallel time = " << 1e-9 * (pt1-pt0).count() << " seconds\n";
+}
@@ -0,0 +1,60 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <chrono>
+#include <cmath>
+#include <iostream>
+#include <limits>
+#include <thread>
+#include <tbb/tbb.h>
+
+#define INCORRECT_VALUE 1
+#define INCORRECT_FUNCTION std::minus<double>()
+
+double calc_pi(int num_intervals) {
+  double dx = 1.0 / num_intervals;
+  double sum = tbb::parallel_reduce(
+    /* STEP 1: fix the upper bound: */ tbb::blocked_range<int>(0, INCORRECT_VALUE), 
+    /* STEP 2: provide a proper identity value for summation */ INCORRECT_VALUE,
+    /* func */ 
+    [=](const tbb::blocked_range<int>& r, double init) -> double {
+      for (int i = r.begin(); i != r.end(); ++i) {
+        // STEP 3: Add the loop body code:
+        //         Hint: it will look a lot like the the sequential code.
+        //               the returned value should be (init + the_partial_sum)
+      }
+      return init;
+    },
+    // STEP 4: provide the reduction function
+    //         Hint, maybe std::plus<double>{}
+    INCORRECT_FUNCTION
+  );
+  double pi = 4 * sum;
+  return pi;
+}
+
+static void warmupTBB() {
+  int num_threads = std::thread::hardware_concurrency();
+  tbb::parallel_for(0, num_threads,
+    [](unsigned int) { 
+      std::this_thread::sleep_for(std::chrono::milliseconds(10)); 
+  });
+}
+
+int main() {
+  const int num_intervals = std::numeric_limits<int>::max();
+  double parallel_time = 0.0;
+  warmupTBB();
+  {
+    auto pt0 = std::chrono::high_resolution_clock::now();
+    double pi = calc_pi(num_intervals);
+    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();
+    std::cout << "parallel pi == " << pi << std::endl;
+  }
+
+  std::cout << "parallel_time == " << parallel_time << " seconds" << std::endl;
+  return 0;
+}
@@ -0,0 +1,36 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <chrono>
+#include <cmath>
+#include <iostream>
+#include <limits>
+
+double calc_pi(int num_intervals) {
+  double dx = 1.0 / num_intervals;
+  double sum = 0.0;
+  for (int i = 0; i < num_intervals; ++i) {
+    double x = (i+0.5)*dx;
+    double h = std::sqrt(1-x*x);
+    sum += h*dx;
+  }
+  double pi = 4 * sum;
+  return pi;
+}
+
+int main() {
+  const int num_intervals = std::numeric_limits<int>::max();
+  double serial_time = 0.0;
+  {
+    auto st0 = std::chrono::high_resolution_clock::now();
+    double pi = calc_pi(num_intervals);
+    serial_time = 1e-9*(std::chrono::high_resolution_clock::now() - st0).count();
+    std::cout << "serial pi == " << pi << std::endl;
+  }
+
+  std::cout << "serial_time == " << serial_time << " seconds" << std::endl;
+  return 0;
+}
@@ -0,0 +1,83 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <algorithm>
+#include <cctype>
+#include <chrono>
+#include <fstream>
+#include <iostream>
+#include <memory>
+#include <string>
+#include <thread>
+
+#include <tbb/tbb.h>
+
+#define INCORRECT_VALUE 1
+#define INCORRECT_MODE_A tbb::filter_mode::serial_in_order
+#define INCORRECT_MODE_B tbb::filter_mode::parallel
+
+using CaseStringPtr = std::shared_ptr<std::string>;
+
+//
+// These functions are defined in common/case.cpp
+//
+void initCaseChange(int num_strings, int string_len, int free_list_size);
+CaseStringPtr getCaseString(std::ofstream& f); 
+void writeCaseString(std::ofstream& f, CaseStringPtr s);
+
+void change_case_parallel(int num_tokens, std::ofstream& caseBeforeFile, std::ofstream& caseAfterFile) {
+  tbb::parallel_pipeline(
+    /* STEP A: set the number of tokens */ INCORRECT_VALUE,
+    tbb::make_filter<void, CaseStringPtr>(
+      tbb::filter_mode::serial_in_order,
+      [&](tbb::flow_control& fc) -> CaseStringPtr {
+        CaseStringPtr s_ptr = getCaseString(caseBeforeFile);
+        if (!s_ptr) 
+          fc.stop();
+        return s_ptr; 
+      }) &
+    tbb::make_filter<CaseStringPtr, CaseStringPtr>(
+      /* STEP B: Change the mode to be parallel */ INCORRECT_MODE_A,
+      [](CaseStringPtr s_ptr) -> CaseStringPtr {
+        /* STEP C: add the body that converts the incoming s_ptr */
+        return s_ptr;
+      }) & // concatenation operation
+    tbb::make_filter<CaseStringPtr, void>(
+      /* STEP D: Change the mode to serial_in_order */ INCORRECT_MODE_B,
+      [&](CaseStringPtr s_ptr) -> void {
+        /* STEP E: add the body to write the output */
+      }) 
+  );
+}
+
+static void warmupTBB() {
+  int num_threads = std::thread::hardware_concurrency();
+  tbb::parallel_for(0, num_threads,
+    [](unsigned int) { 
+      std::this_thread::sleep_for(std::chrono::milliseconds(10)); 
+  });
+}
+
+int main() {
+  int num_tokens = std::thread::hardware_concurrency();
+  int num_strings = 100; 
+  int string_len = 100000;
+  int free_list_size = num_tokens;
+
+  std::ofstream caseBeforeFile("lab/parallel_pipeline_before.txt");
+  std::ofstream caseAfterFile("lab/parallel_pipeline_after.txt");
+  initCaseChange(num_strings, string_len, free_list_size);
+
+  warmupTBB();
+  double parallel_time = 0.0;
+  {
+    auto pt0 = std::chrono::high_resolution_clock::now();
+    change_case_parallel(num_tokens, caseBeforeFile, caseAfterFile);
+    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();
+  }
+  std::cout << "parallel_time == " << parallel_time << " seconds" << std::endl;
+  return 0;
+}
@@ -0,0 +1,57 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <algorithm>
+#include <cctype>
+#include <chrono>
+#include <fstream>
+#include <iostream>
+#include <memory>
+#include <string>
+
+using CaseStringPtr = std::shared_ptr<std::string>;
+
+//
+// These functions are defined in common/case.cpp
+//
+void initCaseChange(int num_strings, int string_len, int free_list_size);
+CaseStringPtr getCaseString(std::ofstream& f); 
+void writeCaseString(std::ofstream& f, CaseStringPtr s);
+    
+void change_case_serial(std::ofstream& caseBeforeFile, std::ofstream& caseAfterFile) {
+  while (CaseStringPtr s_ptr = getCaseString(caseBeforeFile)) {
+    std::transform(s_ptr->begin(), s_ptr->end(), s_ptr->begin(), 
+      [](char c) -> char {
+        if (std::islower(c))
+          return std::toupper(c);
+        else if (std::isupper(c))
+          return std::tolower(c);
+        else
+          return c;
+      }
+    );
+    writeCaseString(caseAfterFile, s_ptr);
+  }
+}
+
+int main() {
+  int num_strings = 100; 
+  int string_len = 100000;
+  int free_list_size = 1;
+
+  std::ofstream caseBeforeFile("lab/serial_pipeline_before.txt");
+  std::ofstream caseAfterFile("lab/serial_pipeline_after.txt");
+  initCaseChange(num_strings, string_len, free_list_size);
+
+  double serial_time = 0.0;
+  {
+    auto pt0 = std::chrono::high_resolution_clock::now();
+    change_case_serial(caseBeforeFile, caseAfterFile);
+    serial_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();
+  }
+  std::cout << "serial_time == " << serial_time << " seconds" << std::endl;
+  return 0;
+}
@@ -0,0 +1 @@
+bin directory needed for notebook
@@ -0,0 +1,55 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <algorithm>
+#include <cctype>
+#include <fstream>
+#include <iostream>
+#include <memory>
+#include <string>
+#include <tbb/tbb.h>
+
+using CaseStringPtr = std::shared_ptr<std::string>;
+static tbb::concurrent_queue<CaseStringPtr> caseFreeList;
+static int numCaseInputs = 0;
+
+void initCaseChange(int num_strings, int string_len, int free_list_size) {
+  numCaseInputs = num_strings;
+  caseFreeList.clear();
+  for (int i = 0; i < free_list_size; ++i) {
+    caseFreeList.push(std::make_shared<std::string>(string_len, ' '));
+  }
+}
+
+CaseStringPtr getCaseString(std::ofstream& f) {
+  std::shared_ptr<std::string> s;
+  if (numCaseInputs > 0) {
+    if (!caseFreeList.try_pop(s) || !s) {
+      return CaseStringPtr{};
+    }
+    int ascii_range = 'z' - 'A' + 2;
+    for (int i = 0; i < s->size(); ++i) {
+      int offset = i%ascii_range;
+      if (offset)
+        (*s)[i] = 'A' + offset - 1; 
+      else
+        (*s)[i] = '\n';
+    } 
+    if (f.good()) {
+      f << *s;
+    }
+    --numCaseInputs;
+  }
+  return s;
+}
+
+void writeCaseString(std::ofstream& f, CaseStringPtr s) {
+  if (f.good()) {
+    f << *s;
+  }
+  caseFreeList.push(s);
+}
+
@@ -0,0 +1 @@
+lab directory needed for notebook
@@ -0,0 +1,776 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# oneTBB generic parallel algorithms"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "##### Sections\n",
+    "- [oneTBB Generic Algorithms](#oneTBB-Generic-Algorithms)\n",
+    "- _Code_: [Calculating pi with tbb::parallel_reduce](#Calculating-pi-with-tbb::parallel_reduce)\n",
+    "- _Code_: [Processing text with tbb::parallel_pipeline](#Processing-text-with-tbb::parallel_pipeline)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Learning Objectives\n",
+    "\n",
+    "* Gain experince with oneTBB generic algorithms \n",
+    "* Use tbb::parallel_reduce to estimate pi as the area of a unit circle\n",
+    "* Use tbb::parallel_pipeline to create a simple text processing pipeline "
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# oneTBB Generic Algorithms"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "While it's possible to implement a parallel application by using oneTBB to specify each individual task that can run\n",
+    "concurrently, it is more common to make use of one of its data parallel generic algorithms. The oneTBB library provides \n",
+    "a number of [generic parallel algorithms](https://spec.oneapi.com/versions/latest/elements/oneTBB/source/algorithms.html),\n",
+    "including `parallel_for`, `parallel_reduce`, `parallel_scan`, `parallel_invoke` and `parallel_pipeline`. These functions \n",
+    "capture many of the common parallel patterns that are key to unlocking multithreaded performance. \n",
+    "\n",
+    "In this section, we provide two exercises that will introduce you to two of these functions."
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Calculating pi with tbb::parallel_reduce\n",
+    "\n",
+    "In this exercise, we calculate pi using the approach shown in the figure below. The idea is to\n",
+    "compute the area of a unit circle, which is equal to pi. We do this by approximating the area of \n",
+    "1/4th of a unit circle, summing up the areas of ``num_intervals`` rectangles that have\n",
+    "a height of ``sqrt(1-x*x)`` and a width of ``dx == 1.0/num_intervals``. This sum is multiplied by \n",
+    "4 to compute the total area of the unit circle, providing us with an approximation for pi.\n",
+    "\n",
+    "![Algorithm to compute pi](img/pi.png)\n",
+    "\n",
+    "### Run the sequential baseline implementation\n",
+    "\n",
+    "Before we add any parallelism, let's validate this approach by running a baseline sequential implementation. Inspect \n",
+    "the sequential code below - there are no modifications necessary. Run the first cell to create the file, then run the \n",
+    "cell below it to compile and execute the code. This represents the baseline sequential result and time for our pi \n",
+    "computation exercise.\n",
+    "\n",
+    "1. Inspect the code cell below, then click run ▶ to save the code to a file\n",
+    "2. Run ▶ the cell in the __Build and Run the baseline__ section below the code snippet to compile and execute the code in the saved file"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/pi-serial.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <cmath>\n",
+    "#include <iostream>\n",
+    "#include <limits>\n",
+    "\n",
+    "double calc_pi(int num_intervals) {\n",
+    "  double dx = 1.0 / num_intervals;\n",
+    "  double sum = 0.0;\n",
+    "  for (int i = 0; i < num_intervals; ++i) {\n",
+    "    double x = (i+0.5)*dx;\n",
+    "    double h = std::sqrt(1-x*x);\n",
+    "    sum += h*dx;\n",
+    "  }\n",
+    "  double pi = 4 * sum;\n",
+    "  return pi;\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  const int num_intervals = std::numeric_limits<int>::max();\n",
+    "  double serial_time = 0.0;\n",
+    "  {\n",
+    "    auto st0 = std::chrono::high_resolution_clock::now();\n",
+    "    double pi = calc_pi(num_intervals);\n",
+    "    serial_time = 1e-9*(std::chrono::high_resolution_clock::now() - st0).count();\n",
+    "    std::cout << \"serial pi == \" << pi << std::endl;\n",
+    "  }\n",
+    "\n",
+    "  std::cout << \"serial_time == \" << serial_time << \" seconds\" << std::endl;\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and Run the baseline\n",
+    "Select the cell below and click Run ▶ to compile and execute the code above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_pi-serial.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_pi-serial.sh; else ./scripts/run_pi-serial.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Implement a parallel version with tbb::parallel_reduce\n",
+    "\n",
+    "Our sequential code accumulates values into a single final sum, making it a reduction operation and a match for ``tbb::parallel_reduce``.\n",
+    "You can find detailed documentation for ``parallel_reduce`` [here](https://software.intel.com/content/www/us/en/develop/documentation/tbb-documentation/top/intel-threading-building-blocks-developer-reference/algorithms/parallelreduce-template-function.html). Briefly though, a ``parallel_reduce`` runs a user-provided function\n",
+    "on chunks of the iteration space, potentially concurrently, resulting in several partial results. In our example, these partial results will be partial sums. These partial results are combined using a user-provided reduction function, in our pi example, `std::plus` might be used (hint). \n",
+    "\n",
+    "The interface of ``parallel_reduce`` needed for this example is shown below:\n",
+    "\n",
+    "```cpp\n",
+    "template<typename Range, typename Value, typename Func, typename Reduction>\n",
+    "Value parallel_reduce( const Range& range, const Value& identity,\n",
+    "                       const Func& func, const Reduction& reduction );\n",
+    "```\n",
+    "\n",
+    "The ``range`` object provides the iteration space, which in our example is 0 to num_intervals - 1. ``identity`` is the identity value for the \n",
+    "operation that is being parallelized; for a summation, the identity value is 0, since ``sum == sum + 0``. We provide a lambda expression for \n",
+    "``func`` to compute the partial results, which in our example will return a partial sum for a given range ``r``, accumulating into the \n",
+    "starting value ``init``. Finally, ``reduction`` is the operation to use to combine the partial results.\n",
+    "\n",
+    "For this exercise, complete the following steps:\n",
+    "\n",
+    "1. Inspect the code cell below and make the following modifications.\n",
+    "  1. Fix the upper bound in the ``tbb::blocked_range``\n",
+    "  2. Fix the identity value\n",
+    "  3. Add the loop body code\n",
+    "  4. Fix the reduction function\n",
+    "2. When the modifications are complete, click run ▶ to save the code to a file.\n",
+    "3. Run ▶ the cell in the __Build and Run the modified code__ section below the code snippet to compile and execute the code in the saved file."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/pi-parallel.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <cmath>\n",
+    "#include <iostream>\n",
+    "#include <limits>\n",
+    "#include <thread>\n",
+    "#include <tbb/tbb.h>\n",
+    "\n",
+    "#define INCORRECT_VALUE 1\n",
+    "#define INCORRECT_FUNCTION std::minus<double>()\n",
+    "\n",
+    "double calc_pi(int num_intervals) {\n",
+    "  double dx = 1.0 / num_intervals;\n",
+    "  double sum = tbb::parallel_reduce(\n",
+    "    /* STEP 1: fix the upper bound: */ tbb::blocked_range<int>(0, INCORRECT_VALUE), \n",
+    "    /* STEP 2: provide a proper identity value for summation */ INCORRECT_VALUE,\n",
+    "    /* func */ \n",
+    "    [=](const tbb::blocked_range<int>& r, double init) -> double {\n",
+    "      for (int i = r.begin(); i != r.end(); ++i) {\n",
+    "        // STEP 3: Add the loop body code:\n",
+    "        //         Hint: it will look a lot like the the sequential code.\n",
+    "        //               the returned value should be (init + the_partial_sum)\n",
+    "      }\n",
+    "      return init;\n",
+    "    },\n",
+    "    // STEP 4: provide the reduction function\n",
+    "    //         Hint, maybe std::plus<double>{}\n",
+    "    INCORRECT_FUNCTION\n",
+    "  );\n",
+    "  double pi = 4 * sum;\n",
+    "  return pi;\n",
+    "}\n",
+    "\n",
+    "static void warmupTBB() {\n",
+    "  int num_threads = std::thread::hardware_concurrency();\n",
+    "  tbb::parallel_for(0, num_threads,\n",
+    "    [](unsigned int) { \n",
+    "      std::this_thread::sleep_for(std::chrono::milliseconds(10)); \n",
+    "  });\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  const int num_intervals = std::numeric_limits<int>::max();\n",
+    "  double parallel_time = 0.0;\n",
+    "  warmupTBB();\n",
+    "  {\n",
+    "    auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "    double pi = calc_pi(num_intervals);\n",
+    "    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();\n",
+    "    std::cout << \"parallel pi == \" << pi << std::endl;\n",
+    "  }\n",
+    "\n",
+    "  std::cout << \"parallel_time == \" << parallel_time << \" seconds\" << std::endl;\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and Run the modified code\n",
+    "\n",
+    "Select the cell below and click Run ▶ to compile and execute the code that you modified above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_pi-parallel.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_pi-parallel.sh; else ./scripts/run_pi-parallel.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Pi Example Solution (Don't peak, unless you have to)"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "%%writefile solutions/pi-parallel.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <cmath>\n",
+    "#include <iostream>\n",
+    "#include <limits>\n",
+    "#include <thread>\n",
+    "#include <tbb/tbb.h>\n",
+    "\n",
+    "double calc_pi(int num_intervals) {\n",
+    "  double dx = 1.0 / num_intervals;\n",
+    "  double sum = tbb::parallel_reduce(\n",
+    "    /* range = */ tbb::blocked_range<int>(0, num_intervals ), \n",
+    "    /* identity = */ 0.0,\n",
+    "    /* func */ \n",
+    "    [=](const tbb::blocked_range<int>& r, double init) -> double {\n",
+    "      for (int i = r.begin(); i != r.end(); ++i) {\n",
+    "        double x = (i+0.5)*dx;\n",
+    "        double h = std::sqrt(1-x*x);\n",
+    "        init += h*dx;\n",
+    "      }\n",
+    "      return init;\n",
+    "    },\n",
+    "    std::plus<double>{}\n",
+    "  );\n",
+    "  double pi = 4 * sum;\n",
+    "  return pi;\n",
+    "}\n",
+    "\n",
+    "static void warmupTBB() {\n",
+    "  int num_threads = std::thread::hardware_concurrency();\n",
+    "  tbb::parallel_for(0, num_threads,\n",
+    "    [](unsigned int) { \n",
+    "      std::this_thread::sleep_for(std::chrono::milliseconds(10)); \n",
+    "  });\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  const int num_intervals = std::numeric_limits<int>::max();\n",
+    "  double parallel_time = 0.0;\n",
+    "  warmupTBB();\n",
+    "  {\n",
+    "    auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "    double pi = calc_pi(num_intervals);\n",
+    "    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();\n",
+    "    std::cout << \"parallel pi == \" << pi << std::endl;\n",
+    "  }\n",
+    "\n",
+    "  std::cout << \"parallel_time == \" << parallel_time << \" seconds\" << std::endl;\n",
+    "  return 0;\n",
+    "}\n"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_pi-solution.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_pi-solution.sh; else ./scripts/run_pi-solution.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### "
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Processing text with tbb::parallel_pipeline\n",
+    "\n",
+    "The loop algorithms provided by oneTBB, like tbb::parallel_reduce, operate over a range that is know in advance. In this section, \n",
+    "we will look at an example where we don't know the range in advance, and instead continue processing data until we dynamically discover \n",
+    "that there is no more data to process. Our example processes a large text file, a chunk at a time, until there are no more chunks to read. \n",
+    "To do so, we will use a simple three-stage pipeline that (1) gets a buffer of text if one exists, (2) processes the text buffer and \n",
+    "(3) writes the processed buffer to an output file.\n",
+    "\n",
+    "![A simple three-stage pipeline](img/pipeline.png)\n",
+    "\n",
+    "### Run the sequential baseline implementation\n",
+    "\n",
+    "Again, let's start with a sequential implementation that serves as a baseline for both a correct result and the execution \n",
+    "time.\n",
+    "\n",
+    "Inspect the sequential code below - there are no modifications necessary. We can notice that the baseline code uses a \n",
+    "``while``-loop since the number of chunks is not known upfront. Run the first cell to create the source file, then run the \n",
+    "cell below it to compile and execute the code. This represents the baseline for the computation that we will convert into \n",
+    "a parallel pipeline.\n",
+    "\n",
+    "1. Inspect the code cell below, then click run ▶ to save the code to a file\n",
+    "2. Run ▶ the cell in the __Build and Run the baseline__ section below the code snippet to compile and execute the code in the saved file\n",
+    "3. Inspect the two new files ``serial_pipeline_before.txt`` and ``serial_pipeline_after.txt`` to see the results of changing the case of the text."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/pipeline-serial.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <algorithm>\n",
+    "#include <cctype>\n",
+    "#include <chrono>\n",
+    "#include <fstream>\n",
+    "#include <iostream>\n",
+    "#include <memory>\n",
+    "#include <string>\n",
+    "\n",
+    "using CaseStringPtr = std::shared_ptr<std::string>;\n",
+    "\n",
+    "//\n",
+    "// These functions are defined in common/case.cpp\n",
+    "//\n",
+    "void initCaseChange(int num_strings, int string_len, int free_list_size);\n",
+    "CaseStringPtr getCaseString(std::ofstream& f); \n",
+    "void writeCaseString(std::ofstream& f, CaseStringPtr s);\n",
+    "    \n",
+    "void change_case_serial(std::ofstream& caseBeforeFile, std::ofstream& caseAfterFile) {\n",
+    "  while (CaseStringPtr s_ptr = getCaseString(caseBeforeFile)) {\n",
+    "    std::transform(s_ptr->begin(), s_ptr->end(), s_ptr->begin(), \n",
+    "      [](char c) -> char {\n",
+    "        if (std::islower(c))\n",
+    "          return std::toupper(c);\n",
+    "        else if (std::isupper(c))\n",
+    "          return std::tolower(c);\n",
+    "        else\n",
+    "          return c;\n",
+    "      }\n",
+    "    );\n",
+    "    writeCaseString(caseAfterFile, s_ptr);\n",
+    "  }\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  int num_strings = 100; \n",
+    "  int string_len = 100000;\n",
+    "  int free_list_size = 1;\n",
+    "\n",
+    "  std::ofstream caseBeforeFile(\"lab/serial_pipeline_before.txt\");\n",
+    "  std::ofstream caseAfterFile(\"lab/serial_pipeline_after.txt\");\n",
+    "  initCaseChange(num_strings, string_len, free_list_size);\n",
+    "\n",
+    "  double serial_time = 0.0;\n",
+    "  {\n",
+    "    auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "    change_case_serial(caseBeforeFile, caseAfterFile);\n",
+    "    serial_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();\n",
+    "  }\n",
+    "  std::cout << \"serial_time == \" << serial_time << \" seconds\" << std::endl;\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and Run the baseline\n",
+    "Select the cell below and click Run ▶ to compile and execute the code above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_pipeline-serial.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_pipeline-serial.sh; else ./scripts/run_pipeline-serial.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Implement a parallel version with tbb::parallel_pipeline\n",
+    "\n",
+    "The oneTBB library provides a function ``tbb::parallel_pipeline`` which is for exactly cases like this. In this section, \n",
+    "we will use a ``tbb::parallel_pipeline`` to create a parallel implementation of the text processing computation.\n",
+    "\n",
+    "You can find detailed documentation for parallel_pipeline [here](https://software.intel.com/content/www/us/en/develop/documentation/tbb-documentation/top/intel-threading-building-blocks-developer-reference/algorithms/parallelpipeline-function.html). But briefly, to create a oneTBB\n",
+    "pipeline, we need to provide a ``max_number_of_live_tokens`` value and a chain of pipeline filter objects. ``max_number_of_live_tokens`` \n",
+    "provides a limit on the number of items that can be in flight in the pipeline concurrently. \n",
+    "\n",
+    "And we create filters using the ``make_filter`` function:\n",
+    "\n",
+    "```cpp\n",
+    "template<typename T, typename U, typename Func> \n",
+    "filter_t<T,U> make_filter(filter::mode mode, const Func& f)\n",
+    "```\n",
+    "\n",
+    "The ``mode`` argument can be ``tbb::filter_mode::serial_in_order``, ``tbb::filter_mode::serial_out_of_order``, or ``tbb::filter_mode::parallel``. \n",
+    "A chain of filters is created by concatenating filters using `operator&`. \n",
+    "\n",
+    "A summary of the interface for ``parallel_pipeline`` needed to complete this exercise is shown below:\n",
+    "\n",
+    "```cpp\n",
+    "void parallel_pipeline( size_t max_number_of_live_tokens, \n",
+    "                        const filter_t<void,void>& filter_chain);\n",
+    "```\n",
+    "\n",
+    "In the code cell below, we provide a skeleton of the solution. You need to complete the following steps:\n",
+    "\n",
+    "1. Inspect the code cell below and make the following modifications.\n",
+    "  1. Set the number of tokens to use in the ``tbb::parallel_pipeline`` to ``num_tokens``\n",
+    "  2. Change the mode of the second filter to ``tbb::filter_mode::parallel``\n",
+    "  3. Add the body code to the second filter\n",
+    "  4. Change the mode of the final filter to ``tbb::filter_mode::serial_in_order``\n",
+    "  5. Add the body code to the final filter\n",
+    "2. When the modifications are complete, click run ▶ to save the code to a file.\n",
+    "3. Run ▶ the cell in the __Build and Run the modified code__ section below the code snippet to compile and execute the code in the saved file.\n",
+    "4. Inspect the two new files ``parallel_pipeline_before.txt`` and ``parallel_pipeline_after.txt`` to see the results of changing the case of the text."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/pipeline-parallel.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <algorithm>\n",
+    "#include <cctype>\n",
+    "#include <chrono>\n",
+    "#include <fstream>\n",
+    "#include <iostream>\n",
+    "#include <memory>\n",
+    "#include <string>\n",
+    "#include <thread>\n",
+    "\n",
+    "#include <tbb/tbb.h>\n",
+    "\n",
+    "#define INCORRECT_VALUE 1\n",
+    "#define INCORRECT_MODE_A tbb::filter_mode::serial_in_order\n",
+    "#define INCORRECT_MODE_B tbb::filter_mode::parallel\n",
+    "\n",
+    "using CaseStringPtr = std::shared_ptr<std::string>;\n",
+    "\n",
+    "//\n",
+    "// These functions are defined in common/case.cpp\n",
+    "//\n",
+    "void initCaseChange(int num_strings, int string_len, int free_list_size);\n",
+    "CaseStringPtr getCaseString(std::ofstream& f); \n",
+    "void writeCaseString(std::ofstream& f, CaseStringPtr s);\n",
+    "\n",
+    "void change_case_parallel(int num_tokens, std::ofstream& caseBeforeFile, std::ofstream& caseAfterFile) {\n",
+    "  tbb::parallel_pipeline(\n",
+    "    /* STEP A: set the number of tokens */ INCORRECT_VALUE,\n",
+    "    tbb::make_filter<void, CaseStringPtr>(\n",
+    "      tbb::filter_mode::serial_in_order,\n",
+    "      [&](tbb::flow_control& fc) -> CaseStringPtr {\n",
+    "        CaseStringPtr s_ptr = getCaseString(caseBeforeFile);\n",
+    "        if (!s_ptr) \n",
+    "          fc.stop();\n",
+    "        return s_ptr; \n",
+    "      }) &\n",
+    "    tbb::make_filter<CaseStringPtr, CaseStringPtr>(\n",
+    "      /* STEP B: Change the mode to be parallel */ INCORRECT_MODE_A,\n",
+    "      [](CaseStringPtr s_ptr) -> CaseStringPtr {\n",
+    "        /* STEP C: add the body that converts the incoming s_ptr */\n",
+    "        return s_ptr;\n",
+    "      }) & // concatenation operation\n",
+    "    tbb::make_filter<CaseStringPtr, void>(\n",
+    "      /* STEP D: Change the mode to serial_in_order */ INCORRECT_MODE_B,\n",
+    "      [&](CaseStringPtr s_ptr) -> void {\n",
+    "        /* STEP E: add the body to write the output */\n",
+    "      }) \n",
+    "  );\n",
+    "}\n",
+    "\n",
+    "static void warmupTBB() {\n",
+    "  int num_threads = std::thread::hardware_concurrency();\n",
+    "  tbb::parallel_for(0, num_threads,\n",
+    "    [](unsigned int) { \n",
+    "      std::this_thread::sleep_for(std::chrono::milliseconds(10)); \n",
+    "  });\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  int num_tokens = std::thread::hardware_concurrency();\n",
+    "  int num_strings = 100; \n",
+    "  int string_len = 100000;\n",
+    "  int free_list_size = num_tokens;\n",
+    "\n",
+    "  std::ofstream caseBeforeFile(\"lab/parallel_pipeline_before.txt\");\n",
+    "  std::ofstream caseAfterFile(\"lab/parallel_pipeline_after.txt\");\n",
+    "  initCaseChange(num_strings, string_len, free_list_size);\n",
+    "\n",
+    "  warmupTBB();\n",
+    "  double parallel_time = 0.0;\n",
+    "  {\n",
+    "    auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "    change_case_parallel(num_tokens, caseBeforeFile, caseAfterFile);\n",
+    "    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();\n",
+    "  }\n",
+    "  std::cout << \"parallel_time == \" << parallel_time << \" seconds\" << std::endl;\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and Run the modified code\n",
+    "\n",
+    "Select the cell below and click Run ▶ to compile and execute the code that you modified above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_pipeline-parallel.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_pipeline-parallel.sh; else ./scripts/run_pipeline-parallel.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Case-Change Pipeline Solution (Don't peak, unless you have to)"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "%%writefile solutions/pipeline-solution.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <algorithm>\n",
+    "#include <cctype>\n",
+    "#include <chrono>\n",
+    "#include <fstream>\n",
+    "#include <iostream>\n",
+    "#include <memory>\n",
+    "#include <string>\n",
+    "#include <thread>\n",
+    "\n",
+    "#include <tbb/tbb.h>\n",
+    "\n",
+    "using CaseStringPtr = std::shared_ptr<std::string>;\n",
+    "\n",
+    "//\n",
+    "// These functions are defined in common/case.cpp\n",
+    "//\n",
+    "void initCaseChange(int num_strings, int string_len, int free_list_size);\n",
+    "CaseStringPtr getCaseString(std::ofstream& f); \n",
+    "void writeCaseString(std::ofstream& f, CaseStringPtr s);\n",
+    "\n",
+    "void change_case_parallel(int num_tokens, std::ofstream& caseBeforeFile, std::ofstream& caseAfterFile) {\n",
+    "  tbb::parallel_pipeline(\n",
+    "    /* tokens */ num_tokens,\n",
+    "    /* the get filter */\n",
+    "    tbb::make_filter<void, CaseStringPtr>(\n",
+    "      /* filter node */ tbb::filter_mode::serial_in_order,\n",
+    "      /* filter body */\n",
+    "      [&](tbb::flow_control& fc) -> CaseStringPtr {\n",
+    "        CaseStringPtr s_ptr = getCaseString(caseBeforeFile);\n",
+    "        if (!s_ptr) \n",
+    "          fc.stop();\n",
+    "        return s_ptr; \n",
+    "      }) & // concatenation operation\n",
+    "    /* make the change case filter */\n",
+    "    tbb::make_filter<CaseStringPtr, CaseStringPtr>(\n",
+    "      /* filter node */ tbb::filter_mode::parallel,\n",
+    "      /* filter body */\n",
+    "      [](CaseStringPtr s_ptr) -> CaseStringPtr {\n",
+    "        std::transform(s_ptr->begin(), s_ptr->end(), s_ptr->begin(), \n",
+    "          [](char c) -> char {\n",
+    "            if (std::islower(c))\n",
+    "              return std::toupper(c);\n",
+    "            else if (std::isupper(c))\n",
+    "              return std::tolower(c);\n",
+    "            else\n",
+    "              return c;\n",
+    "          });\n",
+    "        return s_ptr;\n",
+    "      }) & // concatenation operation\n",
+    "    /* make the write filter */\n",
+    "    tbb::make_filter<CaseStringPtr, void>(\n",
+    "      /* filter node */ tbb::filter_mode::serial_in_order,\n",
+    "      /* filter body */\n",
+    "      [&](CaseStringPtr s_ptr) -> void {\n",
+    "        writeCaseString(caseAfterFile, s_ptr);\n",
+    "      }) \n",
+    "  );\n",
+    "}\n",
+    "\n",
+    "static void warmupTBB() {\n",
+    "  int num_threads = std::thread::hardware_concurrency();\n",
+    "  tbb::parallel_for(0, num_threads,\n",
+    "    [](unsigned int) { \n",
+    "      std::this_thread::sleep_for(std::chrono::milliseconds(10)); \n",
+    "  });\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  int num_tokens = std::thread::hardware_concurrency();\n",
+    "  int num_strings = 100; \n",
+    "  int string_len = 100000;\n",
+    "  int free_list_size = num_tokens;\n",
+    "\n",
+    "  std::ofstream caseBeforeFile(\"solutions/solution_pipeline_before.txt\");\n",
+    "  std::ofstream caseAfterFile(\"solutions/solution_pipeline_after.txt\");\n",
+    "  initCaseChange(num_strings, string_len, free_list_size);\n",
+    "\n",
+    "  warmupTBB();\n",
+    "  double parallel_time = 0.0;\n",
+    "  {\n",
+    "    auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "    change_case_parallel(num_tokens, caseBeforeFile, caseAfterFile);\n",
+    "    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();\n",
+    "  }\n",
+    "  std::cout << \"parallel_time == \" << parallel_time << \" seconds\" << std::endl;\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_pipeline-solution.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_pipeline-solution.sh; else ./scripts/run_pipeline-solution.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Next steps\n",
+    "\n",
+    "If you are ready, go to [the next module](../03_oneTBB_flow_graph/oneTBB_flow_graph.ipynb)."
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3.7 (Intel® oneAPI)",
+   "language": "python",
+   "name": "c009-intel_distribution_of_python_3_oneapi-beta05-python"
+  },
+  "language_info": {
+   "codemirror_mode": {
+    "name": "ipython",
+    "version": 3
+   },
+   "file_extension": ".py",
+   "mimetype": "text/x-python",
+   "name": "python",
+   "nbconvert_exporter": "python",
+   "pygments_lexer": "ipython3",
+   "version": "3.7.7"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 4
+}
@@ -0,0 +1,56 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+# Script to submit job in Intel(R) DevCloud
+# Version: 0.7
+#==========================================
+if [ -z "$1" ]; then
+    echo "Missing script argument, Usage: ./q run.sh"
+elif [ ! -f "$1" ]; then
+    echo "File $1 does not exist"
+else
+    echo "Job has been submitted to Intel(R) DevCloud and will execute soon."
+    echo ""
+    echo " If you do not see result in 60 seconds, please restart the Jupyter kernel:"
+    echo " Kernel -> 'Restart Kernel and Clear All Outputs...' and then try again"
+    echo ""
+    script=$1
+    shift
+    args=$@
+    # Remove old output files
+    rm *.sh.* > /dev/null 2>&1
+    # Submit job using qsub
+    qsub_id=`qsub -l nodes=1:ppn=2 -d . $script $@`
+    job_id="$(cut -d'.' -f1 <<<"$qsub_id")"
+    # Print qstat output
+    qstat 
+    # Wait for output file to be generated and display
+    echo ""
+    echo -ne "Waiting for Output "
+    script_name=${script/scripts\//}
+    echo $script_name.o$job_id
+    until [ -f $script_name.o$job_id ]; do
+        sleep 1
+        echo -ne "█"
+        ((timeout++))
+        # Timeout if no output file generated within 60 seconds
+        if [ $timeout == 60 ]; then
+            echo ""
+            echo ""
+            echo "TimeOut 60 seconds: Job is still queued for execution, check for output file later ($script.o$job_id)"
+            echo ""
+            break
+        fi
+    done
+    # Print output and error file content if exist
+    if [ -n "$(find -name '*.sh.o'$job_id)" ]; then
+        echo " Done⬇"
+        cat $script_name.o$job_id
+        cat $script_name.e$job_id
+        echo "Job Completed in $timeout seconds."
+        rm *.sh.*$job_id > /dev/null 2>&1
+    fi
+fi
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction exercise pi-parallel 
+dpcpp -std=c++17 -O2 lab/pi-parallel.cpp -o bin/pi-parallel -tbb
+bin/pi-parallel
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction exercise pi-serial 
+dpcpp -std=c++17 -O2 lab/pi-serial.cpp -o bin/pi-serial
+bin/pi-serial
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction sample pi-parallel 
+dpcpp -std=c++17 -O2 solutions/pi-parallel.cpp -o bin/pi-parallel-solution -tbb
+bin/pi-parallel-solution
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction exercise pipeline-parallel 
+dpcpp -std=c++17 -O2 lab/pipeline-parallel.cpp common/case.cpp -o bin/pipeline-parallel -tbb
+bin/pipeline-parallel
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction exercise pipeline-serial 
+dpcpp -std=c++17 -O2 lab/pipeline-serial.cpp common/case.cpp -o bin/pipeline-serial -tbb
+bin/pipeline-serial
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction exercise pipeline-solution 
+dpcpp -std=c++17 -O2 solutions/pipeline-solution.cpp common/case.cpp -o bin/pipeline-solution -tbb
+bin/pipeline-solution
@@ -0,0 +1 @@
+solutions directory needed for notebook
@@ -0,0 +1,55 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <chrono>
+#include <cmath>
+#include <iostream>
+#include <limits>
+#include <thread>
+#include <tbb/tbb.h>
+
+double calc_pi(int num_intervals) {
+  double dx = 1.0 / num_intervals;
+  double sum = tbb::parallel_reduce(
+    /* range = */ tbb::blocked_range<int>(0, num_intervals ), 
+    /* identity = */ 0.0,
+    /* func */ 
+    [=](const tbb::blocked_range<int>& r, double init) -> double {
+      for (int i = r.begin(); i != r.end(); ++i) {
+        double x = (i+0.5)*dx;
+        double h = std::sqrt(1-x*x);
+        init += h*dx;
+      }
+      return init;
+    },
+    std::plus<double>{}
+  );
+  double pi = 4 * sum;
+  return pi;
+}
+
+static void warmupTBB() {
+  int num_threads = std::thread::hardware_concurrency();
+  tbb::parallel_for(0, num_threads,
+    [](unsigned int) { 
+      std::this_thread::sleep_for(std::chrono::milliseconds(10)); 
+  });
+}
+
+int main() {
+  const int num_intervals = std::numeric_limits<int>::max();
+  double parallel_time = 0.0;
+  warmupTBB();
+  {
+    auto pt0 = std::chrono::high_resolution_clock::now();
+    double pi = calc_pi(num_intervals);
+    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();
+    std::cout << "parallel pi == " << pi << std::endl;
+  }
+
+  std::cout << "parallel_time == " << parallel_time << " seconds" << std::endl;
+  return 0;
+}
@@ -0,0 +1,93 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <algorithm>
+#include <cctype>
+#include <chrono>
+#include <fstream>
+#include <iostream>
+#include <memory>
+#include <string>
+#include <thread>
+
+#include <tbb/tbb.h>
+
+using CaseStringPtr = std::shared_ptr<std::string>;
+
+//
+// These functions are defined in common/case.cpp
+//
+void initCaseChange(int num_strings, int string_len, int free_list_size);
+CaseStringPtr getCaseString(std::ofstream& f); 
+void writeCaseString(std::ofstream& f, CaseStringPtr s);
+
+void change_case_parallel(int num_tokens, std::ofstream& caseBeforeFile, std::ofstream& caseAfterFile) {
+  tbb::parallel_pipeline(
+    /* tokens */ num_tokens,
+    /* the get filter */
+    tbb::make_filter<void, CaseStringPtr>(
+      /* filter node */ tbb::filter_mode::serial_in_order,
+      /* filter body */
+      [&](tbb::flow_control& fc) -> CaseStringPtr {
+        CaseStringPtr s_ptr = getCaseString(caseBeforeFile);
+        if (!s_ptr) 
+          fc.stop();
+        return s_ptr; 
+      }) & // concatenation operation
+    /* make the change case filter */
+    tbb::make_filter<CaseStringPtr, CaseStringPtr>(
+      /* filter node */ tbb::filter_mode::parallel,
+      /* filter body */
+      [](CaseStringPtr s_ptr) -> CaseStringPtr {
+        std::transform(s_ptr->begin(), s_ptr->end(), s_ptr->begin(), 
+          [](char c) -> char {
+            if (std::islower(c))
+              return std::toupper(c);
+            else if (std::isupper(c))
+              return std::tolower(c);
+            else
+              return c;
+          });
+        return s_ptr;
+      }) & // concatenation operation
+    /* make the write filter */
+    tbb::make_filter<CaseStringPtr, void>(
+      /* filter node */ tbb::filter_mode::serial_in_order,
+      /* filter body */
+      [&](CaseStringPtr s_ptr) -> void {
+        writeCaseString(caseAfterFile, s_ptr);
+      }) 
+  );
+}
+
+static void warmupTBB() {
+  int num_threads = std::thread::hardware_concurrency();
+  tbb::parallel_for(0, num_threads,
+    [](unsigned int) { 
+      std::this_thread::sleep_for(std::chrono::milliseconds(10)); 
+  });
+}
+
+int main() {
+  int num_tokens = std::thread::hardware_concurrency();
+  int num_strings = 100; 
+  int string_len = 100000;
+  int free_list_size = num_tokens;
+
+  std::ofstream caseBeforeFile("solutions/solution_pipeline_before.txt");
+  std::ofstream caseAfterFile("solutions/solution_pipeline_after.txt");
+  initCaseChange(num_strings, string_len, free_list_size);
+
+  warmupTBB();
+  double parallel_time = 0.0;
+  {
+    auto pt0 = std::chrono::high_resolution_clock::now();
+    change_case_parallel(num_tokens, caseBeforeFile, caseAfterFile);
+    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();
+  }
+  std::cout << "parallel_time == " << parallel_time << " seconds" << std::endl;
+  return 0;
+}
@@ -0,0 +1,83 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <chrono>
+#include <iostream>
+#include <thread>
+#include <vector>
+
+#include <tbb/tbb.h>
+
+
+// defined in common/fwd_sub.cpp
+std::vector<double> init_fwd_sub(std::vector<double>& x,
+                                 std::vector<double>& a,
+                                 std::vector<double>& b); 
+
+void check_fwd_sub(const std::vector<double>& x,
+                   const std::vector<double>& x_gold); 
+
+using Node = tbb::flow::continue_node<tbb::flow::continue_msg>;
+using NodePtr = std::shared_ptr<Node>;
+
+NodePtr create_node(tbb::flow::graph& g, int r, int c, int block_size, 
+                   std::vector<double>& x, const std::vector<double>& a, 
+                   std::vector<double>& b) {
+
+  const int N = x.size();
+  return std::make_shared<Node>(g,
+    [r, c, block_size, N, &x, &a, &b] (const tbb::flow::continue_msg& msg) {
+      // STEP A: Add node body
+      return msg;
+    }
+  );
+}
+
+void add_edges(std::vector<NodePtr>& nodes, int r, int c, int block_size, int num_blocks) {
+  NodePtr np = nodes[r*num_blocks + c];
+  if (c + 1 < num_blocks && r != c)
+    tbb::flow::make_edge(*np, *nodes[r*num_blocks + c + 1]);
+  if (r + 1 < num_blocks)
+    tbb::flow::make_edge(*np, *nodes[(r + 1)*num_blocks + c]);
+}
+
+void fwd_sub_parallel(std::vector<double>& x, const std::vector<double>& a, 
+              std::vector<double>& b) {
+  const int N = x.size();
+  const int block_size = 1024;
+  const int num_blocks = N / block_size;
+
+  std::vector<NodePtr> nodes(num_blocks*num_blocks);
+  tbb::flow::graph g;
+  for (int r = num_blocks - 1; r >= 0; --r) {
+    for (int c = r; c >= 0; --c) {
+      nodes[r*num_blocks + c] = create_node(g, r, c, block_size, x, a, b);;
+      // STEP B: Add a call to function add_edges
+    }
+  }
+  // STEP C: send a tbb::flow::continue_msg{} to nodes[0]
+  // STEP D: wait for the graph to complete
+}
+
+int main() {
+  const int N = 32768;
+
+  std::vector<double> a(N*N);
+  std::vector<double> b(N);
+  std::vector<double> x(N);
+
+  auto x_gold = init_fwd_sub(x,a,b);
+
+  double parallel_time = 0.0;
+  {
+    auto pt0 = std::chrono::high_resolution_clock::now();
+    fwd_sub_parallel(x,a,b);
+    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();
+  }
+  check_fwd_sub(x, x_gold);
+  std::cout << "parallel_time == " << parallel_time << " seconds" << std::endl;
+  return 0;
+}
@@ -0,0 +1,61 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <chrono>
+#include <iostream>
+#include <thread>
+#include <vector>
+
+
+// defined in common/fwd_sub.cpp
+std::vector<double> init_fwd_sub(std::vector<double>& x,
+                                 std::vector<double>& a,
+                                 std::vector<double>& b); 
+
+void check_fwd_sub(const std::vector<double>& x,
+                   const std::vector<double>& x_gold); 
+        
+void fwd_sub_serial(std::vector<double>& x, const std::vector<double>& a, std::vector<double>& b) {
+  const int N = x.size();
+  const int block_size = 512;
+  const int num_blocks = N / block_size;
+
+  for ( int r = 0; r < num_blocks; ++r ) {
+    for ( int c = 0; c <= r; ++c ) {
+      int i_start = r*block_size, i_end = i_start + block_size;
+      int j_start = c*block_size, j_max = j_start + block_size - 1;
+      for (int i = i_start; i < i_end; ++i) {
+        int j_end = (i <= j_max) ? i : j_max + 1;
+        for (int j = j_start; j < j_end; ++j) {
+          b[i] -= a[j + i*N] * x[j];
+        }
+        if (j_end == i) {
+          x[i] = b[i] / a[i + i*N];
+        }
+      }
+    }
+  }
+}
+
+int main() {
+  const int N = 32768;
+
+  std::vector<double> a(N*N);
+  std::vector<double> b(N);
+  std::vector<double> x(N);
+
+  auto x_gold = init_fwd_sub(x,a,b);
+
+  double serial_time = 0.0;
+  {
+    auto st0 = std::chrono::high_resolution_clock::now();
+    fwd_sub_serial(x,a,b);
+    serial_time = 1e-9*(std::chrono::high_resolution_clock::now() - st0).count();
+  }
+  check_fwd_sub(x, x_gold);
+  std::cout << "serial_time == " << serial_time << " seconds" << std::endl;
+  return 0;
+}
@@ -0,0 +1 @@
+bin directory needed for notebook
@@ -0,0 +1,44 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <vector>
+#include <iostream>
+
+std::vector<double> init_fwd_sub(std::vector<double>& x, 
+                                 std::vector<double>& a, 
+                                 std::vector<double>& b) {
+  const int N = x.size();
+  for (int i = 0; i < N; ++i) {
+    x[i] = 0;
+    b[i] = i*i;
+    for (int j = 0; j <= i; ++j) {
+      a[j + i*N] = 1 + j*i;
+    }
+  }
+
+  std::vector<double> b_tmp = b;
+  std::vector<double> x_gold = x;
+  for (int i = 0; i < N; ++i) {
+    for (int j = 0; j < i; ++j) {
+      b_tmp[i] -= a[j + i*N] * x_gold[j];
+    }
+    x_gold[i] = b_tmp[i] / a[i + i*N];
+  }
+  return x_gold;
+}
+
+void check_fwd_sub(const std::vector<double>& x,
+                   const std::vector<double>& x_gold) {
+  if (x.size() != x_gold.size()) {
+    std::cerr << "ERROR: Invalid result" << std::endl;
+  }
+  for (int i = 0; i < x.size(); ++i) {
+    if (x[i] > 1.1*x_gold[i] || x[i] < 0.9*x_gold[i]) {
+      std::cerr << "ERROR: Invalid result" 
+                << "  at " << i << " " << x[i] << " != " << x_gold[i] << std::endl;
+    }
+  }
+}  
@@ -0,0 +1 @@
+lab directory needed for notebook
@@ -0,0 +1,472 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# oneTBB flow graph"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "##### Sections\n",
+    "- [oneTBB Flow Graph](#oneTBB-Flow-Graph)\n",
+    "- _Code_: [Forward substitution using a oneTBB flow graph](#Forward-substitution-using-a-oneTBB-flow-graph)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Learning Objectives\n",
+    "* Understand the basic steps of constructing and executing a oneTBB flow graph"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# oneTBB Flow Graph\n",
+    "\n",
+    "The oneTBB library provides functions and classes for applications that can be expressed as graphs\n",
+    "of computations. These classes and functions are in namespace ``tbb::flow``. A flow graph is used when you want to be precise about the dependencies in your code, or if you have a streaming\n",
+    "application that requires more than just the simple linear pipeline supported by ``tbb::parallel_pipeline``.\n",
+    "You can read more about the oneTBB flow graph interfaces in the \n",
+    "[flow graph](https://spec.oneapi.com/versions/latest/elements/oneTBB/source/flow_graph.html) section of \n",
+    "the oneAPI specification.  \n",
+    "\n",
+    "## Forward substitution using a oneTBB flow graph\n",
+    "\n",
+    "Here we implement the forward substitution method of solving a set of \n",
+    "equations ``Ax = b``, where ``A`` is an NxN lower triangular matrix. To do so, we create a\n",
+    "simple flow graph example that uses only a single node type, ``continue_node``. \n",
+    "\n",
+    "As shown below in the diagram of the dependencies in forward substitution, once the value of \n",
+    "element a<sub>i</sub><sub>,j</sub> is known, both a<sub>i+1</sub><sub>,j</sub> and a<sub>i</sub><sub>,j+1</sub>\n",
+    "can be solved for.  These dependencies are shown in figure (a) below. To reduce scheduling overheads in\n",
+    "parallel implementations, we can aggregate elements into blocks to schedule as a task, where once the \n",
+    "elements in block i,j are available, then all the elements in blocks i+1,j and i,j+1 can be solved for."
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "![A tiled dependency graph implementation of forward substitution](img/fwd_sub.png)\n",
+    "\n",
+    "### Run the sequential baseline implementation\n",
+    "\n",
+    "Inspect the sequential implementation of forward substitution below - there are no modifications necessary. This code implements \n",
+    "a serial blocked implementation of forward substitution. Run the first cell to create the file, then run the cell below it to compile \n",
+    "and execute the code. This represents the baseline for the computation that we will convert into a oneTBB flow graph.\n",
+    "\n",
+    "1. Inspect the code cell below, then click run ▶ to save the code to a file\n",
+    "2. Run ▶ the cell in the __Build and Run the baseline__ section below the code snippet to compile and execute the code in the saved file\n",
+    "3. Inspect the two new files ``serial_pipeline_before.txt`` and ``serial_pipeline_after.txt`` to see the results of changing the case of the text."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/fwd-sub-serial.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <iostream>\n",
+    "#include <thread>\n",
+    "#include <vector>\n",
+    "\n",
+    "\n",
+    "// defined in common/fwd_sub.cpp\n",
+    "std::vector<double> init_fwd_sub(std::vector<double>& x,\n",
+    "                                 std::vector<double>& a,\n",
+    "                                 std::vector<double>& b); \n",
+    "\n",
+    "void check_fwd_sub(const std::vector<double>& x,\n",
+    "                   const std::vector<double>& x_gold); \n",
+    "        \n",
+    "void fwd_sub_serial(std::vector<double>& x, const std::vector<double>& a, std::vector<double>& b) {\n",
+    "  const int N = x.size();\n",
+    "  const int block_size = 512;\n",
+    "  const int num_blocks = N / block_size;\n",
+    "\n",
+    "  for ( int r = 0; r < num_blocks; ++r ) {\n",
+    "    for ( int c = 0; c <= r; ++c ) {\n",
+    "      int i_start = r*block_size, i_end = i_start + block_size;\n",
+    "      int j_start = c*block_size, j_max = j_start + block_size - 1;\n",
+    "      for (int i = i_start; i < i_end; ++i) {\n",
+    "        int j_end = (i <= j_max) ? i : j_max + 1;\n",
+    "        for (int j = j_start; j < j_end; ++j) {\n",
+    "          b[i] -= a[j + i*N] * x[j];\n",
+    "        }\n",
+    "        if (j_end == i) {\n",
+    "          x[i] = b[i] / a[i + i*N];\n",
+    "        }\n",
+    "      }\n",
+    "    }\n",
+    "  }\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  const int N = 32768;\n",
+    "\n",
+    "  std::vector<double> a(N*N);\n",
+    "  std::vector<double> b(N);\n",
+    "  std::vector<double> x(N);\n",
+    "\n",
+    "  auto x_gold = init_fwd_sub(x,a,b);\n",
+    "\n",
+    "  double serial_time = 0.0;\n",
+    "  {\n",
+    "    auto st0 = std::chrono::high_resolution_clock::now();\n",
+    "    fwd_sub_serial(x,a,b);\n",
+    "    serial_time = 1e-9*(std::chrono::high_resolution_clock::now() - st0).count();\n",
+    "  }\n",
+    "  check_fwd_sub(x, x_gold);\n",
+    "  std::cout << \"serial_time == \" << serial_time << \" seconds\" << std::endl;\n",
+    "  return 0;\n",
+    "}\n"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and Run the baseline code\n",
+    "\n",
+    "Select the cell below and click Run ▶ to compile and execute the code that you modified above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_fwd-sub-serial.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_fwd-sub-serial.sh; else ./scripts/run_fwd-sub-serial.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Implement a parallel version with tbb::flow::graph\n",
+    "\n",
+    "In this section, you will use a ``tbb::flow::graph`` to create a parallel implementation of forward substitution. The classes and functions\n",
+    "for creating oneTBB flow graphs are in namespace ``tbb::flow``.\n",
+    "\n",
+    "There are five general steps for creating and running a graph.\n",
+    "\n",
+    "1. Create an instance of class ``tbb::flow::graph``. In our example, we use an object named ``g``.\n",
+    "2. Create the node objects. In this example, we will only use ``tbb::flow::continue_node`` objects. One object will be created per block of elements.\n",
+    "3. Make edges between the nodes from predecessor to successor using ``tbb::flow::make_edge(predecessor_node, successor_node)``. In this example, each node will be connected to nodes i+1,j and i,j+1. This enforces the ordering necessary for a legal execution of forward substitution.\n",
+    "4. Start the activity in the graph. In this example, we will invoke try_put on the top corner of the graph, i.e. \"nodes[0]->try_put(tbb::flow::continue_msg{})``.\n",
+    "5. Wait for the graph to complete, by calling ``g.wait_for_all()``.\n",
+    "\n",
+    "You can find more information about tbb::flow::continue_node [here](https://spec.oneapi.com/versions/latest/elements/oneTBB/source/flow_graph/continue_node_cls.html). The key points for our example\n",
+    "are:\n",
+    "\n",
+    "The constructor arguments:\n",
+    "\n",
+    "```cpp\n",
+    "template<typename Body>\n",
+    "continue_node( graph &g, Body body );\n",
+    "```\n",
+    "\n",
+    "And the [requirements on the body](https://spec.oneapi.com/versions/latest/elements/oneTBB/source/named_requirements/flow_graph/continue_node_body.html):\n",
+    "\n",
+    "```cpp\n",
+    "continue_msg Body::operator()(const continue_msg &v) const\n",
+    "```\n",
+    "\n",
+    "A `continue_msg` is an empty message, differentiating it from other value-carrying messages supported by the flow graph \n",
+    "interfaces. \n",
+    "\n",
+    "For this exercise, complete the following steps:\n",
+    "\n",
+    "1. Inspect the code cell below and make the following modifications.\n",
+    "  1. Add the body to the ``continue_node`` created in function ``create_node``. The function ``create_node`` is called repeatedly in our example.\n",
+    "  2. Add a call to create_edges in the loop in function ``fwd_sub_parallel``.\n",
+    "  3. Invoke ``nodes[0]->try_put(tbb::flow::continue_msg{})`` to start the graph.\n",
+    "  4. Add a call to ``g.wait_for_all()`` after the graph has been started.\n",
+    "2. When the modifications are complete, click run ▶ to save the code to a file.\n",
+    "3. Run ▶ the cell in the __Build and Run the modified code__ section below the code snippet to compile and execute the code in the saved file."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/fwd-sub-parallel.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <iostream>\n",
+    "#include <thread>\n",
+    "#include <vector>\n",
+    "\n",
+    "#include <tbb/tbb.h>\n",
+    "\n",
+    "\n",
+    "// defined in common/fwd_sub.cpp\n",
+    "std::vector<double> init_fwd_sub(std::vector<double>& x,\n",
+    "                                 std::vector<double>& a,\n",
+    "                                 std::vector<double>& b); \n",
+    "\n",
+    "void check_fwd_sub(const std::vector<double>& x,\n",
+    "                   const std::vector<double>& x_gold); \n",
+    "\n",
+    "using Node = tbb::flow::continue_node<tbb::flow::continue_msg>;\n",
+    "using NodePtr = std::shared_ptr<Node>;\n",
+    "\n",
+    "NodePtr create_node(tbb::flow::graph& g, int r, int c, int block_size, \n",
+    "                   std::vector<double>& x, const std::vector<double>& a, \n",
+    "                   std::vector<double>& b) {\n",
+    "\n",
+    "  const int N = x.size();\n",
+    "  return std::make_shared<Node>(g,\n",
+    "    [r, c, block_size, N, &x, &a, &b] (const tbb::flow::continue_msg& msg) {\n",
+    "      // STEP A: Add node body\n",
+    "      return msg;\n",
+    "    }\n",
+    "  );\n",
+    "}\n",
+    "\n",
+    "void add_edges(std::vector<NodePtr>& nodes, int r, int c, int block_size, int num_blocks) {\n",
+    "  NodePtr np = nodes[r*num_blocks + c];\n",
+    "  if (c + 1 < num_blocks && r != c)\n",
+    "    tbb::flow::make_edge(*np, *nodes[r*num_blocks + c + 1]);\n",
+    "  if (r + 1 < num_blocks)\n",
+    "    tbb::flow::make_edge(*np, *nodes[(r + 1)*num_blocks + c]);\n",
+    "}\n",
+    "\n",
+    "void fwd_sub_parallel(std::vector<double>& x, const std::vector<double>& a, \n",
+    "              std::vector<double>& b) {\n",
+    "  const int N = x.size();\n",
+    "  const int block_size = 1024;\n",
+    "  const int num_blocks = N / block_size;\n",
+    "\n",
+    "  std::vector<NodePtr> nodes(num_blocks*num_blocks);\n",
+    "  tbb::flow::graph g;\n",
+    "  for (int r = num_blocks - 1; r >= 0; --r) {\n",
+    "    for (int c = r; c >= 0; --c) {\n",
+    "      nodes[r*num_blocks + c] = create_node(g, r, c, block_size, x, a, b);;\n",
+    "      // STEP B: Add a call to function add_edges\n",
+    "    }\n",
+    "  }\n",
+    "  // STEP C: send a tbb::flow::continue_msg{} to nodes[0]\n",
+    "  // STEP D: wait for the graph to complete\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  const int N = 32768;\n",
+    "\n",
+    "  std::vector<double> a(N*N);\n",
+    "  std::vector<double> b(N);\n",
+    "  std::vector<double> x(N);\n",
+    "\n",
+    "  auto x_gold = init_fwd_sub(x,a,b);\n",
+    "\n",
+    "  double parallel_time = 0.0;\n",
+    "  {\n",
+    "    auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "    fwd_sub_parallel(x,a,b);\n",
+    "    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();\n",
+    "  }\n",
+    "  check_fwd_sub(x, x_gold);\n",
+    "  std::cout << \"parallel_time == \" << parallel_time << \" seconds\" << std::endl;\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and Run the modified code\n",
+    "\n",
+    "Select the cell below and click Run ▶ to compile and execute the code that you modified above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_fwd-sub-parallel.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_fwd-sub-parallel.sh; else ./scripts/run_fwd-sub-parallel.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Forward Substitution Solution (Don't peak unless you have to)"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "%%writefile solutions/fwd-sub-solution.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <chrono>\n",
+    "#include <iostream>\n",
+    "#include <thread>\n",
+    "#include <vector>\n",
+    "\n",
+    "#include <tbb/tbb.h>\n",
+    "\n",
+    "\n",
+    "// defined in common/fwd_sub.cpp\n",
+    "std::vector<double> init_fwd_sub(std::vector<double>& x,\n",
+    "                                 std::vector<double>& a,\n",
+    "                                 std::vector<double>& b); \n",
+    "\n",
+    "void check_fwd_sub(const std::vector<double>& x,\n",
+    "                   const std::vector<double>& x_gold); \n",
+    "\n",
+    "using Node = tbb::flow::continue_node<tbb::flow::continue_msg>;\n",
+    "using NodePtr = std::shared_ptr<Node>;\n",
+    "\n",
+    "NodePtr create_node(tbb::flow::graph& g, int r, int c, int block_size, \n",
+    "                   std::vector<double>& x, const std::vector<double>& a, \n",
+    "                   std::vector<double>& b) {\n",
+    "\n",
+    "  const int N = x.size();\n",
+    "  return std::make_shared<Node>(g,\n",
+    "    [r, c, block_size, N, &x, &a, &b] (const tbb::flow::continue_msg& msg) {\n",
+    "      int i_start = r*block_size, i_end = i_start + block_size;\n",
+    "      int j_start = c*block_size, j_max = j_start + block_size - 1;\n",
+    "      for (int i = i_start; i < i_end; ++i) {\n",
+    "        int j_end = (i <= j_max) ? i : j_max + 1;\n",
+    "        for (int j = j_start; j < j_end; ++j) {\n",
+    "          b[i] -= a[j + i*N] * x[j];\n",
+    "        }\n",
+    "        if (j_end == i) {\n",
+    "          x[i] = b[i] / a[i + i*N];\n",
+    "        }\n",
+    "      }\n",
+    "      return msg;\n",
+    "    }\n",
+    "  );\n",
+    "}\n",
+    "\n",
+    "void add_edges(std::vector<NodePtr>& nodes, int r, int c, int block_size, int num_blocks) {\n",
+    "  NodePtr np = nodes[r*num_blocks + c];\n",
+    "  if (c + 1 < num_blocks && r != c)\n",
+    "    tbb::flow::make_edge(*np, *nodes[r*num_blocks + c + 1]);\n",
+    "  if (r + 1 < num_blocks)\n",
+    "    tbb::flow::make_edge(*np, *nodes[(r + 1)*num_blocks + c]);\n",
+    "}\n",
+    "\n",
+    "void fwd_sub_parallel(std::vector<double>& x, const std::vector<double>& a, \n",
+    "              std::vector<double>& b) {\n",
+    "  const int N = x.size();\n",
+    "  const int block_size = 1024;\n",
+    "  const int num_blocks = N / block_size;\n",
+    "\n",
+    "  std::vector<NodePtr> nodes(num_blocks*num_blocks);\n",
+    "  tbb::flow::graph g;\n",
+    "  for (int r = num_blocks - 1; r >= 0; --r) {\n",
+    "    for (int c = r; c >= 0; --c) {\n",
+    "      nodes[r*num_blocks + c] = create_node(g, r, c, block_size, x, a, b);;\n",
+    "      add_edges(nodes, r, c, block_size, num_blocks);\n",
+    "    }\n",
+    "  }\n",
+    "  nodes[0]->try_put(tbb::flow::continue_msg{});\n",
+    "  g.wait_for_all();\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  const int N = 32768;\n",
+    "\n",
+    "  std::vector<double> a(N*N);\n",
+    "  std::vector<double> b(N);\n",
+    "  std::vector<double> x(N);\n",
+    "\n",
+    "  auto x_gold = init_fwd_sub(x,a,b);\n",
+    "\n",
+    "  double parallel_time = 0.0;\n",
+    "  {\n",
+    "    auto pt0 = std::chrono::high_resolution_clock::now();\n",
+    "    fwd_sub_parallel(x,a,b);\n",
+    "    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();\n",
+    "  }\n",
+    "  check_fwd_sub(x, x_gold);\n",
+    "  std::cout << \"parallel_time == \" << parallel_time << \" seconds\" << std::endl;\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_fwd-sub-solution.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_fwd-sub-solution.sh; else ./scripts/run_fwd-sub-solution.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Next steps\n",
+    "\n",
+    "If you are ready, go to [the next module](../04_oneTBB_concurrent_containers/oneTBB_concurrent_containers.ipynb)."
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3.7 (Intel® oneAPI)",
+   "language": "python",
+   "name": "c009-intel_distribution_of_python_3_oneapi-beta05-python"
+  },
+  "language_info": {
+   "codemirror_mode": {
+    "name": "ipython",
+    "version": 3
+   },
+   "file_extension": ".py",
+   "mimetype": "text/x-python",
+   "name": "python",
+   "nbconvert_exporter": "python",
+   "pygments_lexer": "ipython3",
+   "version": "3.7.7"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 4
+}
@@ -0,0 +1,56 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+# Script to submit job in Intel(R) DevCloud
+# Version: 0.7
+#==========================================
+if [ -z "$1" ]; then
+    echo "Missing script argument, Usage: ./q run.sh"
+elif [ ! -f "$1" ]; then
+    echo "File $1 does not exist"
+else
+    echo "Job has been submitted to Intel(R) DevCloud and will execute soon."
+    echo ""
+    echo " If you do not see result in 60 seconds, please restart the Jupyter kernel:"
+    echo " Kernel -> 'Restart Kernel and Clear All Outputs...' and then try again"
+    echo ""
+    script=$1
+    shift
+    args=$@
+    # Remove old output files
+    rm *.sh.* > /dev/null 2>&1
+    # Submit job using qsub
+    qsub_id=`qsub -l nodes=1:ppn=2 -d . $script $@`
+    job_id="$(cut -d'.' -f1 <<<"$qsub_id")"
+    # Print qstat output
+    qstat 
+    # Wait for output file to be generated and display
+    echo ""
+    echo -ne "Waiting for Output "
+    script_name=${script/scripts\//}
+    echo $script_name.o$job_id
+    until [ -f $script_name.o$job_id ]; do
+        sleep 1
+        echo -ne "█"
+        ((timeout++))
+        # Timeout if no output file generated within 60 seconds
+        if [ $timeout == 60 ]; then
+            echo ""
+            echo ""
+            echo "TimeOut 60 seconds: Job is still queued for execution, check for output file later ($script.o$job_id)"
+            echo ""
+            break
+        fi
+    done
+    # Print output and error file content if exist
+    if [ -n "$(find -name '*.sh.o'$job_id)" ]; then
+        echo " Done⬇"
+        cat $script_name.o$job_id
+        cat $script_name.e$job_id
+        echo "Job Completed in $timeout seconds."
+        rm *.sh.*$job_id > /dev/null 2>&1
+    fi
+fi
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction exercise fwd-sub-parallel 
+dpcpp -std=c++17 -O2 lab/fwd-sub-parallel.cpp common/fwd_sub.cpp -o bin/fwd-sub-parallel -tbb
+bin/fwd-sub-parallel
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction exercise fwd-sub-serial 
+dpcpp -std=c++17 -O2 lab/fwd-sub-serial.cpp common/fwd_sub.cpp -o bin/fwd-sub-serial -tbb
+bin/fwd-sub-serial
@@ -0,0 +1,11 @@
+#!/bin/bash
+#==========================================
+# Copyright (c) 2020 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+#==========================================
+
+source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
+/bin/echo "##" $(whoami) is compiling oneTBB Introduction exercise fwd-sub-solution 
+dpcpp -std=c++17 -O2 solutions/fwd-sub-solution.cpp common/fwd_sub.cpp -o bin/fwd-sub-solution -tbb
+bin/fwd-sub-solution
@@ -0,0 +1 @@
+solutions directory needed for notebook
@@ -0,0 +1,93 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <chrono>
+#include <iostream>
+#include <thread>
+#include <vector>
+
+#include <tbb/tbb.h>
+
+
+// defined in common/fwd_sub.cpp
+std::vector<double> init_fwd_sub(std::vector<double>& x,
+                                 std::vector<double>& a,
+                                 std::vector<double>& b); 
+
+void check_fwd_sub(const std::vector<double>& x,
+                   const std::vector<double>& x_gold); 
+
+using Node = tbb::flow::continue_node<tbb::flow::continue_msg>;
+using NodePtr = std::shared_ptr<Node>;
+
+NodePtr create_node(tbb::flow::graph& g, int r, int c, int block_size, 
+                   std::vector<double>& x, const std::vector<double>& a, 
+                   std::vector<double>& b) {
+
+  const int N = x.size();
+  return std::make_shared<Node>(g,
+    [r, c, block_size, N, &x, &a, &b] (const tbb::flow::continue_msg& msg) {
+      int i_start = r*block_size, i_end = i_start + block_size;
+      int j_start = c*block_size, j_max = j_start + block_size - 1;
+      for (int i = i_start; i < i_end; ++i) {
+        int j_end = (i <= j_max) ? i : j_max + 1;
+        for (int j = j_start; j < j_end; ++j) {
+          b[i] -= a[j + i*N] * x[j];
+        }
+        if (j_end == i) {
+          x[i] = b[i] / a[i + i*N];
+        }
+      }
+      return msg;
+    }
+  );
+}
+
+void add_edges(std::vector<NodePtr>& nodes, int r, int c, int block_size, int num_blocks) {
+  NodePtr np = nodes[r*num_blocks + c];
+  if (c + 1 < num_blocks && r != c)
+    tbb::flow::make_edge(*np, *nodes[r*num_blocks + c + 1]);
+  if (r + 1 < num_blocks)
+    tbb::flow::make_edge(*np, *nodes[(r + 1)*num_blocks + c]);
+}
+
+void fwd_sub_parallel(std::vector<double>& x, const std::vector<double>& a, 
+              std::vector<double>& b) {
+  const int N = x.size();
+  const int block_size = 1024;
+  const int num_blocks = N / block_size;
+
+  std::vector<NodePtr> nodes(num_blocks*num_blocks);
+  tbb::flow::graph g;
+  for (int r = num_blocks - 1; r >= 0; --r) {
+    for (int c = r; c >= 0; --c) {
+      nodes[r*num_blocks + c] = create_node(g, r, c, block_size, x, a, b);;
+      add_edges(nodes, r, c, block_size, num_blocks);
+    }
+  }
+  nodes[0]->try_put(tbb::flow::continue_msg{});
+  g.wait_for_all();
+}
+
+int main() {
+  const int N = 32768;
+
+  std::vector<double> a(N*N);
+  std::vector<double> b(N);
+  std::vector<double> x(N);
+
+  auto x_gold = init_fwd_sub(x,a,b);
+
+  double parallel_time = 0.0;
+  {
+    auto pt0 = std::chrono::high_resolution_clock::now();
+    fwd_sub_parallel(x,a,b);
+    parallel_time = 1e-9*(std::chrono::high_resolution_clock::now() - pt0).count();
+  }
+  check_fwd_sub(x, x_gold);
+  std::cout << "parallel_time == " << parallel_time << " seconds" << std::endl;
+  return 0;
+}
@@ -0,0 +1,63 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <iostream>
+#include <queue>
+#include <thread>
+#include <tbb/tbb.h>
+
+#define INCORRECT_QUEUE_TYPE std::priority_queue<int>;
+
+// STEP A: Replace std::priority_queue with tbb::concurrent_priority_queue
+using queue_type = INCORRECT_QUEUE_TYPE;
+        
+int consume(queue_type &myPQ) {
+  int local_sum = 0;
+  int value = -1;
+    
+  bool consuming = true;
+  while(consuming) {
+    // STEP B: replace the pattern of empty, top and pop 
+    //         with a single call to myPQ.try_pop(value)
+    if (!myPQ.empty()) {
+      value = myPQ.top();
+      myPQ.pop();
+      if (value == -1)
+        consuming = false;
+      else
+        local_sum += value;
+    }
+  }
+  return local_sum;
+}
+
+int main() {
+  int sum (0);
+  int item;
+
+  queue_type myPQ;
+ 
+  std::thread producer([&]() {
+    for(int i=0; i<10001; i+=1) {
+      myPQ.push(i);
+    }
+    // to signal the end to the two consumers
+    myPQ.push(-1);
+    myPQ.push(-1);
+  });
+
+  int local_sum1 = 0, local_sum2 = 0;
+  std::thread consumer1([&]() { local_sum1 = consume(myPQ); });
+  std::thread consumer2([&]() { local_sum2 = consume(myPQ); });
+
+  producer.join();
+  consumer1.join();
+  consumer2.join();
+    
+  // prints "total: 50005000" (for 0,10001,1)
+  std::cout << "total: " << local_sum1 + local_sum2 << '\n';
+  return 0;
+}
@@ -0,0 +1,28 @@
+//==============================================================
+// Copyright (c) 2020 Intel Corporation
+//
+// SPDX-License-Identifier: Apache-2.0
+// =============================================================
+
+#include <iostream>
+#include <queue>
+
+int main() {
+  int sum (0);
+  int item;
+
+  std::priority_queue<int> myPQ;
+ 
+  for(int i=0; i<10001; i+=1) {
+    myPQ.push(i);
+  }
+ 
+  while( !myPQ.empty() ) {
+    sum += myPQ.top();
+    myPQ.pop();
+  }
+
+  // prints "total: 50005000" (for 0,10001,1)
+  std::cout << "total: " << sum << '\n';
+  return 0;
+}
@@ -0,0 +1 @@
+bin directory needed for notebook
@@ -0,0 +1 @@
+lab directory needed for notebook
@@ -0,0 +1,368 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# oneTBB Concurrent Containers"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "##### Sections\n",
+    "- [oneTBB Concurrent Containers](#oneTBB-Concurrent-Containers)\n",
+    "- _Code_: [A Producer-Consumer Application with tbb::concurrent_queue](#A-Producer-Consumer-Application-with-tbb::concurrent_queue)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Learning Objectives\n",
+    "* Learn how thread-unsafe uses of a standard container might be addressed by using a oneTBB concurrent container"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# oneTBB Concurrent Containers"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "The oneTBB library provides a number of \n",
+    "[concurrent containers](https://spec.oneapi.com/versions/latest/elements/oneTBB/source/containers.html),\n",
+    "including `concurrent_vector`, `concurrent_queue`, `concurrent_bounded_queue`, `concurrent_priority_queue`,\n",
+    "`concurrent_hash_map`, `concurrent_unordered_map` and more. These classes provide optimized containers that \n",
+    "permit multiple threads to simultaneously invoke certain functions on the same container."
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## A Producer-Consumer Application with tbb::concurrent_priority_queue\n",
+    "\n",
+    "In this section, we will implement a very simple producer-consumer application that uses a shared \n",
+    "priority queue. A producer thread will generate 1000 items, putting them into a queue and a consumer \n",
+    "thread will read the items from the queue, generating a sum of the values as a result.\n",
+    "\n",
+    "### Run the sequential baseline implementation\n",
+    "\n",
+    "In our baseline implementation, the producer thread and the consumer thread are the same thread, the main thread.\n",
+    "There is no concurrency expressed in our serial code, but if there was, we would have a race condition since we\n",
+    "use a `std::priority_queue`. In the `while`-loop, we call `empty` before calling `pop`. If more than one thread\n",
+    "was simultaneously popping from this shared queue, the queue might appear non-empty at the test in the `while`-loop\n",
+    "but be empty by the time `pop` is called. This problem does not exist in our serial code since only one thread is\n",
+    "used. In the next section, we will add concurrency and the potential for such a race condition.\n",
+    "\n",
+    "In this section, just inspect the sequential code below - there are no modifications necessary. Run the first cell \n",
+    "to create the file, then run the cell below it to compile and execute the code. This represents the baseline for \n",
+    "our producer-consumer exercise.\n",
+    "\n",
+    "1. Inspect the code cell below, then click run ▶ to save the code to a file\n",
+    "2. Run ▶ the cell in the __Build and Run the baseline__ section below the code snippet to compile and execute the code in the saved file"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/q-serial.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <iostream>\n",
+    "#include <queue>\n",
+    "\n",
+    "int main() {\n",
+    "  int sum (0);\n",
+    "  int item;\n",
+    "\n",
+    "  std::priority_queue<int> myPQ;\n",
+    " \n",
+    "  for(int i=0; i<10001; i+=1) {\n",
+    "    myPQ.push(i);\n",
+    "  }\n",
+    " \n",
+    "  while( !myPQ.empty() ) {\n",
+    "    sum += myPQ.top();\n",
+    "    myPQ.pop();\n",
+    "  }\n",
+    "\n",
+    "  // prints \"total: 50005000\" (for 0,10001,1)\n",
+    "  std::cout << \"total: \" << sum << '\\n';\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and Run the baseline\n",
+    "Select the cell below and click Run ▶ to compile and execute the code above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_q-serial.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_q-serial.sh; else ./scripts/run_q-serial.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Implement a parallel version with tbb::concurrent_priority_queue\n",
+    "\n",
+    "In this section, we modify the example to create one producer thread\n",
+    "and two consumer threads that run concurrently. To eliminate the potential\n",
+    "race on the call to `empty` and `pop`, we replace the `std::priority_queue` \n",
+    "with `tbb::concurrent_priority_queue`. In fact, if you run the code below before\n",
+    "making the required modifications, the output might include a segfault.\n",
+    "\n",
+    "The key interfaces in `tbb::concurrent_priority_queue` needed for this exercise \n",
+    "are reproduced below:\n",
+    "\n",
+    "```cpp\n",
+    "namespace tbb {\n",
+    "  template <typename T, typename Compare = std::less<T>,\n",
+    "            typename Allocator = cache_aligned_allocator<T>>\n",
+    "  class concurrent_priority_queue {\n",
+    "  public:\n",
+    "    concurrent_priority_queue();\n",
+    "    void push( const value_type& value );\n",
+    "    bool try_pop( value_type& value );\n",
+    "  };\n",
+    "```\n",
+    "\n",
+    "You can find detailed documentation for ``concurrent_queue`` [here](https://spec.oneapi.com/versions/latest/elements/oneTBB/source/containers/concurrent_priority_queue_cls.html).\n",
+    "\n",
+    "For this exercise, complete the following steps:\n",
+    "\n",
+    "1. Inspect the code cell below and make the following modifications.\n",
+    "  1. Replace the type `std::priority_queue` with `tbb::concurrent_priority_queue`.\n",
+    "  2. Replace the invocation of `empty`, `top` and `pop` with a single call to `try_pop`.\n",
+    "2. When the modifications are complete, click run ▶ to save the code to a file. \n",
+    "3. Run ▶ the cell in the __Build and Run the modified code__ section below the code snippet to compile and execute the code in the saved file."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "%%writefile lab/q-parallel.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <iostream>\n",
+    "#include <queue>\n",
+    "#include <thread>\n",
+    "#include <tbb/tbb.h>\n",
+    "\n",
+    "#define INCORRECT_QUEUE_TYPE std::priority_queue<int>;\n",
+    "\n",
+    "// STEP A: Replace std::priority_queue with tbb::concurrent_priority_queue\n",
+    "using queue_type = INCORRECT_QUEUE_TYPE;\n",
+    "        \n",
+    "int consume(queue_type &myPQ) {\n",
+    "  int local_sum = 0;\n",
+    "  int value = -1;\n",
+    "    \n",
+    "  bool consuming = true;\n",
+    "  while(consuming) {\n",
+    "    // STEP B: replace the pattern of empty, top and pop \n",
+    "    //         with a single call to myPQ.try_pop(value)\n",
+    "    if (!myPQ.empty()) {\n",
+    "      value = myPQ.top();\n",
+    "      myPQ.pop();\n",
+    "      if (value == -1)\n",
+    "        consuming = false;\n",
+    "      else\n",
+    "        local_sum += value;\n",
+    "    }\n",
+    "  }\n",
+    "  return local_sum;\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  int sum (0);\n",
+    "  int item;\n",
+    "\n",
+    "  queue_type myPQ;\n",
+    " \n",
+    "  std::thread producer([&]() {\n",
+    "    for(int i=0; i<10001; i+=1) {\n",
+    "      myPQ.push(i);\n",
+    "    }\n",
+    "    // to signal the end to the two consumers\n",
+    "    myPQ.push(-1);\n",
+    "    myPQ.push(-1);\n",
+    "  });\n",
+    "\n",
+    "  int local_sum1 = 0, local_sum2 = 0;\n",
+    "  std::thread consumer1([&]() { local_sum1 = consume(myPQ); });\n",
+    "  std::thread consumer2([&]() { local_sum2 = consume(myPQ); });\n",
+    "\n",
+    "  producer.join();\n",
+    "  consumer1.join();\n",
+    "  consumer2.join();\n",
+    "    \n",
+    "  // prints \"total: 50005000\" (for 0,10001,1)\n",
+    "  std::cout << \"total: \" << local_sum1 + local_sum2 << '\\n';\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### Build and Run the modified code\n",
+    "\n",
+    "Select the cell below and click Run ▶ to compile and execute the code that you modified above:"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_q-parallel.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_q-parallel.sh; else ./scripts/run_q-parallel.sh; fi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Producer-Consumer Solution (Don't peak unless you have to)"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "%%writefile solutions/q-parallel.cpp\n",
+    "//==============================================================\n",
+    "// Copyright (c) 2020 Intel Corporation\n",
+    "//\n",
+    "// SPDX-License-Identifier: Apache-2.0\n",
+    "// =============================================================\n",
+    "\n",
+    "#include <iostream>\n",
+    "#include <queue>\n",
+    "#include <thread>\n",
+    "#include <tbb/tbb.h>\n",
+    "\n",
+    "// STEP A: Replace std::priority_queue with tbb::concurrent_priority_queue\n",
+    "using queue_type = tbb::concurrent_priority_queue<int>;\n",
+    "        \n",
+    "int consume(queue_type &myPQ) {\n",
+    "  int local_sum = 0;\n",
+    "  int value = -1;\n",
+    "    \n",
+    "  bool consuming = true;\n",
+    "  while(consuming) {\n",
+    "    // STEP B: replace the pattern of empty, top and pop \n",
+    "    //         with a single call to myPQ.try_pop(value)\n",
+    "    if (myPQ.try_pop(value)) {\n",
+    "      if (value == -1)\n",
+    "        consuming = false;\n",
+    "      else\n",
+    "        local_sum += value;\n",
+    "    }\n",
+    "  }\n",
+    "  return local_sum;\n",
+    "}\n",
+    "\n",
+    "int main() {\n",
+    "  int sum (0);\n",
+    "  int item;\n",
+    "\n",
+    "  queue_type myPQ;\n",
+    " \n",
+    "  std::thread producer([&]() {\n",
+    "    for(int i=0; i<10001; i+=1) {\n",
+    "      myPQ.push(i);\n",
+    "    }\n",
+    "    // to signal the end to the two consumers\n",
+    "    myPQ.push(-1);\n",
+    "    myPQ.push(-1);\n",
+    "  });\n",
+    "\n",
+    "  int local_sum1 = 0, local_sum2 = 0;\n",
+    "  std::thread consumer1([&]() { local_sum1 = consume(myPQ); });\n",
+    "  std::thread consumer2([&]() { local_sum2 = consume(myPQ); });\n",
+    "\n",
+    "  producer.join();\n",
+    "  consumer1.join();\n",
+    "  consumer2.join();\n",
+    "    \n",
+    "  // prints \"total: 50005000\" (for 0,10001,1)\n",
+    "  std::cout << \"total: \" << local_sum1 + local_sum2 << '\\n';\n",
+    "  return 0;\n",
+    "}"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {
+    "jupyter": {
+     "source_hidden": true
+    }
+   },
+   "outputs": [],
+   "source": [
+    "! chmod 755 q; chmod 755 ./scripts/run_q-solution.sh; if [ -x \"$(command -v qsub)\" ]; then ./q scripts/run_q-solution.sh; else ./scripts/run_q-solution.sh; fi"
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3.7 (Intel® oneAPI)",
+   "language": "python",
+   "name": "c009-intel_distribution_of_python_3_oneapi-beta05-python"
+  },
+  "language_info": {
+   "codemirror_mode": {
+    "name": "ipython",
+    "version": 3
+   },
+   "file_extension": ".py",
+   "mimetype": "text/x-python",
+   "name": "python",
+   "nbconvert_exporter": "python",
+   "pygments_lexer": "ipython3",
+   "version": "3.7.7"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 4
+}
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
tbbdev	1482c0712a	Add bin and lab subfolders to examples/SC20	2020-11-06 10:45:56 +03:00
tbbdev	69d0954ba9	Committing oneTBB tutorials for Supercomputing Conference 2020	2020-11-05 15:30:59 +03:00
tbbdev	46d58dc838	Committing oneTBB tutorials CppCon2020	2020-09-18 17:57:53 +03:00