Extend scanline comparisons for image padding (#2619)

Extend compare_scanlines to handle INT10X6/12X4/14X2 padding and
per-channel SNORM INT8/INT16 comparisons, and switch pixel loads to
memcpy to avoid unaligned access. Rework image_from_buffer_fill_check to
compare rows via compare_scanlines.

Signed-off-by: Xin Jin <xin.jin@arm.com>

Author: xinjin01

test_common/harness/imageHelpers.cpp

@@ -493,43 +493,119 @@ size_t compare_scanlines(const image_descriptor *imageInfo, const char *aPtr,
             // If the data type is 101010, then ignore bits 31 and 32 when
             // comparing the row
             case CL_UNORM_INT_101010: {
-                cl_uint aPixel = *(cl_uint *)aPtr;
-                cl_uint bPixel = *(cl_uint *)bPtr;
+                cl_uint aPixel = 0;
+                cl_uint bPixel = 0;
+                memcpy(&aPixel, aPtr, sizeof(aPixel));
+                memcpy(&bPixel, bPtr, sizeof(bPixel));
                 if ((aPixel & 0x3fffffff) != (bPixel & 0x3fffffff))
                     return column;
             }
             break;
 
             // If the data type is 555, ignore bit 15 when comparing the row
             case CL_UNORM_SHORT_555: {
-                cl_ushort aPixel = *(cl_ushort *)aPtr;
-                cl_ushort bPixel = *(cl_ushort *)bPtr;
+                cl_ushort aPixel = 0;
+                cl_ushort bPixel = 0;
+                memcpy(&aPixel, aPtr, sizeof(aPixel));
+                memcpy(&bPixel, bPtr, sizeof(bPixel));
                 if ((aPixel & 0x7fff) != (bPixel & 0x7fff)) return column;
             }
             break;
 
+            // 16-bit per-channel formats with LSB padding (per spec); compare
+            // defined bits only.
+            case CL_UNSIGNED_INT10X6_EXT:
+            case CL_UNORM_INT10X6_EXT: {
+                const size_t channel_count =
+                    get_format_channel_count(imageInfo->format);
+                for (size_t chan = 0; chan < channel_count; ++chan)
+                {
+                    cl_ushort aChanVal = 0;
+                    cl_ushort bChanVal = 0;
+                    const char *aChan = aPtr + chan * sizeof(cl_ushort);
+                    const char *bChan = bPtr + chan * sizeof(cl_ushort);
+                    memcpy(&aChanVal, aChan, sizeof(aChanVal));
+                    memcpy(&bChanVal, bChan, sizeof(bChanVal));
+                    if ((aChanVal & 0xffc0) != (bChanVal & 0xffc0))
+                        return column;
+                }
+            }
+            break;
+            case CL_UNSIGNED_INT12X4_EXT:
+            case CL_UNORM_INT12X4_EXT: {
+                const size_t channel_count =
+                    get_format_channel_count(imageInfo->format);
+                for (size_t chan = 0; chan < channel_count; ++chan)
+                {
+                    cl_ushort aChanVal = 0;
+                    cl_ushort bChanVal = 0;
+                    const char *aChan = aPtr + chan * sizeof(cl_ushort);
+                    const char *bChan = bPtr + chan * sizeof(cl_ushort);
+                    memcpy(&aChanVal, aChan, sizeof(aChanVal));
+                    memcpy(&bChanVal, bChan, sizeof(bChanVal));
+                    if ((aChanVal & 0xfff0) != (bChanVal & 0xfff0))
+                        return column;
+                }
+            }
+            break;
+            case CL_UNSIGNED_INT14X2_EXT:
+            case CL_UNORM_INT14X2_EXT: {
+                const size_t channel_count =
+                    get_format_channel_count(imageInfo->format);
+                for (size_t chan = 0; chan < channel_count; ++chan)
+                {
+                    cl_ushort aChanVal = 0;
+                    cl_ushort bChanVal = 0;
+                    const char *aChan = aPtr + chan * sizeof(cl_ushort);
+                    const char *bChan = bPtr + chan * sizeof(cl_ushort);
+                    memcpy(&aChanVal, aChan, sizeof(aChanVal));
+                    memcpy(&bChanVal, bChan, sizeof(bChanVal));
+                    if ((aChanVal & 0xfffc) != (bChanVal & 0xfffc))
+                        return column;
+                }
+            }
+            break;
+
             case CL_SNORM_INT8: {
-                cl_uchar aPixel = *(cl_uchar *)aPtr;
-                cl_uchar bPixel = *(cl_uchar *)bPtr;
-                // -1.0 is defined as 0x80 and 0x81
-                aPixel = (aPixel == 0x80) ? 0x81 : aPixel;
-                bPixel = (bPixel == 0x80) ? 0x81 : bPixel;
-                if (aPixel != bPixel)
+                const size_t channel_count =
+                    get_format_channel_count(imageInfo->format);
+                for (size_t chan = 0; chan < channel_count; ++chan)
                 {
-                    return column;
+                    cl_uchar aChanVal = 0;
+                    cl_uchar bChanVal = 0;
+                    const char *aChan = aPtr + chan * sizeof(cl_uchar);
+                    const char *bChan = bPtr + chan * sizeof(cl_uchar);
+                    memcpy(&aChanVal, aChan, sizeof(aChanVal));
+                    memcpy(&bChanVal, bChan, sizeof(bChanVal));
+                    // -1.0 is defined as 0x80 and 0x81
+                    aChanVal = (aChanVal == 0x80) ? 0x81 : aChanVal;
+                    bChanVal = (bChanVal == 0x80) ? 0x81 : bChanVal;
+                    if (aChanVal != bChanVal)
+                    {
+                        return column;
+                    }
                 }
             }
             break;
 
             case CL_SNORM_INT16: {
-                cl_ushort aPixel = *(cl_ushort *)aPtr;
-                cl_ushort bPixel = *(cl_ushort *)bPtr;
-                // -1.0 is defined as 0x8000 and 0x8001
-                aPixel = (aPixel == 0x8000) ? 0x8001 : aPixel;
-                bPixel = (bPixel == 0x8000) ? 0x8001 : bPixel;
-                if (aPixel != bPixel)
+                const size_t channel_count =
+                    get_format_channel_count(imageInfo->format);
+                for (size_t chan = 0; chan < channel_count; ++chan)
                 {
-                    return column;
+                    cl_ushort aChanVal = 0;
+                    cl_ushort bChanVal = 0;
+                    const char *aChan = aPtr + chan * sizeof(cl_ushort);
+                    const char *bChan = bPtr + chan * sizeof(cl_ushort);
+                    memcpy(&aChanVal, aChan, sizeof(aChanVal));
+                    memcpy(&bChanVal, bChan, sizeof(bChanVal));
+                    // -1.0 is defined as 0x8000 and 0x8001
+                    aChanVal = (aChanVal == 0x8000) ? 0x8001 : aChanVal;
+                    bChanVal = (bChanVal == 0x8000) ? 0x8001 : bChanVal;
+                    if (aChanVal != bChanVal)
+                    {
+                        return column;
+                    }
                 }
             }
             break;
@@ -544,6 +620,7 @@ size_t compare_scanlines(const image_descriptor *imageInfo, const char *aPtr,
     }
 
     // If we didn't find a difference, return the width of the image
+    assert(column == imageInfo->width);
     return column;
 }
 
@@ -1208,19 +1285,21 @@ void escape_inf_nan_subnormal_values(char *data, size_t allocSize)
 {
     // filter values with 8 not-quite-highest bits
     unsigned int *intPtr = (unsigned int *)data;
-    for (size_t i = 0; i<allocSize>> 2; i++)
+    for (size_t i = 0; i < (allocSize >> 2); i++)
     {
-        if ((intPtr[i] & 0x7F800000) == 0x7F800000) intPtr[i] ^= 0x40000000;
+        if ((intPtr[i] & 0x7F800000) == 0x7F800000)
+            intPtr[i] ^= 0x40000000;
         else if ((intPtr[i] & 0x7F800000) == 0)
             intPtr[i] ^= 0x40000000;
     }
 
     // Ditto with half floats (16-bit numbers with the 5 not-quite-highest bits
     // = 0x7C00 are special)
     unsigned short *shortPtr = (unsigned short *)data;
-    for (size_t i = 0; i<allocSize>> 1; i++)
+    for (size_t i = 0; i < (allocSize >> 1); i++)
     {
-        if ((shortPtr[i] & 0x7C00) == 0x7C00) shortPtr[i] ^= 0x4000;
+        if ((shortPtr[i] & 0x7C00) == 0x7C00)
+            shortPtr[i] ^= 0x4000;
         else if ((shortPtr[i] & 0x7C00) == 0)
             shortPtr[i] ^= 0x4000;
     }
@@ -3191,12 +3270,12 @@ void pack_image_pixel_error(const float *srcVector,
         case CL_UNSIGNED_INT32: {
             const cl_uint *ptr = (const cl_uint *)results;
             for (unsigned int i = 0; i < channelCount; i++)
-                errors[i] = (cl_float)(
-                    (cl_long)ptr[i]
-                    - (cl_long)CONVERT_UINT(
-                        srcVector[i],
-                        MAKE_HEX_FLOAT(0x1.fffffep31f, 0x1fffffe, 31 - 23),
-                        CL_UINT_MAX));
+                errors[i] = (cl_float)((cl_long)ptr[i]
+                                       - (cl_long)CONVERT_UINT(
+                                           srcVector[i],
+                                           MAKE_HEX_FLOAT(0x1.fffffep31f,
+                                                          0x1fffffe, 31 - 23),
+                                           CL_UINT_MAX));
             break;
         }
         case CL_UNSIGNED_INT10X6_EXT: {

test_conformance/extensions/cl_khr_external_memory_ahb/debug_ahb.cpp

@@ -60,7 +60,7 @@ ahardwareBufferDecodeUsageFlagsToString(const AHardwareBuffer_UsageFlags flags)
 
     return std::accumulate(active_flags.begin() + 1, active_flags.end(),
                            active_flags.front(),
-                           [](std::string acc, const std::string& flag) {
+                           [](std::string acc, const std::string &flag) {
                                return std::move(acc) + "|" + flag;
                            });
 }

test_conformance/extensions/cl_khr_external_semaphore/test_external_semaphore.cpp

@@ -52,7 +52,7 @@
         }                                                                      \
     } while (false)
 
-static const char *source = "__kernel void empty() {}";
+static const char* source = "__kernel void empty() {}";
 
 static void log_info_semaphore_type(
     VulkanExternalSemaphoreHandleType vkExternalSemaphoreHandleType)
@@ -64,7 +64,7 @@ static void log_info_semaphore_type(
     log_info("%s", semaphore_type_description.str().c_str());
 }
 
-static int init_vulkan_device(cl_uint num_devices, cl_device_id *deviceIds)
+static int init_vulkan_device(cl_uint num_devices, cl_device_id* deviceIds)
 {
     cl_platform_id platform = nullptr;
 
@@ -84,7 +84,7 @@ static int init_vulkan_device(cl_uint num_devices, cl_device_id *deviceIds)
 
 static cl_int get_device_semaphore_handle_types(
     cl_device_id deviceID, cl_device_info param,
-    std::vector<cl_external_semaphore_handle_type_khr> &handle_types)
+    std::vector<cl_external_semaphore_handle_type_khr>& handle_types)
 {
     int err = CL_SUCCESS;
     // Query for export support

test_conformance/images/kernel_read_write/test_cl_ext_image_from_buffer.cpp

@@ -642,11 +642,12 @@ int image_from_small_buffer_negative(cl_device_id device, cl_context context,
 }
 
 static int image_from_buffer_fill_check(cl_command_queue queue, cl_mem image,
-                                        size_t* region, size_t element_size,
-                                        char pattern)
+                                        const cl_image_format& format,
+                                        size_t* region, char pattern)
 {
     /* read the image from buffer and check the pattern */
-    const size_t image_size = region[0] * region[1] * region[2] * element_size;
+    const size_t pixel_size = get_pixel_size(&format);
+    const size_t image_size = region[0] * region[1] * region[2] * pixel_size;
     size_t origin[3] = { 0, 0, 0 };
     std::vector<char> read_buffer(image_size);
 
@@ -655,21 +656,29 @@ static int image_from_buffer_fill_check(cl_command_queue queue, cl_mem image,
                            read_buffer.data(), 0, nullptr, nullptr);
     test_error(error, "Error clEnqueueReadImage");
 
-    for (size_t line = 0; line < region[0]; line++)
+    const size_t row_pitch = region[0] * pixel_size;
+    const size_t slice_pitch = row_pitch * region[1];
+
+    image_descriptor image_info = {};
+    image_info.width = region[0];
+    image_info.format = &format;
+
+    std::vector<char> expected_row(row_pitch, pattern);
+
+    for (size_t depth = 0; depth < region[2]; depth++)
     {
         for (size_t row = 0; row < region[1]; row++)
         {
-            for (size_t depth = 0; depth < region[2]; depth++)
+            const char* actual_ptr =
+                read_buffer.data() + depth * slice_pitch + row * row_pitch;
+            size_t where =
+                compare_scanlines(&image_info, expected_row.data(), actual_ptr);
+            // compare_scanlines returns width when rows match; < width
+            // indicates mismatch.
+            if (where < region[0])
             {
-                for (size_t elmt = 0; elmt < element_size; elmt++)
-                {
-                    size_t index = line * row * depth * elmt;
-
-                    if (read_buffer[index] != pattern)
-                    {
-                        test_fail("Image pattern check failed\n");
-                    }
-                }
+                test_fail("Image pattern check failed (z=%zu, y=%zu, x=%zu)\n",
+                          depth, row, where);
             }
         }
     }
@@ -827,14 +836,14 @@ int image_from_buffer_fill_positive(cl_device_id device, cl_context context,
                 test_error(err, "Error clFinish");
 
                 int fill_error = image_from_buffer_fill_check(
-                    queue, image_from_buffer, region, element_size, pattern);
+                    queue, image_from_buffer, format, region, pattern);
                 if (TEST_PASS != fill_error)
                 {
                     return fill_error;
                 }
 
-                fill_error = image_from_buffer_fill_check(
-                    queue, image, region, element_size, pattern);
+                fill_error = image_from_buffer_fill_check(queue, image, format,
+                                                          region, pattern);
                 if (TEST_PASS != fill_error)
                 {
                     return fill_error;