Merge pull request #160 from codeplaysoftware/imageconv-perf-fix

Make image conv memory coalescence example more intuitive

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);
@@ -84,6 +84,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);
@@ -84,7 +84,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);
@@ -92,7 +92,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);