Make image conv memory coalescence example more intuitive

The previous arrangement used a rather devious (32, 1) group size to
enforce a requirement for flipping the globalId inside the
kernel. This is fine but it makes it seem like a sycl ID is row major
but a sycl::buffer is col major.

The exercise is now kind of reversed - we keep an intuitive work-group
shape in reference.cpp but we add the (now rogue) globalId flip so
that memory access isn't coalesced. The 'solution' for the coalescence
exercise now is to remove the globalId flip.

Author: joeatodd

Code_Exercises/Exercise_15_Image_Convolution/reference.cpp

@@ -61,7 +61,7 @@ TEST_CASE("image_convolution_naive", "image_convolution_reference") {
     auto halo = filter.half_width();
 
     auto globalRange = sycl::range(inputImgWidth, inputImgHeight);
-    auto localRange = sycl::range(32, 1);
+    auto localRange = sycl::range(1, 32);
     auto ndRange = sycl::nd_range(globalRange, localRange);
 
     auto inBufRange = (inputImgWidth + (halo * 2)) * sycl::range(1, channels);
@@ -86,6 +86,7 @@ TEST_CASE("image_convolution_naive", "image_convolution_reference") {
               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 channelsStride = sycl::range(1, channels);
                     auto haloOffset = sycl::id(halo, halo);

Code_Exercises/Exercise_16_Coalesced_Global_Memory/solution.cpp

@@ -61,7 +61,7 @@ TEST_CASE("image_convolution_coalesced", "coalesced_global_memory_solution") {
     auto halo = filter.half_width();
 
     auto globalRange = sycl::range(inputImgWidth, inputImgHeight);
-    auto localRange = sycl::range(32, 1);
+    auto localRange = sycl::range(1, 32);
     auto ndRange = sycl::nd_range(globalRange, localRange);
 
     auto inBufRange = (inputImgWidth + (halo * 2)) * sycl::range(1, channels);
@@ -86,7 +86,6 @@ TEST_CASE("image_convolution_coalesced", "coalesced_global_memory_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 channelsStride = sycl::range(1, channels);
                     auto haloOffset = sycl::id(halo, halo);

Code_Exercises/Exercise_17_Vectors/solution.cpp

@@ -62,7 +62,7 @@ TEST_CASE("image_convolution_vectorized", "vectors_solution") {
     auto halo = filter.half_width();
 
     auto globalRange = sycl::range(inputImgWidth, inputImgHeight);
-    auto localRange = sycl::range(32, 1);
+    auto localRange = sycl::range(1, 32);
     auto ndRange = sycl::nd_range(globalRange, localRange);
 
     auto inBufRange = (inputImgWidth + (halo * 2)) * sycl::range(1, channels);
@@ -95,7 +95,6 @@ TEST_CASE("image_convolution_vectorized", "vectors_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 haloOffset = sycl::id(halo, halo);
                     auto src = (globalId + haloOffset);