Applying same fix to ex 19

Author: Hugh Delaney

Code_Exercises/Exercise_19_Work_Group_Sizes/solution.cpp

@@ -23,18 +23,15 @@
 #include <benchmark.h>
 #include <image_conv.h>
 
-
 class image_convolution;
 
 inline constexpr util::filter_type filterType = util::filter_type::blur;
 inline constexpr int filterWidth = 11;
 inline constexpr int halo = filterWidth / 2;
 
-TEST_CASE("image_convolution_work_group_sizes", "work_group_sizes_solution") {
-  const char* inputImageFile =
-      "../Images/dogs.png";
-  const char* outputImageFile =
-      "../Images/blurred_dogs.png";
+TEST_CASE("image_convolution_tiled", "local_memory_tiling_solution") {
+  constexpr auto inputImageFile = "../Images/dogs.png";
+  constexpr auto outputImageFile = "../Images/blurred_dogs.png";
 
   auto inputImage = util::read_image(inputImageFile, halo);
 
@@ -84,8 +81,8 @@ TEST_CASE("image_convolution_work_group_sizes", "work_group_sizes_solution") {
           filterRange / sycl::range(1, channels));
 
       util::benchmark(
-          [&]() {
-            myQueue.submit([&](sycl::handler& cgh) {
+          [&] {
+            myQueue.submit([&](sycl::handler &cgh) {
               auto inputAcc =
                   inBufVec.get_access<sycl::access::mode::read>(cgh);
               auto outputAcc =
@@ -101,31 +98,31 @@ TEST_CASE("image_convolution_work_group_sizes", "work_group_sizes_solution") {
               cgh.parallel_for<image_convolution>(
                   ndRange, [=](sycl::nd_item<2> item) {
                     auto globalId = item.get_global_id();
-                    globalId = sycl::id{globalId[1], globalId[0]};
+                    auto groupId = item.get_group().get_group_id();
                     auto localId = item.get_local_id();
+                    auto globalGroupOffset = groupId * localRange;
+
+                    for (auto i = localId[0]; i < scratchpadRange[0];
+                         i += localRange[0]) {
+                      for (auto j = localId[1]; j < scratchpadRange[1];
+                           j += localRange[1]) {
+                        scratchpad[i][j] =
+                            inputAcc[globalGroupOffset + sycl::range(i, j)];
+                      }
+                    }
 
-                    auto haloOffset = sycl::id(halo, halo);
-                    auto src = (globalId + haloOffset);
-                    auto dest = globalId;
-                    auto temp = localId + halo;
-
-                    scratchpad[temp] = inputAcc[src];
-
-                    sycl::group_barrier(item.get_group());
+                    item.barrier();
 
                     auto sum = sycl::float4{0.0f, 0.0f, 0.0f, 0.0f};
 
                     for (int r = 0; r < filterWidth; ++r) {
                       for (int c = 0; c < filterWidth; ++c) {
-                        auto srcOffset = sycl::id(temp[0] + (r - halo),
-                                                  temp[1] + ((c - halo)));
-                        auto filterOffset = sycl::id(r, c);
-
-                        sum += scratchpad[srcOffset] * filterAcc[filterOffset];
+                        auto idx = sycl::range(r, c);
+                        sum += scratchpad[localId + idx] * filterAcc[idx];
                       }
                     }
 
-                    outputAcc[dest] = sum;
+                    outputAcc[globalId] = sum;
                   });
             });