Merge pull request #158 from codeplaysoftware/imageconv-range-fix

Fix range dimension order to reflect row-major

Author: rodburns

Code_Exercises/Exercise_15_Image_Convolution/reference.cpp

@@ -64,8 +64,13 @@ TEST_CASE("image_convolution_naive", "image_convolution_reference") {
     auto localRange = sycl::range(1, 32);
     auto ndRange = sycl::nd_range(globalRange, localRange);
 
-    auto inBufRange = (inputImgWidth + (halo * 2)) * sycl::range(1, channels);
-    auto outBufRange = inputImgHeight * sycl::range(1, channels);
+    auto inBufRange =
+        sycl::range(inputImgHeight + (halo * 2), inputImgWidth + (halo * 2)) *
+        sycl::range(1, channels);
+    auto outBufRange =
+        sycl::range(inputImgHeight, inputImgWidth) * sycl::range(1, channels);
+
+
     auto filterRange = filterWidth * sycl::range(1, channels);
 
     {

Code_Exercises/Exercise_16_Coalesced_Global_Memory/solution.cpp

@@ -64,8 +64,11 @@ TEST_CASE("image_convolution_coalesced", "coalesced_global_memory_solution") {
     auto localRange = sycl::range(1, 32);
     auto ndRange = sycl::nd_range(globalRange, localRange);
 
-    auto inBufRange = (inputImgWidth + (halo * 2)) * sycl::range(1, channels);
-    auto outBufRange = inputImgHeight * sycl::range(1, channels);
+    auto inBufRange =
+        sycl::range(inputImgHeight + (halo * 2), inputImgWidth + (halo * 2)) *
+        sycl::range(1, channels);
+    auto outBufRange =
+        sycl::range(inputImgHeight, inputImgWidth) * sycl::range(1, channels);
     auto filterRange = filterWidth * sycl::range(1, channels);
 
     {

Code_Exercises/Exercise_17_Vectors/solution.cpp

@@ -65,8 +65,11 @@ TEST_CASE("image_convolution_vectorized", "vectors_solution") {
     auto localRange = sycl::range(1, 32);
     auto ndRange = sycl::nd_range(globalRange, localRange);
 
-    auto inBufRange = (inputImgWidth + (halo * 2)) * sycl::range(1, channels);
-    auto outBufRange = inputImgHeight * sycl::range(1, channels);
+    auto inBufRange =
+        sycl::range(inputImgHeight + (halo * 2), inputImgWidth + (halo * 2)) *
+        sycl::range(1, channels);
+    auto outBufRange =
+        sycl::range(inputImgHeight, inputImgWidth) * sycl::range(1, channels);
     auto filterRange = filterWidth * sycl::range(1, channels);
 
     {

Code_Exercises/Exercise_18_Local_Memory_Tiling/solution.cpp

@@ -58,12 +58,15 @@ TEST_CASE("image_convolution_tiled", "local_memory_tiling_solution") {
     auto filterWidth = filter.width();
     auto halo = filter.half_width();
 
-    auto globalRange = sycl::range(inputImgWidth, inputImgHeight);
+    auto globalRange = sycl::range(inputImgHeight, inputImgWidth);
     auto localRange = sycl::range(8, 8);
     auto ndRange = sycl::nd_range(globalRange, localRange);
 
-    auto inBufRange = (inputImgWidth + (halo * 2)) * sycl::range(1, channels);
-    auto outBufRange = inputImgHeight * sycl::range(1, channels);
+     auto inBufRange =
+        sycl::range(inputImgHeight + (halo * 2), inputImgWidth + (halo * 2)) *
+        sycl::range(1, channels);
+    auto outBufRange =
+        sycl::range(inputImgHeight, inputImgWidth) * sycl::range(1, channels);
     auto filterRange = filterWidth * sycl::range(1, channels);
     auto scratchpadRange = localRange + sycl::range(halo * 2, halo * 2);
 

Code_Exercises/Exercise_19_Work_Group_Sizes/solution.cpp

@@ -58,12 +58,15 @@ TEST_CASE("image_convolution_tiled", "local_memory_tiling_solution") {
     auto filterWidth = filter.width();
     auto halo = filter.half_width();
 
-    auto globalRange = sycl::range(inputImgWidth, inputImgHeight);
+    auto globalRange = sycl::range(inputImgHeight, inputImgWidth);
     auto localRange = sycl::range(16, 16);
     auto ndRange = sycl::nd_range(globalRange, localRange);
 
-    auto inBufRange = (inputImgWidth + (halo * 2)) * sycl::range(1, channels);
-    auto outBufRange = inputImgHeight * sycl::range(1, channels);
+    auto inBufRange =
+        sycl::range(inputImgHeight + (halo * 2), inputImgWidth + (halo * 2)) *
+        sycl::range(1, channels);
+    auto outBufRange =
+        sycl::range(inputImgHeight, inputImgWidth) * sycl::range(1, channels);
     auto filterRange = filterWidth * sycl::range(1, channels);
     auto scratchpadRange = localRange + sycl::range(halo * 2, halo * 2);