update snippets

snippets/snip-UnoAPI-f-implementation.tex

@@ -1,5 +1,5 @@
-\begin{cppcode*}{firstnumber=4}
+\begin{cppcode*}{firstnumber=6}
 double f(const double x) {
-    return cos(x) * cos(x) + sin(x) * sin(x);
+    return 3 * pow(x, 2);
 }
 \end{cppcode*}

snippets/snip-UnoAPI-f-interface.tex

@@ -1,5 +1,5 @@
 \begin{ccode*}{firstnumber=6}
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 
 SYCL_EXTERNAL double f(double x);
 \end{ccode*}

snippets/snip-UnoAPI-integration-test-f1.tex

@@ -1,5 +1,5 @@
 \begin{cppcode*}{firstnumber=42}
 TEST_F(IntegrationTest, F1) {
-    EXPECT_NEAR(f(0.5), 1, EPS);
+    EXPECT_NEAR(f(0.5), 0.75, EPS);
 }
 \end{cppcode*}

snippets/snip-UnoAPI-main-cli-setup-and-parse.tex

@@ -1,14 +1,16 @@
-\begin{cppcode*}{firstnumber=36}
+\begin{cppcode*}{firstnumber=63}
     CLI::App app{"Trapezoidal integration"};
     app.option_defaults()->always_capture_default(true);
-    app.add_option("-n,--trapezoids", number_of_trapezoids, "number of trapezoids")->check(CLI::PositiveNumber.description(" >= 1"));
     app.add_option("-l,--lower,--xmin", x_min, "x min value");
     app.add_option("-u,--upper,--xmax", x_max, "x max value");
-    app.add_flag("-v,--show-function-values", show_function_values);
+    app.add_option("-n,--total-workload", total_workload, "total workload (number of trapezoids)")->check(CLI::PositiveNumber.description(" >= 1"));
+    app.add_option("-g,--grain-size", grain_size, "number of inner (sequential) trapezoids (for each outer trapezoid)")->check(CLI::PositiveNumber.description(" >= 1"));
     app.add_flag("-s,--sequential", run_sequentially);
     app.add_flag("-c,--cpu-only", run_cpuonly);
+    app.add_flag("-v,--show-function-values", show_function_values);
     app.add_option("-x,--x-format-precision", x_precision, "decimal precision for x values")->check(CLI::PositiveNumber.description(" >= 1"));
     app.add_option("-y,--y-format-precision", y_precision, "decimal precision for y (function) values")->check(CLI::PositiveNumber.description(" >= 1"));
+    app.add_option("-p,--perfdata-output-file", perf_output, "output file for performance data");
 
     CLI11_PARSE(app, argc, argv);
 \end{cppcode*}

snippets/snip-UnoAPI-main-compute-outer-trapezoid.tex

@@ -0,0 +1,19 @@
+\begin{cppcode*}{firstnumber=27}
+// common function to compute a single outer trapezoid
+// from as many inner trapezoids as the grain size
+double compute_outer_trapezoid(
+    const int grain_size,
+    const double x_pos,
+    const double dx_inner,
+    const double half_dx_inner
+) {
+    auto area{0.0};
+    auto y_left{f(x_pos)};
+    for (auto j{0UL}; j < grain_size; j++) {
+        auto y_right{f(x_pos + (j + 1) * dx_inner)};
+        area += trapezoid(y_left, y_right, half_dx_inner);
+        y_left = y_right;
+    }
+    return area;
+}
+\end{cppcode*}

snippets/snip-UnoAPI-main-declarations.tex

@@ -1,11 +1,14 @@
-\begin{cppcode*}{firstnumber=24}
-    constexpr size_t DEFAULT_NUMBER_OF_TRAPEZOIDS{10};
-    size_t number_of_trapezoids{DEFAULT_NUMBER_OF_TRAPEZOIDS};
+\begin{cppcode*}{firstnumber=48}
+    size_t total_workload{1000};
+    uint grain_size{100};
     double x_min{0.0};
     double x_max{1.0};
     bool show_function_values{false};
     bool run_sequentially{false};
     bool run_cpuonly{false};
     uint x_precision{1};
     uint y_precision{1};
+    std::string perf_output;
+    ts_vector timestamps;
+    std::string device_name;
 \end{cppcode*}

snippets/snip-UnoAPI-main-domain-setup.tex

@@ -1,5 +1,8 @@
-\begin{cppcode*}{firstnumber=57}
+\begin{cppcode*}{firstnumber=89}
+    size_t number_of_trapezoids{total_workload / grain_size};
     const auto size{number_of_trapezoids + 1};
     const auto dx{(x_max - x_min) / number_of_trapezoids};
     const auto half_dx{0.5 * dx}; // precomputed for area calculation
+    const auto dx_inner{dx / grain_size};
+    const auto half_dx_inner{0.5 * dx_inner};
 \end{cppcode*}

snippets/snip-UnoAPI-main-parallel-buffers.tex

@@ -1,4 +1,5 @@
-\begin{cppcode*}{firstnumber=92}
+\begin{cppcode*}{firstnumber=134}
         sycl::buffer<double> v_buf{sycl::range<1>{size}};
+        sycl::buffer<double> t_buf{sycl::range<1>{number_of_trapezoids}};
         sycl::buffer<double> r_buf{sycl::range<1>{1}};
 \end{cppcode*}

snippets/snip-UnoAPI-main-parallel-devices.tex

@@ -1,3 +1,3 @@
-\begin{cppcode*}{firstnumber=99}
+\begin{cppcode*}{firstnumber=143}
         sycl::device device { run_cpuonly ? sycl::cpu_selector_v : sycl::default_selector_v };
 \end{cppcode*}

snippets/snip-UnoAPI-main-parallel-gather-on-host.tex

@@ -1,5 +1,6 @@
-\begin{cppcode*}{firstnumber=138}
+\begin{cppcode*}{firstnumber=189}
         const sycl::host_accessor result{r_buf};
+        mark_time(timestamps,"Integration");
         spdlog::info("result should be available now");
         fmt::print("result = {}\n", result[0]);
 \end{cppcode*}

snippets/snip-UnoAPI-main-parallel-inorder-q.tex

@@ -1,4 +1,6 @@
-\begin{cppcode*}{firstnumber=110}
-        sycl::queue q{device, dpc_common::exception_handler, sycl::property::queue::in_order()};
-        spdlog::info("Device: {}", q.get_device().get_info<sycl::info::device::name>());
+\begin{cppcode*}{firstnumber=148}
+        sycl::queue q{device, dpc_common::exception_handler};
+        mark_time(timestamps,"Queue creation");
+        device_name = q.get_device().get_info<sycl::info::device::name>();
+        spdlog::info("Device: {}", device_name);
 \end{cppcode*}

snippets/snip-UnoAPI-main-parallel-show-results-log.tex

@@ -1,8 +1,10 @@
-\begin{cppcode*}{firstnumber=144}
+\begin{cppcode*}{firstnumber=196}
         if (show_function_values) {
             spdlog::info("preparing function values");
             const sycl::host_accessor values{v_buf};
+            mark_time(timestamps,"Host data access");
             spdlog::info("showing function values");
             print_function_values(values, x_min, dx, x_precision, y_precision);
+            mark_time(timestamps,"Output");
         }
 \end{cppcode*}

snippets/snip-UnoAPI-main-parallel-submit-parallel-for.tex

@@ -1,4 +1,4 @@
-\begin{cppcode*}{firstnumber=116}
+\begin{cppcode*}{firstnumber=156}
         q.submit([&](auto & h) {
             const sycl::accessor v{v_buf, h};
             h.parallel_for(size, [=](const auto & index) {

snippets/snip-UnoAPI-main-parallel-submit-reduce.tex

@@ -1,9 +1,9 @@
-\begin{cppcode*}{firstnumber=126}
+\begin{cppcode*}{firstnumber=177}
         q.submit([&](auto & h) {
-            const sycl::accessor v{v_buf, h};
+            const sycl::accessor t{t_buf, h};
             const auto sum_reduction{sycl::reduction(r_buf, h, sycl::plus<>())};
             h.parallel_for(sycl::range<1>{number_of_trapezoids}, sum_reduction, [=](const auto & index, auto & sum) {
-                sum.combine(trapezoid(v[index], v[index + 1], half_dx));
+                sum.combine(t[index]);
             });
         }); // end of command group
 \end{cppcode*}

snippets/snip-UnoAPI-main-print-function-values.tex

@@ -1,6 +1,7 @@
-\begin{cppcode*}{firstnumber=13}
+\begin{cppcode*}{firstnumber=16}
 // function template to avoid code repetition (DRY)
-// works as long as the first argument has a size() method and an indexing operator[]
+// works as long as the first argument has a size()
+// method and an indexing operator[]
 template <class Indexable> void print_function_values(const Indexable & values, const double x_min, const double dx, const uint x_precision, const uint y_precision) {
     for (auto i{0UL}; i < values.size(); i++) {
         fmt::print("{}: f({:.{}f}) = {:.{}f}\n", i, x_min + i * dx, x_precision, values[i], y_precision);

snippets/snip-UnoAPI-main-sequential-option.tex

@@ -1,23 +1,29 @@
-\begin{cppcode*}{firstnumber=65}
+\begin{cppcode*}{firstnumber=101}
     if (run_sequentially) {
+        device_name = "sequential";
         std::vector values(size, 0.0);
-        double result{0.0};
+        auto result{0.0};
 
+        mark_time(timestamps,"Memory allocation");
         spdlog::info("starting sequential integration");
 
         // populate vector with function values and add trapezoid area to result
-        values[0] = f(x_min);
+        // the inner loop performs a finer-grained calculation
+        values[0] += f(x_min);
         for (auto i{0UL}; i < number_of_trapezoids; i++) {
-            values[i + 1] = f(x_min + i * dx);
-            result += trapezoid(values[i], values[i + 1], half_dx);
+            const auto x_pos{x_min + i * dx};
+            result += compute_outer_trapezoid(grain_size, x_pos, dx_inner, half_dx_inner);
+            values[i + 1] = f(x_pos);
         }
 
+        mark_time(timestamps, "Integration");
         spdlog::info("result should be available now");
         fmt::print("result = {}\n", result);
 
         if (show_function_values) {
             spdlog::info("showing function values");
             print_function_values(values, x_min, dx, x_precision, y_precision);
+            mark_time(timestamps, "Output");
         }
     }
 \end{cppcode*}

snippets/snip-UnoAPI-timestamps-mark-time.tex

@@ -0,0 +1,5 @@
+\begin{cppcode*}{firstnumber=7}
+void mark_time(ts_vector & timestamps, const std::string_view label) {
+    timestamps.push_back(std::pair(label.data(), std::chrono::steady_clock::now()));
+}
+\end{cppcode*}

snippets/snip-UnoAPI-timestamps-print-timestamps.tex

@@ -0,0 +1,24 @@
+\begin{cppcode*}{firstnumber=13}
+void print_timestamps(const ts_vector & timestamps, const std::string_view filename, const std::string_view device_name) {
+    using std::chrono::nanoseconds;
+    using std::chrono::duration_cast;
+
+    constexpr auto ROW_HEADER{"TIME,DELTA,UNIT,DEVICE,PHASE\n"};
+    constexpr auto ROW_FORMAT{"{},{},{},{},{}\n"};
+    constexpr auto TIME_UNIT{"ns"};
+
+    const auto & start = timestamps.front().second;
+    auto outfile = filename.empty() ? stdout : std::fopen(filename.data(), "w");
+    fmt::print(outfile, ROW_HEADER);
+    fmt::print(outfile, ROW_FORMAT, duration_cast<nanoseconds>(start.time_since_epoch()).count(), 0, TIME_UNIT, device_name, timestamps.front().first);
+    for (auto t = timestamps.begin() + 1; t != timestamps.end(); t++) {
+        const auto dur{duration_cast<nanoseconds>(t->second - (t - 1)->second).count()};
+        fmt::print(outfile, ROW_FORMAT, duration_cast<nanoseconds>(t->second.time_since_epoch()).count(), dur, TIME_UNIT, device_name, t->first);
+    }
+    const auto & stop{timestamps.back().second};
+    const auto total{duration_cast<nanoseconds>(stop - start).count()};
+    fmt::print(outfile, ROW_FORMAT, duration_cast<nanoseconds>(stop.time_since_epoch()).count(), total, TIME_UNIT, device_name, "TOTAL");
+    if (! filename.empty())
+        std::fclose(outfile);
+}
+\end{cppcode*}