Test IMAGE1D_BUFFER in more scenario (#1806)

* cl_copy_images

* cl_get_info

* cl_fill_image

* cl_read_write_image

* kernel_image_methods

* IMAGE1D_BUFFER cannot be created with (USE_|ALLOC_|COPY_)_HOST_PTR

* do not allow mipmap with 1D buffer

* adjust M to be within maximum_sizes and max_pixels

* remove unused variables

* make sure M will never be 0

* fix region[0] after refactoring removing mipmap

* fix formatting

* format with clang-format-11

* fix image1d_buffer creation with gEnablePitch

* add missing case in switch

* use align_malloc when CL version is at least 2.0

* use CL_DEVICE_NUMERIC_VERSION and align_free

* fix free of pitch buffer

* fix formatting

* fix formatting

* fix data->is_aligned

test_common/harness/imageHelpers.cpp

@@ -769,10 +769,14 @@ void get_max_sizes(
 
     (*numberOfSizes) = 0;
 
-    if (image_type == CL_MEM_OBJECT_IMAGE1D)
+    if (image_type == CL_MEM_OBJECT_IMAGE1D
+        || image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER)
     {
 
         size_t M = maximum_sizes[0];
+        size_t A = max_pixels;
+
+        M = static_cast<size_t>(fmax(1, fmin(A / M, M)));
 
         // Store the size
         sizes[(*numberOfSizes)][0] = M;
@@ -860,6 +864,7 @@ void get_max_sizes(
     {
         switch (image_type)
         {
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
             case CL_MEM_OBJECT_IMAGE1D:
                 log_info(" size[%d] = [%zu] (%g MB image)\n", j, sizes[j][0],
                          raw_pixel_size * sizes[j][0] * sizes[j][1]
@@ -1080,6 +1085,7 @@ cl_ulong get_image_size(image_descriptor const *imageInfo)
     {
         switch (imageInfo->type)
         {
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
             case CL_MEM_OBJECT_IMAGE1D: imageSize = imageInfo->rowPitch; break;
             case CL_MEM_OBJECT_IMAGE2D:
                 imageSize = imageInfo->height * imageInfo->rowPitch;
@@ -2317,6 +2323,7 @@ int debug_find_vector_in_image(void *imagePtr, image_descriptor *imageInfo,
 
     switch (imageInfo->type)
     {
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
         case CL_MEM_OBJECT_IMAGE1D:
             width = (imageInfo->width >> lod) ? (imageInfo->width >> lod) : 1;
             height = 1;
@@ -3513,6 +3520,7 @@ void copy_image_data(image_descriptor *srcImageInfo,
 
         switch (srcImageInfo->type)
         {
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
             case CL_MEM_OBJECT_IMAGE1D:
                 src_lod = sourcePos[1];
                 sourcePos_lod[1] = sourcePos_lod[2] = 0;
@@ -3558,6 +3566,7 @@ void copy_image_data(image_descriptor *srcImageInfo,
         size_t dst_height_lod = 1 /*dstImageInfo->height*/;
         switch (dstImageInfo->type)
         {
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
             case CL_MEM_OBJECT_IMAGE1D:
                 dst_lod = destPos[1];
                 destPos_lod[1] = destPos_lod[2] = 0;
@@ -4022,6 +4031,7 @@ cl_ulong compute_mipmapped_image_size(image_descriptor imageInfo)
                 retSize += (cl_ulong)curr_width * curr_height
                     * get_pixel_size(imageInfo.format);
                 break;
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
             case CL_MEM_OBJECT_IMAGE1D:
                 retSize +=
                     (cl_ulong)curr_width * get_pixel_size(imageInfo.format);
@@ -4043,6 +4053,7 @@ cl_ulong compute_mipmapped_image_size(image_descriptor imageInfo)
             case CL_MEM_OBJECT_IMAGE2D:
             case CL_MEM_OBJECT_IMAGE2D_ARRAY:
                 curr_height = curr_height >> 1 ? curr_height >> 1 : 1;
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
             case CL_MEM_OBJECT_IMAGE1D:
             case CL_MEM_OBJECT_IMAGE1D_ARRAY:
                 curr_width = curr_width >> 1 ? curr_width >> 1 : 1;
@@ -4080,6 +4091,7 @@ size_t compute_mip_level_offset(image_descriptor *imageInfo, size_t lod)
                 retOffset +=
                     (size_t)width * height * get_pixel_size(imageInfo->format);
                 break;
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
             case CL_MEM_OBJECT_IMAGE1D:
                 retOffset += (size_t)width * get_pixel_size(imageInfo->format);
                 break;
@@ -4092,6 +4104,7 @@ size_t compute_mip_level_offset(image_descriptor *imageInfo, size_t lod)
             case CL_MEM_OBJECT_IMAGE2D:
             case CL_MEM_OBJECT_IMAGE2D_ARRAY:
                 height = (height >> 1) ? (height >> 1) : 1;
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
             case CL_MEM_OBJECT_IMAGE1D_ARRAY:
             case CL_MEM_OBJECT_IMAGE1D: width = (width >> 1) ? (width >> 1) : 1;
         }

test_conformance/images/clCopyImage/CMakeLists.txt

@@ -10,6 +10,7 @@ set(${MODULE_NAME}_SOURCES
     test_copy_2D_array.cpp
     test_copy_3D.cpp
     test_copy_3D_2D_array.cpp
+    test_copy_1D_buffer.cpp
     test_copy_generic.cpp
     test_loops.cpp
     ../common.cpp

test_conformance/images/clCopyImage/main.cpp

@@ -45,6 +45,21 @@ int test_3D(cl_device_id device, cl_context context, cl_command_queue queue, int
 {
     return test_image_set( device, context, queue, k3D );
 }
+int test_1Dbuffer(cl_device_id device, cl_context context,
+                  cl_command_queue queue, int num_elements)
+{
+    return test_image_set(device, context, queue, k1DBuffer);
+}
+int test_1DTo1Dbuffer(cl_device_id device, cl_context context,
+                      cl_command_queue queue, int num_elements)
+{
+    return test_image_set(device, context, queue, k1DTo1DBuffer);
+}
+int test_1DbufferTo1D(cl_device_id device, cl_context context,
+                      cl_command_queue queue, int num_elements)
+{
+    return test_image_set(device, context, queue, k1DBufferTo1D);
+}
 int test_1Darray(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
 {
     return test_image_set( device, context, queue, k1DArray );
@@ -79,17 +94,11 @@ int test_3Dto2Darray(cl_device_id device, cl_context context, cl_command_queue q
 }
 
 test_definition test_list[] = {
-    ADD_TEST( 1D ),
+    ADD_TEST(1D),           ADD_TEST(2D),           ADD_TEST(3D),
-    ADD_TEST( 2D ),
+    ADD_TEST(1Darray),      ADD_TEST(2Darray),      ADD_TEST(2Dto3D),
-    ADD_TEST( 3D ),
+    ADD_TEST(3Dto2D),       ADD_TEST(2Darrayto2D),  ADD_TEST(2Dto2Darray),
-    ADD_TEST( 1Darray ),
+    ADD_TEST(2Darrayto3D),  ADD_TEST(3Dto2Darray),  ADD_TEST(1Dbuffer),
-    ADD_TEST( 2Darray ),
+    ADD_TEST(1DTo1Dbuffer), ADD_TEST(1DbufferTo1D),
-    ADD_TEST( 2Dto3D ),
-    ADD_TEST( 3Dto2D ),
-    ADD_TEST( 2Darrayto2D ),
-    ADD_TEST( 2Dto2Darray ),
-    ADD_TEST( 2Darrayto3D ),
-    ADD_TEST( 3Dto2Darray ),
 };
 
 const int test_num = ARRAY_SIZE( test_list );

test_conformance/images/clCopyImage/test_copy_1D_buffer.cpp

@@ -0,0 +1,516 @@
+//
+// Copyright (c) 2023 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "../testBase.h"
+
+extern int test_copy_image_generic(cl_context context, cl_command_queue queue,
+                                   image_descriptor *srcImageInfo,
+                                   image_descriptor *dstImageInfo,
+                                   const size_t sourcePos[],
+                                   const size_t destPos[],
+                                   const size_t regionSize[], MTdata d);
+
+int test_copy_image_size_1D_buffer(cl_context context, cl_command_queue queue,
+                                   image_descriptor *srcImageInfo,
+                                   image_descriptor *dstImageInfo, MTdata d)
+{
+    size_t sourcePos[3], destPos[3], regionSize[3];
+    int ret = 0, retCode;
+    size_t width_lod = srcImageInfo->width;
+
+    // First, try just a full covering region
+    sourcePos[0] = sourcePos[1] = sourcePos[2] = 0;
+    destPos[0] = destPos[1] = destPos[2] = 0;
+    regionSize[0] = srcImageInfo->width;
+    regionSize[1] = 1;
+    regionSize[2] = 1;
+
+    retCode =
+        test_copy_image_generic(context, queue, srcImageInfo, dstImageInfo,
+                                sourcePos, destPos, regionSize, d);
+    if (retCode < 0)
+        return retCode;
+    else
+        ret += retCode;
+
+    // Now try a sampling of different random regions
+    for (int i = 0; i < 8; i++)
+    {
+        // Pick a random size
+        regionSize[0] = (width_lod > 8)
+            ? (size_t)random_in_range(8, (int)width_lod - 1, d)
+            : width_lod;
+
+        // Now pick positions within valid ranges
+        sourcePos[0] = (width_lod > regionSize[0]) ? (size_t)random_in_range(
+                           0, (int)(width_lod - regionSize[0] - 1), d)
+                                                   : 0;
+        destPos[0] = (width_lod > regionSize[0]) ? (size_t)random_in_range(
+                         0, (int)(width_lod - regionSize[0] - 1), d)
+                                                 : 0;
+
+
+        // Go for it!
+        retCode =
+            test_copy_image_generic(context, queue, srcImageInfo, srcImageInfo,
+                                    sourcePos, destPos, regionSize, d);
+        if (retCode < 0)
+            return retCode;
+        else
+            ret += retCode;
+    }
+
+    return ret;
+}
+
+int test_copy_image_set_1D_buffer(cl_device_id device, cl_context context,
+                                  cl_command_queue queue,
+                                  cl_image_format *format)
+{
+    size_t maxWidth;
+    cl_ulong maxAllocSize, memSize;
+    image_descriptor imageInfo = { 0 };
+    RandomSeed seed(gRandomSeed);
+    size_t pixelSize;
+
+    if (gTestMipmaps)
+    {
+        // 1D image buffers don't support mipmaps
+        // https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html#cl_khr_mipmap_image
+        return 0;
+    }
+
+    imageInfo.format = format;
+    imageInfo.height = imageInfo.depth = imageInfo.arraySize =
+        imageInfo.slicePitch = 0;
+    imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+    pixelSize = get_pixel_size(imageInfo.format);
+
+    int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE,
+                                sizeof(maxWidth), &maxWidth, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
+                             sizeof(maxAllocSize), &maxAllocSize, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize),
+                             &memSize, NULL);
+    test_error(error, "Unable to get max image 1D buffer size from device");
+
+    if (memSize > (cl_ulong)SIZE_MAX)
+    {
+        memSize = (cl_ulong)SIZE_MAX;
+        maxAllocSize = (cl_ulong)SIZE_MAX;
+    }
+
+    if (gTestSmallImages)
+    {
+        for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++)
+        {
+            size_t rowPadding = gEnablePitch ? 48 : 0;
+            imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+            if (gEnablePitch)
+            {
+                do
+                {
+                    rowPadding++;
+                    imageInfo.rowPitch =
+                        imageInfo.width * pixelSize + rowPadding;
+                } while ((imageInfo.rowPitch % pixelSize) != 0);
+            }
+
+            if (gDebugTrace) log_info("   at size %d\n", (int)imageInfo.width);
+
+            int ret = test_copy_image_size_1D_buffer(context, queue, &imageInfo,
+                                                     &imageInfo, seed);
+            if (ret) return -1;
+        }
+    }
+    else if (gTestMaxImages)
+    {
+        // Try a specific set of maximum sizes
+        size_t numbeOfSizes;
+        size_t sizes[100][3];
+
+        get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1,
+                      maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER,
+                      imageInfo.format);
+
+        for (size_t idx = 0; idx < numbeOfSizes; idx++)
+        {
+            size_t rowPadding = gEnablePitch ? 48 : 0;
+            imageInfo.width = sizes[idx][0];
+            imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+            if (gEnablePitch)
+            {
+                do
+                {
+                    rowPadding++;
+                    imageInfo.rowPitch =
+                        imageInfo.width * pixelSize + rowPadding;
+                } while ((imageInfo.rowPitch % pixelSize) != 0);
+            }
+
+            log_info("Testing %d\n", (int)sizes[idx][0]);
+            if (gDebugTrace)
+                log_info("   at max size %d\n", (int)sizes[idx][0]);
+            if (test_copy_image_size_1D_buffer(context, queue, &imageInfo,
+                                               &imageInfo, seed))
+                return -1;
+        }
+    }
+    else
+    {
+        for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++)
+        {
+            cl_ulong size;
+            size_t rowPadding = gEnablePitch ? 48 : 0;
+            // Loop until we get a size that a) will fit in the max alloc size
+            // and b) that an allocation of that image, the result array, plus
+            // offset arrays, will fit in the global ram space
+            do
+            {
+                imageInfo.width =
+                    (size_t)random_log_in_range(16, (int)maxWidth / 32, seed);
+
+                imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+                if (gEnablePitch)
+                {
+                    do
+                    {
+                        rowPadding++;
+                        imageInfo.rowPitch =
+                            imageInfo.width * pixelSize + rowPadding;
+                    } while ((imageInfo.rowPitch % pixelSize) != 0);
+                }
+
+                size = (size_t)imageInfo.rowPitch * 4;
+            } while (size > maxAllocSize || (size * 3) > memSize);
+
+            if (gDebugTrace)
+            {
+                log_info("   at size %d (row pitch %d) out of %d\n",
+                         (int)imageInfo.width, (int)imageInfo.rowPitch,
+                         (int)maxWidth);
+            }
+
+            int ret = test_copy_image_size_1D_buffer(context, queue, &imageInfo,
+                                                     &imageInfo, seed);
+            if (ret) return -1;
+        }
+    }
+
+    return 0;
+}
+
+int test_copy_image_set_1D_1D_buffer(cl_device_id device, cl_context context,
+                                     cl_command_queue queue,
+                                     cl_image_format *format)
+{
+    size_t maxWidth;
+    cl_ulong maxAllocSize, memSize;
+    image_descriptor imageInfo = { 0 };
+    RandomSeed seed(gRandomSeed);
+    size_t pixelSize;
+
+    if (gTestMipmaps)
+    {
+        // 1D image buffers don't support mipmaps
+        // https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html#cl_khr_mipmap_image
+        return 0;
+    }
+
+    imageInfo.format = format;
+    imageInfo.height = imageInfo.depth = imageInfo.arraySize =
+        imageInfo.slicePitch = 0;
+    imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+    pixelSize = get_pixel_size(imageInfo.format);
+
+    int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH,
+                                sizeof(maxWidth), &maxWidth, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
+                             sizeof(maxAllocSize), &maxAllocSize, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize),
+                             &memSize, NULL);
+    test_error(error, "Unable to get max image 1D buffer size from device");
+
+    if (memSize > (cl_ulong)SIZE_MAX)
+    {
+        memSize = (cl_ulong)SIZE_MAX;
+        maxAllocSize = (cl_ulong)SIZE_MAX;
+    }
+
+    if (gTestSmallImages)
+    {
+        for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++)
+        {
+            size_t rowPadding = gEnablePitch ? 48 : 0;
+            imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+            if (gEnablePitch)
+            {
+                do
+                {
+                    rowPadding++;
+                    imageInfo.rowPitch =
+                        imageInfo.width * pixelSize + rowPadding;
+                } while ((imageInfo.rowPitch % pixelSize) != 0);
+            }
+
+            if (gDebugTrace) log_info("   at size %d\n", (int)imageInfo.width);
+
+            image_descriptor srcImageInfo = imageInfo;
+            srcImageInfo.type = CL_MEM_OBJECT_IMAGE1D;
+
+            int ret = test_copy_image_size_1D_buffer(
+                context, queue, &srcImageInfo, &imageInfo, seed);
+            if (ret) return -1;
+        }
+    }
+    else if (gTestMaxImages)
+    {
+        // Try a specific set of maximum sizes
+        size_t numbeOfSizes;
+        size_t sizes[100][3];
+
+        get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1,
+                      maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER,
+                      imageInfo.format);
+
+        for (size_t idx = 0; idx < numbeOfSizes; idx++)
+        {
+            size_t rowPadding = gEnablePitch ? 48 : 0;
+            imageInfo.width = sizes[idx][0];
+            imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+            if (gEnablePitch)
+            {
+                do
+                {
+                    rowPadding++;
+                    imageInfo.rowPitch =
+                        imageInfo.width * pixelSize + rowPadding;
+                } while ((imageInfo.rowPitch % pixelSize) != 0);
+            }
+
+            log_info("Testing %d\n", (int)sizes[idx][0]);
+            if (gDebugTrace)
+                log_info("   at max size %d\n", (int)sizes[idx][0]);
+
+            image_descriptor srcImageInfo = imageInfo;
+            srcImageInfo.type = CL_MEM_OBJECT_IMAGE1D;
+
+            if (test_copy_image_size_1D_buffer(context, queue, &srcImageInfo,
+                                               &imageInfo, seed))
+                return -1;
+        }
+    }
+    else
+    {
+        for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++)
+        {
+            cl_ulong size;
+            size_t rowPadding = gEnablePitch ? 48 : 0;
+            // Loop until we get a size that a) will fit in the max alloc size
+            // and b) that an allocation of that image, the result array, plus
+            // offset arrays, will fit in the global ram space
+            do
+            {
+                imageInfo.width =
+                    (size_t)random_log_in_range(16, (int)maxWidth / 32, seed);
+
+                imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+                if (gEnablePitch)
+                {
+                    do
+                    {
+                        rowPadding++;
+                        imageInfo.rowPitch =
+                            imageInfo.width * pixelSize + rowPadding;
+                    } while ((imageInfo.rowPitch % pixelSize) != 0);
+                }
+
+                size = (size_t)imageInfo.rowPitch * 4;
+            } while (size > maxAllocSize || (size * 3) > memSize);
+
+            if (gDebugTrace)
+            {
+                log_info("   at size %d (row pitch %d) out of %d\n",
+                         (int)imageInfo.width, (int)imageInfo.rowPitch,
+                         (int)maxWidth);
+            }
+
+            image_descriptor srcImageInfo = imageInfo;
+            srcImageInfo.type = CL_MEM_OBJECT_IMAGE1D;
+
+            int ret = test_copy_image_size_1D_buffer(
+                context, queue, &srcImageInfo, &imageInfo, seed);
+            if (ret) return -1;
+        }
+    }
+
+    return 0;
+}
+
+int test_copy_image_set_1D_buffer_1D(cl_device_id device, cl_context context,
+                                     cl_command_queue queue,
+                                     cl_image_format *format)
+{
+    size_t maxWidth;
+    cl_ulong maxAllocSize, memSize;
+    image_descriptor imageInfo = { 0 };
+    RandomSeed seed(gRandomSeed);
+    size_t pixelSize;
+
+    if (gTestMipmaps)
+    {
+        // 1D image buffers don't support mipmaps
+        // https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html#cl_khr_mipmap_image
+        return 0;
+    }
+
+    imageInfo.format = format;
+    imageInfo.height = imageInfo.depth = imageInfo.arraySize =
+        imageInfo.slicePitch = 0;
+    imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+    pixelSize = get_pixel_size(imageInfo.format);
+
+    int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH,
+                                sizeof(maxWidth), &maxWidth, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
+                             sizeof(maxAllocSize), &maxAllocSize, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize),
+                             &memSize, NULL);
+    test_error(error, "Unable to get max image 1D buffer size from device");
+
+    if (memSize > (cl_ulong)SIZE_MAX)
+    {
+        memSize = (cl_ulong)SIZE_MAX;
+        maxAllocSize = (cl_ulong)SIZE_MAX;
+    }
+
+    if (gTestSmallImages)
+    {
+        for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++)
+        {
+            size_t rowPadding = gEnablePitch ? 48 : 0;
+            imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+            if (gEnablePitch)
+            {
+                do
+                {
+                    rowPadding++;
+                    imageInfo.rowPitch =
+                        imageInfo.width * pixelSize + rowPadding;
+                } while ((imageInfo.rowPitch % pixelSize) != 0);
+            }
+
+            if (gDebugTrace) log_info("   at size %d\n", (int)imageInfo.width);
+
+            image_descriptor dstImageInfo = imageInfo;
+            dstImageInfo.type = CL_MEM_OBJECT_IMAGE1D;
+
+            int ret = test_copy_image_size_1D_buffer(context, queue, &imageInfo,
+                                                     &dstImageInfo, seed);
+            if (ret) return -1;
+        }
+    }
+    else if (gTestMaxImages)
+    {
+        // Try a specific set of maximum sizes
+        size_t numbeOfSizes;
+        size_t sizes[100][3];
+
+        get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1,
+                      maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER,
+                      imageInfo.format);
+
+        for (size_t idx = 0; idx < numbeOfSizes; idx++)
+        {
+            size_t rowPadding = gEnablePitch ? 48 : 0;
+            imageInfo.width = sizes[idx][0];
+            imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+            if (gEnablePitch)
+            {
+                do
+                {
+                    rowPadding++;
+                    imageInfo.rowPitch =
+                        imageInfo.width * pixelSize + rowPadding;
+                } while ((imageInfo.rowPitch % pixelSize) != 0);
+            }
+
+            log_info("Testing %d\n", (int)sizes[idx][0]);
+            if (gDebugTrace)
+                log_info("   at max size %d\n", (int)sizes[idx][0]);
+
+            image_descriptor dstImageInfo = imageInfo;
+            dstImageInfo.type = CL_MEM_OBJECT_IMAGE1D;
+
+            if (test_copy_image_size_1D_buffer(context, queue, &imageInfo,
+                                               &dstImageInfo, seed))
+                return -1;
+        }
+    }
+    else
+    {
+        for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++)
+        {
+            cl_ulong size;
+            size_t rowPadding = gEnablePitch ? 48 : 0;
+            // Loop until we get a size that a) will fit in the max alloc size
+            // and b) that an allocation of that image, the result array, plus
+            // offset arrays, will fit in the global ram space
+            do
+            {
+                imageInfo.width =
+                    (size_t)random_log_in_range(16, (int)maxWidth / 32, seed);
+
+                imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+                if (gEnablePitch)
+                {
+                    do
+                    {
+                        rowPadding++;
+                        imageInfo.rowPitch =
+                            imageInfo.width * pixelSize + rowPadding;
+                    } while ((imageInfo.rowPitch % pixelSize) != 0);
+                }
+
+                size = (size_t)imageInfo.rowPitch * 4;
+            } while (size > maxAllocSize || (size * 3) > memSize);
+
+            if (gDebugTrace)
+            {
+                log_info("   at size %d (row pitch %d) out of %d\n",
+                         (int)imageInfo.width, (int)imageInfo.rowPitch,
+                         (int)maxWidth);
+            }
+
+            image_descriptor dstImageInfo = imageInfo;
+            dstImageInfo.type = CL_MEM_OBJECT_IMAGE1D;
+
+            int ret = test_copy_image_size_1D_buffer(context, queue, &imageInfo,
+                                                     &dstImageInfo, seed);
+            if (ret) return -1;
+        }
+    }
+
+    return 0;
+}

test_conformance/images/clCopyImage/test_copy_generic.cpp

@@ -14,12 +14,18 @@
 // limitations under the License.
 //
 #include "../testBase.h"
+#include <CL/cl.h>
 
 static void CL_CALLBACK free_pitch_buffer( cl_mem image, void *buf )
 {
     free( buf );
 }
 
+static void CL_CALLBACK release_cl_buffer(cl_mem image, void *buf)
+{
+    clReleaseMemObject((cl_mem)buf);
+}
+
 cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr<char>& data, image_descriptor *imageInfo, int *error )
 {
     cl_mem img;
@@ -69,6 +75,71 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
             if ( gEnablePitch )
                 host_ptr = malloc( imageInfo->arraySize * imageInfo->slicePitch );
             break;
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+            if (gDebugTrace)
+                log_info(" - Creating 1D buffer image %d ...\n",
+                         (int)imageInfo->width);
+            {
+                cl_int err;
+                cl_mem_flags buffer_flags = CL_MEM_READ_WRITE;
+                if (gEnablePitch)
+                {
+                    cl_device_id device;
+                    err =
+                        clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE,
+                                              sizeof(device), &device, nullptr);
+                    if (err != CL_SUCCESS)
+                    {
+                        log_error(
+                            "Error: Could not get CL_QUEUE_DEVICE from queue");
+                        return NULL;
+                    }
+                    char major_version;
+                    err = clGetDeviceInfo(device, CL_DEVICE_VERSION,
+                                          sizeof(major_version), &major_version,
+                                          nullptr);
+                    if (err != CL_SUCCESS)
+                    {
+                        log_error("Error: Could not get CL_DEVICE_VERSION from "
+                                  "device");
+                        return NULL;
+                    }
+                    if (major_version == '1')
+                    {
+                        host_ptr = malloc(imageInfo->rowPitch);
+                    }
+                    else
+                    {
+                        cl_uint base_address_alignment = 0;
+                        err = clGetDeviceInfo(
+                            device, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT,
+                            sizeof(base_address_alignment),
+                            &base_address_alignment, nullptr);
+                        if (err != CL_SUCCESS)
+                        {
+                            log_error("ERROR: Could not get "
+                                      "CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT "
+                                      "from device");
+                            return NULL;
+                        }
+                        host_ptr = align_malloc(imageInfo->rowPitch,
+                                                base_address_alignment);
+                    }
+                    buffer_flags |= CL_MEM_USE_HOST_PTR;
+                }
+
+                cl_mem buffer = clCreateBuffer(context, buffer_flags,
+                                               imageInfo->rowPitch, NULL, &err);
+                if (err != CL_SUCCESS)
+                {
+                    log_error("ERROR: Could not create buffer for 1D buffer "
+                              "image. %ld bytes\n",
+                              imageInfo->width);
+                    return NULL;
+                }
+                imageDesc.buffer = buffer;
+            }
+            break;
     }
 
     if ( gDebugTrace && gTestMipmaps )
@@ -81,8 +152,11 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
             log_error( "ERROR: Unable to create backing store for pitched 3D image. %ld bytes\n",  imageInfo->depth * imageInfo->slicePitch );
             return NULL;
         }
+        if (imageInfo->type != CL_MEM_OBJECT_IMAGE1D_BUFFER)
+        {
             mem_flags = CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR;
         }
+    }
 
     img = clCreateImage(context, mem_flags, imageInfo->format, &imageDesc, host_ptr, error);
 
@@ -103,6 +177,21 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
             free(host_ptr);
     }
 
+    if (imageDesc.buffer != NULL)
+    {
+        int callbackError = clSetMemObjectDestructorCallback(
+            img, release_cl_buffer, imageDesc.buffer);
+        if (callbackError != CL_SUCCESS)
+        {
+            log_error("Error: Unable to attach destructor callback to 1d "
+                      "buffer image. Err: %d\n",
+                      callbackError);
+            clReleaseMemObject(imageDesc.buffer);
+            clReleaseMemObject(img);
+            return NULL;
+        }
+    }
+
     if ( *error != CL_SUCCESS )
     {
         long long unsigned imageSize = get_image_size_mb(imageInfo);
@@ -141,6 +230,12 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
                           (int)imageInfo->arraySize, imageSize,
                           IGetErrorString(*error));
                 break;
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+                log_error(
+                    "ERROR: Unable to create 1D buffer image of size %d (%llu "
+                    "MB):(%s)",
+                    (int)imageInfo->width, imageSize, IGetErrorString(*error));
+                break;
         }
         log_error("ERROR: and %llu mip levels\n", (unsigned long long) imageInfo->num_mip_levels);
         return NULL;
@@ -161,6 +256,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
             height = imageInfo->arraySize;
             depth = 1;
             break;
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
         case CL_MEM_OBJECT_IMAGE1D:
             height = depth = 1;
             break;
@@ -196,6 +292,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
                 case CL_MEM_OBJECT_IMAGE1D_ARRAY:
                     origin[ 2 ] =  lod;
                     break;
+                case CL_MEM_OBJECT_IMAGE1D_BUFFER:
                 case CL_MEM_OBJECT_IMAGE1D:
                     origin[ 1 ] = lod;
                     break;
@@ -210,6 +307,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
                 case CL_MEM_OBJECT_IMAGE2D:
                     height = ( imageInfo->height >> lod ) ? (imageInfo->height >> lod) : 1;
                 case CL_MEM_OBJECT_IMAGE1D_ARRAY:
+                case CL_MEM_OBJECT_IMAGE1D_BUFFER:
                 case CL_MEM_OBJECT_IMAGE1D:
                     width = ( imageInfo->width >> lod ) ? (imageInfo->width >> lod) : 1;
             }
@@ -260,6 +358,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
                     break;
                 case CL_MEM_OBJECT_IMAGE1D_ARRAY:
                 case CL_MEM_OBJECT_IMAGE1D:
+                case CL_MEM_OBJECT_IMAGE1D_BUFFER:
                     dstPitch2D = mappedSlice;
                     break;
             }
@@ -391,8 +490,7 @@ int test_copy_image_generic( cl_context context, cl_command_queue queue, image_d
         switch(dstImageInfo->type)
         {
             case CL_MEM_OBJECT_IMAGE1D:
-                dst_lod = destPos[1];
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER: dst_lod = destPos[1]; break;
-                break;
             case CL_MEM_OBJECT_IMAGE1D_ARRAY:
             case CL_MEM_OBJECT_IMAGE2D:
                 dst_lod = destPos[2];
@@ -407,6 +505,7 @@ int test_copy_image_generic( cl_context context, cl_command_queue queue, image_d
     }
     switch (dstImageInfo->type)
     {
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
         case CL_MEM_OBJECT_IMAGE1D:
             if( gTestMipmaps )
                 origin[ 1 ] = dst_lod;
@@ -542,6 +641,7 @@ int test_copy_image_generic( cl_context context, cl_command_queue queue, image_d
             secondDim = dstImageInfo->height;
             break;
         }
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
         case CL_MEM_OBJECT_IMAGE1D: {
             break;
         }

test_conformance/images/clCopyImage/test_loops.cpp

@@ -24,6 +24,18 @@ extern int test_copy_image_set_2D_array( cl_device_id device, cl_context context
 extern int test_copy_image_set_2D_3D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, bool reverse );
 extern int test_copy_image_set_2D_2D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, bool reverse );
 extern int test_copy_image_set_3D_2D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, bool reverse );
+extern int test_copy_image_set_1D_buffer(cl_device_id device,
+                                         cl_context context,
+                                         cl_command_queue queue,
+                                         cl_image_format *format);
+extern int test_copy_image_set_1D_1D_buffer(cl_device_id device,
+                                            cl_context context,
+                                            cl_command_queue queue,
+                                            cl_image_format *format);
+extern int test_copy_image_set_1D_buffer_1D(cl_device_id device,
+                                            cl_context context,
+                                            cl_command_queue queue,
+                                            cl_image_format *format);
 
 int test_image_type( cl_device_id device, cl_context context, cl_command_queue queue, MethodsToTest testMethod, cl_mem_flags flags )
 {
@@ -87,6 +99,18 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q
             name = "3D -> 2D array";
             imageType = CL_MEM_OBJECT_IMAGE3D;
             break;
+        case k1DBuffer:
+            name = "1D buffer -> 1D buffer";
+            imageType = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+            break;
+        case k1DTo1DBuffer:
+            name = "1D -> 1D buffer";
+            imageType = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+            break;
+        case k1DBufferTo1D:
+            name = "1D buffer -> 1D";
+            imageType = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+            break;
     }
 
     if(gTestMipmaps)
@@ -138,6 +162,16 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q
             test_return = test_copy_image_set_3D_2D_array( device, context, queue, &formatList[ i ], true);
         else if( testMethod == k3DTo2DArray)
             test_return = test_copy_image_set_3D_2D_array( device, context, queue, &formatList[ i ], false);
+        else if (testMethod == k1DBuffer)
+            test_return = test_copy_image_set_1D_buffer(device, context, queue,
+                                                        &formatList[i]);
+        else if (testMethod == k1DBufferTo1D)
+            test_return = test_copy_image_set_1D_buffer_1D(
+                device, context, queue, &formatList[i]);
+        else if (testMethod == k1DTo1DBuffer)
+            test_return = test_copy_image_set_1D_1D_buffer(
+                device, context, queue, &formatList[i]);
+
 
         if (test_return) {
             gFailCount++;

test_conformance/images/clFillImage/CMakeLists.txt

@@ -4,6 +4,7 @@ set(MODULE_NAME CL_FILL_IMAGES)
 set(${MODULE_NAME}_SOURCES
     main.cpp
     test_fill_1D.cpp
+    test_fill_1D_buffer.cpp
     test_fill_1D_array.cpp
     test_fill_2D.cpp
     test_fill_2D_array.cpp

test_conformance/images/clFillImage/main.cpp

@@ -52,13 +52,15 @@ int test_2Darray(cl_device_id device, cl_context context, cl_command_queue queue
 {
     return test_image_set(device, context, queue, k2DArray);
 }
+int test_1Dbuffer(cl_device_id device, cl_context context,
+                  cl_command_queue queue, int num_elements)
+{
+    return test_image_set(device, context, queue, k1DBuffer);
+}
 
 test_definition test_list[] = {
-    ADD_TEST( 1D ),
+    ADD_TEST(1D),      ADD_TEST(2D),      ADD_TEST(3D),
-    ADD_TEST( 2D ),
+    ADD_TEST(1Darray), ADD_TEST(2Darray), ADD_TEST(1Dbuffer),
-    ADD_TEST( 3D ),
-    ADD_TEST( 1Darray ),
-    ADD_TEST( 2Darray ),
 };
 
 const int test_num = ARRAY_SIZE( test_list );

test_conformance/images/clFillImage/test_fill_1D_buffer.cpp

@@ -0,0 +1,203 @@
+//
+// Copyright (c) 2023 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "../testBase.h"
+
+// Defined in test_fill_2D_3D.cpp
+extern int test_fill_image_generic(cl_context context, cl_command_queue queue,
+                                   image_descriptor *imageInfo,
+                                   const size_t origin[], const size_t region[],
+                                   ExplicitType outputType, MTdata d);
+
+
+int test_fill_image_size_1D_buffer(cl_context context, cl_command_queue queue,
+                                   image_descriptor *imageInfo,
+                                   ExplicitType outputType, MTdata d)
+{
+    size_t origin[3], region[3];
+    int ret = 0, retCode;
+
+    // First, try just a full covering region fill
+    origin[0] = origin[1] = origin[2] = 0;
+    region[0] = imageInfo->width;
+    region[1] = 1;
+    region[2] = 1;
+
+    retCode = test_fill_image_generic(context, queue, imageInfo, origin, region,
+                                      outputType, d);
+    if (retCode < 0)
+        return retCode;
+    else
+        ret += retCode;
+
+    // Now try a sampling of different random regions
+    for (int i = 0; i < 8; i++)
+    {
+        // Pick a random size
+        region[0] = (imageInfo->width > 8)
+            ? (size_t)random_in_range(8, (int)imageInfo->width - 1, d)
+            : imageInfo->width;
+
+        // Now pick positions within valid ranges
+        origin[0] = (imageInfo->width > region[0]) ? (size_t)random_in_range(
+                        0, (int)(imageInfo->width - region[0] - 1), d)
+                                                   : 0;
+
+        // Go for it!
+        retCode = test_fill_image_generic(context, queue, imageInfo, origin,
+                                          region, outputType, d);
+        if (retCode < 0)
+            return retCode;
+        else
+            ret += retCode;
+    }
+
+    return ret;
+}
+
+
+int test_fill_image_set_1D_buffer(cl_device_id device, cl_context context,
+                                  cl_command_queue queue,
+                                  cl_image_format *format,
+                                  ExplicitType outputType)
+{
+    size_t maxWidth;
+    cl_ulong maxAllocSize, memSize;
+    image_descriptor imageInfo = { 0 };
+    RandomSeed seed(gRandomSeed);
+    const size_t rowPadding_default = 48;
+    size_t rowPadding = gEnablePitch ? rowPadding_default : 0;
+    size_t pixelSize;
+
+    memset(&imageInfo, 0x0, sizeof(image_descriptor));
+    imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+    imageInfo.format = format;
+    pixelSize = get_pixel_size(imageInfo.format);
+
+    int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE,
+                                sizeof(maxWidth), &maxWidth, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
+                             sizeof(maxAllocSize), &maxAllocSize, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize),
+                             &memSize, NULL);
+    test_error(error, "Unable to get max image 2D size from device");
+
+    if (memSize > (cl_ulong)SIZE_MAX)
+    {
+        memSize = (cl_ulong)SIZE_MAX;
+        maxAllocSize = (cl_ulong)SIZE_MAX;
+    }
+
+    if (gTestSmallImages)
+    {
+        for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++)
+        {
+            imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+            if (gEnablePitch)
+            {
+                rowPadding = rowPadding_default;
+                do
+                {
+                    rowPadding++;
+                    imageInfo.rowPitch =
+                        imageInfo.width * pixelSize + rowPadding;
+                } while ((imageInfo.rowPitch % pixelSize) != 0);
+            }
+
+            if (gDebugTrace)
+                log_info("   at size %d,%d\n", (int)imageInfo.width,
+                         (int)imageInfo.height);
+
+            int ret = test_fill_image_size_1D_buffer(context, queue, &imageInfo,
+                                                     outputType, seed);
+            if (ret) return -1;
+        }
+    }
+    else if (gTestMaxImages)
+    {
+        // Try a specific set of maximum sizes
+        size_t numbeOfSizes;
+        size_t sizes[100][3];
+
+        get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1,
+                      maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER,
+                      imageInfo.format);
+
+        for (size_t idx = 0; idx < numbeOfSizes; idx++)
+        {
+            imageInfo.width = sizes[idx][0];
+            imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+            if (gEnablePitch)
+            {
+                rowPadding = rowPadding_default;
+                do
+                {
+                    rowPadding++;
+                    imageInfo.rowPitch =
+                        imageInfo.width * pixelSize + rowPadding;
+                } while ((imageInfo.rowPitch % pixelSize) != 0);
+            }
+
+            log_info("Testing %d\n", (int)sizes[idx][0]);
+            if (gDebugTrace)
+                log_info("   at max size %d\n", (int)sizes[idx][0]);
+            if (test_fill_image_size_1D_buffer(context, queue, &imageInfo,
+                                               outputType, seed))
+                return -1;
+        }
+    }
+    else
+    {
+        for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++)
+        {
+            cl_ulong size;
+            // Loop until we get a size that a) will fit in the max alloc size
+            // and b) that an allocation of that image, the result array, plus
+            // offset arrays, will fit in the global ram space
+            do
+            {
+                imageInfo.width =
+                    (size_t)random_log_in_range(16, (int)maxWidth / 32, seed);
+
+                imageInfo.rowPitch = imageInfo.width * pixelSize + rowPadding;
+
+                if (gEnablePitch)
+                {
+                    rowPadding = rowPadding_default;
+                    do
+                    {
+                        rowPadding++;
+                        imageInfo.rowPitch =
+                            imageInfo.width * pixelSize + rowPadding;
+                    } while ((imageInfo.rowPitch % pixelSize) != 0);
+                }
+
+                size = (size_t)imageInfo.rowPitch * 4;
+            } while (size > maxAllocSize || (size * 3) > memSize);
+
+            if (gDebugTrace)
+                log_info("   at size %d (row pitch %d) out of %d\n",
+                         (int)imageInfo.width, (int)imageInfo.rowPitch,
+                         (int)maxWidth);
+            int ret = test_fill_image_size_1D_buffer(context, queue, &imageInfo,
+                                                     outputType, seed);
+            if (ret) return -1;
+        }
+    }
+
+    return 0;
+}

test_conformance/images/clFillImage/test_fill_generic.cpp

@@ -17,9 +17,27 @@
 
 extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo, int x, int y, int z, float *outData );
 
-static void CL_CALLBACK free_pitch_buffer( cl_mem image, void *buf )
+struct pitch_buffer_data
 {
-    free( buf );
+    void *buf;
+    bool is_aligned;
+};
+static void CL_CALLBACK free_pitch_buffer(cl_mem image, void *data)
+{
+    struct pitch_buffer_data *d = (struct pitch_buffer_data *)data;
+    if (d->is_aligned)
+    {
+        align_free(d->buf);
+    }
+    else
+    {
+        free(d->buf);
+    }
+    free(d);
+}
+static void CL_CALLBACK release_cl_buffer(cl_mem image, void *buf)
+{
+    clReleaseMemObject((cl_mem)buf);
 }
 
 cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr<char>& data, image_descriptor *imageInfo, int *error )
@@ -38,6 +56,26 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
     imageDesc.image_row_pitch = gEnablePitch ? imageInfo->rowPitch : 0;
     imageDesc.image_slice_pitch = gEnablePitch ? imageInfo->slicePitch : 0;
 
+    cl_version version;
+    cl_device_id device;
+    {
+        cl_int err = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE,
+                                           sizeof(device), &device, nullptr);
+        if (err != CL_SUCCESS)
+        {
+            log_error("Error: Could not get CL_QUEUE_DEVICE from queue");
+            return NULL;
+        }
+        err = clGetDeviceInfo(device, CL_DEVICE_NUMERIC_VERSION,
+                              sizeof(version), &version, nullptr);
+        if (err != CL_SUCCESS)
+        {
+            log_error("Error: Could not get CL_DEVICE_NUMERIC_VERSION from "
+                      "device");
+            return NULL;
+        }
+    }
+
     switch (imageInfo->type)
     {
         case CL_MEM_OBJECT_IMAGE1D:
@@ -70,6 +108,50 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
             if ( gEnablePitch )
                 host_ptr = malloc( imageInfo->arraySize * imageInfo->slicePitch );
             break;
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+            if (gDebugTrace)
+                log_info(" - Creating 1D buffer image %d ...\n",
+                         (int)imageInfo->width);
+            {
+                cl_int err;
+                cl_mem_flags buffer_flags = CL_MEM_READ_WRITE;
+                if (gEnablePitch)
+                {
+                    if (CL_VERSION_MAJOR(version) == 1)
+                    {
+                        host_ptr = malloc(imageInfo->rowPitch);
+                    }
+                    else
+                    {
+                        cl_uint base_address_alignment = 0;
+                        err = clGetDeviceInfo(
+                            device, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT,
+                            sizeof(base_address_alignment),
+                            &base_address_alignment, nullptr);
+                        if (err != CL_SUCCESS)
+                        {
+                            log_error("ERROR: Could not get "
+                                      "CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT "
+                                      "from device");
+                            return NULL;
+                        }
+                        host_ptr = align_malloc(imageInfo->rowPitch,
+                                                base_address_alignment);
+                    }
+                    buffer_flags |= CL_MEM_USE_HOST_PTR;
+                }
+                cl_mem buffer = clCreateBuffer(
+                    context, buffer_flags, imageInfo->rowPitch, host_ptr, &err);
+                if (err != CL_SUCCESS)
+                {
+                    log_error("ERROR: Could not create buffer for 1D buffer "
+                              "image. %ld bytes\n",
+                              imageInfo->rowPitch);
+                    return NULL;
+                }
+                imageDesc.buffer = buffer;
+            }
+            break;
     }
 
     if (gEnablePitch)
@@ -79,26 +161,63 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
             log_error( "ERROR: Unable to create backing store for pitched 3D image. %ld bytes\n",  imageInfo->depth * imageInfo->slicePitch );
             return NULL;
         }
+        if (imageInfo->type != CL_MEM_OBJECT_IMAGE1D_BUFFER)
+        {
             mem_flags = CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR;
         }
+    }
 
-    img = clCreateImage(context, mem_flags, imageInfo->format, &imageDesc, host_ptr, error);
+    if (imageInfo->type != CL_MEM_OBJECT_IMAGE1D_BUFFER)
+    {
+        img = clCreateImage(context, mem_flags, imageInfo->format, &imageDesc,
+                            host_ptr, error);
+    }
+    else
+    {
+        img = clCreateImage(context, mem_flags, imageInfo->format, &imageDesc,
+                            nullptr, error);
+    }
 
     if (gEnablePitch)
     {
+        struct pitch_buffer_data *data = (struct pitch_buffer_data *)malloc(
+            sizeof(struct pitch_buffer_data));
+        data->buf = host_ptr;
+        data->is_aligned = (CL_VERSION_MAJOR(version) != 1)
+            && (imageInfo->type == CL_MEM_OBJECT_IMAGE1D_BUFFER);
         if (*error == CL_SUCCESS)
         {
-            int callbackError = clSetMemObjectDestructorCallback( img, free_pitch_buffer, host_ptr );
+            int callbackError =
+                clSetMemObjectDestructorCallback(img, free_pitch_buffer, data);
             if (CL_SUCCESS != callbackError)
             {
-                free( host_ptr );
+                free_pitch_buffer(img, data);
-                log_error( "ERROR: Unable to attach destructor callback to pitched 3D image. Err: %d\n", callbackError );
+                log_error("ERROR: Unable to attach destructor callback to "
+                          "pitched 3D image. Err: %d\n",
+                          callbackError);
                 clReleaseMemObject(img);
                 return NULL;
             }
         }
         else
-            free(host_ptr);
+        {
+            free_pitch_buffer(img, data);
+        }
+    }
+
+    if (imageDesc.buffer != NULL)
+    {
+        int callbackError = clSetMemObjectDestructorCallback(
+            img, release_cl_buffer, imageDesc.buffer);
+        if (callbackError != CL_SUCCESS)
+        {
+            log_error("Error: Unable to attach destructor callback to 1d "
+                      "buffer image. Err: %d\n",
+                      callbackError);
+            clReleaseMemObject(imageDesc.buffer);
+            clReleaseMemObject(img);
+            return NULL;
+        }
     }
 
     if ( *error != CL_SUCCESS )
@@ -122,6 +241,12 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
             case CL_MEM_OBJECT_IMAGE2D_ARRAY:
                 log_error( "ERROR: Unable to create 2D image array of size %d x %d x %d (%llu MB): %s\n", (int)imageInfo->width, (int)imageInfo->height, (int)imageInfo->arraySize, imageSize, IGetErrorString( *error ) );
                 break;
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+                log_error(
+                    "ERROR: Unable to create 1D buffer image of size %d (%llu "
+                    "MB):(%s)",
+                    (int)imageInfo->width, imageSize, IGetErrorString(*error));
+                break;
         }
         return NULL;
     }
@@ -139,6 +264,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
             depth = 1;
             imageSize = imageInfo->rowPitch * imageInfo->arraySize;
             break;
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
         case CL_MEM_OBJECT_IMAGE1D:
             height = depth = 1;
             imageSize = imageInfo->rowPitch;
@@ -194,8 +320,7 @@ cl_mem create_image( cl_context context, cl_command_queue queue, BufferOwningPtr
                 break;
             case CL_MEM_OBJECT_IMAGE1D_ARRAY:
             case CL_MEM_OBJECT_IMAGE1D:
-                dstPitch2D = mappedSlice;
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER: dstPitch2D = mappedSlice; break;
-                break;
         }
 
         for ( size_t z = 0; z < depth; z++ )
@@ -286,6 +411,9 @@ int test_fill_image_generic( cl_context context, cl_command_queue queue, image_d
         case CL_MEM_OBJECT_IMAGE2D_ARRAY:
             dataBytes = imageInfo->arraySize * imageInfo->slicePitch;
             break;
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+            dataBytes = imageInfo->rowPitch;
+            break;
     }
 
     if (dataBytes > imgData.getSize())
@@ -398,6 +526,7 @@ int test_fill_image_generic( cl_context context, cl_command_queue queue, image_d
     size_t imageRegion[ 3 ] = { imageInfo->width, 1, 1 };
     switch (imageInfo->type)
     {
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
         case CL_MEM_OBJECT_IMAGE1D:
             break;
         case CL_MEM_OBJECT_IMAGE2D:
@@ -437,6 +566,7 @@ int test_fill_image_generic( cl_context context, cl_command_queue queue, image_d
     size_t secondDim = 1;
 
     switch (imageInfo->type) {
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
         case CL_MEM_OBJECT_IMAGE1D:
             secondDim = 1;
             thirdDim = 1;
@@ -458,7 +588,8 @@ int test_fill_image_generic( cl_context context, cl_command_queue queue, image_d
             thirdDim = imageInfo->arraySize;
             break;
         default:
-        log_error("Test error: unhandled image type at %s:%d\n",__FILE__,__LINE__);
+            log_error("Test error: unhandled image type at %s:%d\n", __FILE__,
+                      __LINE__);
     };
 
     // Count the number of bytes successfully matched
@@ -485,7 +616,9 @@ int test_fill_image_generic( cl_context context, cl_command_queue queue, image_d
 
             total_matched += scanlineSize;
             sourcePtr += imageInfo->rowPitch;
-            if((imageInfo->type == CL_MEM_OBJECT_IMAGE1D_ARRAY || imageInfo->type == CL_MEM_OBJECT_IMAGE1D))
+            if ((imageInfo->type == CL_MEM_OBJECT_IMAGE1D_ARRAY
+                 || imageInfo->type == CL_MEM_OBJECT_IMAGE1D
+                 || imageInfo->type == CL_MEM_OBJECT_IMAGE1D_BUFFER))
                 destPtr += mappedSlice;
             else
             destPtr += mappedRow;

test_conformance/images/clFillImage/test_loops.cpp

@@ -23,6 +23,11 @@ extern int test_fill_image_set_2D( cl_device_id device, cl_context context, cl_c
 extern int test_fill_image_set_3D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType outputType );
 extern int test_fill_image_set_1D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType outputType );
 extern int test_fill_image_set_2D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType outputType );
+extern int test_fill_image_set_1D_buffer(cl_device_id device,
+                                         cl_context context,
+                                         cl_command_queue queue,
+                                         cl_image_format *format,
+                                         ExplicitType outputType);
 typedef int (*test_func)(cl_device_id device, cl_context context,
                          cl_command_queue queue, cl_image_format *format,
                          ExplicitType outputType);
@@ -60,6 +65,11 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q
             imageType = CL_MEM_OBJECT_IMAGE3D;
             test_fn = &test_fill_image_set_3D;
             break;
+        case k1DBuffer:
+            name = "1D Image Buffer Fill";
+            imageType = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+            test_fn = &test_fill_image_set_1D_buffer;
+            break;
         default: log_error("Unhandled method\n"); return -1;
     }
 

test_conformance/images/clGetInfo/CMakeLists.txt

@@ -7,6 +7,7 @@ set(${MODULE_NAME}_SOURCES
     test_2D.cpp
     test_loops.cpp
     test_3D.cpp
+    test_1D_buffer.cpp
     ../common.cpp
 )
 

test_conformance/images/clGetInfo/main.cpp

@@ -54,13 +54,15 @@ int test_2Darray(cl_device_id device, cl_context context, cl_command_queue queue
 {
     return test_image_set( device, context, CL_MEM_OBJECT_IMAGE2D_ARRAY );
 }
+int test_1Dbuffer(cl_device_id device, cl_context context,
+                  cl_command_queue queue, int num_elements)
+{
+    return test_image_set(device, context, CL_MEM_OBJECT_IMAGE1D_BUFFER);
+}
 
 test_definition test_list[] = {
-    ADD_TEST( 1D ),
+    ADD_TEST(1D),      ADD_TEST(2D),      ADD_TEST(3D),
-    ADD_TEST( 2D ),
+    ADD_TEST(1Darray), ADD_TEST(2Darray), ADD_TEST(1Dbuffer),
-    ADD_TEST( 3D ),
-    ADD_TEST( 1Darray ),
-    ADD_TEST( 2Darray ),
 };
 
 const int test_num = ARRAY_SIZE( test_list );

test_conformance/images/clGetInfo/test_1D_buffer.cpp

@@ -0,0 +1,128 @@
+//
+// Copyright (c) 2023 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "../testBase.h"
+#include <CL/cl.h>
+
+extern int test_get_image_info_single(cl_context context,
+                                      image_descriptor *imageInfo, MTdata d,
+                                      cl_mem_flags flags, size_t row_pitch,
+                                      size_t slice_pitch);
+
+
+int test_get_image_info_1D_buffer(cl_device_id device, cl_context context,
+                                  cl_image_format *format, cl_mem_flags flags)
+{
+    size_t maxWidth;
+    cl_ulong maxAllocSize, memSize;
+    image_descriptor imageInfo = { 0 };
+    RandomSeed seed(gRandomSeed);
+    size_t pixelSize;
+
+    memset(&imageInfo, 0x0, sizeof(image_descriptor));
+    imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+    imageInfo.format = format;
+    pixelSize = get_pixel_size(imageInfo.format);
+
+    int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE,
+                                sizeof(maxWidth), &maxWidth, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
+                             sizeof(maxAllocSize), &maxAllocSize, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize),
+                             &memSize, NULL);
+    test_error(error, "Unable to get max image 1D size from device");
+
+    if (memSize > (cl_ulong)SIZE_MAX)
+    {
+        memSize = (cl_ulong)SIZE_MAX;
+        maxAllocSize = (cl_ulong)SIZE_MAX;
+    }
+
+    if (gTestSmallImages)
+    {
+        for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++)
+        {
+            imageInfo.rowPitch = imageInfo.width * pixelSize;
+            if (gDebugTrace)
+                log_info("   at size %d (flags 0x%x pitch %d)\n",
+                         (int)imageInfo.width, (unsigned int)flags,
+                         (int)imageInfo.rowPitch);
+            if (test_get_image_info_single(context, &imageInfo, seed, flags, 0,
+                                           0))
+                return -1;
+        }
+    }
+    else if (gTestMaxImages)
+    {
+        // Try a specific set of maximum sizes
+        size_t numbeOfSizes;
+        size_t sizes[100][3];
+
+        get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1,
+                      maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER,
+                      imageInfo.format);
+
+        for (size_t idx = 0; idx < numbeOfSizes; idx++)
+        {
+            imageInfo.width = sizes[idx][0];
+            imageInfo.rowPitch = imageInfo.width * pixelSize;
+            log_info("Testing %d x 1\n", (int)sizes[idx][0]);
+            if (gDebugTrace)
+                log_info("   at max size %d (flags 0x%x pitch %d)\n",
+                         (int)imageInfo.width, (unsigned int)flags,
+                         (int)imageInfo.rowPitch);
+            if (test_get_image_info_single(context, &imageInfo, seed, flags, 0,
+                                           0))
+                return -1;
+        }
+    }
+    else
+    {
+        for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++)
+        {
+            cl_ulong size;
+            // Loop until we get a size that a) will fit in the max alloc size
+            // and b) that an allocation of that image, the result array, plus
+            // offset arrays, will fit in the global ram space
+            do
+            {
+                imageInfo.width =
+                    (size_t)random_log_in_range(16, (int)maxWidth / 32, seed);
+
+                imageInfo.rowPitch = imageInfo.width * pixelSize;
+                size_t extraWidth = (int)random_log_in_range(0, 64, seed);
+                imageInfo.rowPitch += extraWidth;
+
+                do
+                {
+                    extraWidth++;
+                    imageInfo.rowPitch += extraWidth;
+                } while ((imageInfo.rowPitch % pixelSize) != 0);
+
+                size = (cl_ulong)imageInfo.rowPitch * 4;
+            } while (size > maxAllocSize || (size * 3) > memSize);
+
+            if (gDebugTrace)
+                log_info("   at size %d (flags 0x%x pitch %d) out of %d\n",
+                         (int)imageInfo.width, (unsigned int)flags,
+                         (int)imageInfo.rowPitch, (int)maxWidth);
+            if (test_get_image_info_single(context, &imageInfo, seed, flags, 0,
+                                           0))
+                return -1;
+        }
+    }
+
+    return 0;
+}

test_conformance/images/clGetInfo/test_2D.cpp

@@ -19,6 +19,7 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo,
 {
     int error;
     clMemWrapper image;
+    clMemWrapper buffer;
     cl_image_desc imageDesc;
     void *host_ptr = NULL;
 
@@ -69,6 +70,24 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo,
             if ( gDebugTrace )
                 log_info( " - Creating 2D image array %d by %d by %d with flags=0x%lx row_pitch=%d slice_pitch=%d host_ptr=%p...\n", (int)imageInfo->width, (int)imageInfo->height, (int)imageInfo->arraySize, (unsigned long)flags, (int)row_pitch, (int)slice_pitch, host_ptr );
             break;
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+            if (gDebugTrace)
+                log_info(" - Creating 1D buffer image %d with flags=0x%lx "
+                         "row_pitch=%d slice_pitch=%d host_ptr=%p...\n",
+                         (int)imageInfo->width, (unsigned long)flags,
+                         (int)row_pitch, (int)slice_pitch, host_ptr);
+            int err;
+            buffer = clCreateBuffer(context, flags, imageInfo->rowPitch,
+                                    host_ptr, &err);
+            if (err != CL_SUCCESS)
+            {
+                log_error("ERROR: Unable to create buffer for 1D image buffer "
+                          "of size %d (%s)",
+                          (int)imageInfo->rowPitch, IGetErrorString(err));
+                return -1;
+            }
+            imageDesc.buffer = imageInfo->buffer = buffer;
+            break;
     }
 
     image = clCreateImage(context, flags, imageInfo->format, &imageDesc, host_ptr, &error);
@@ -92,6 +111,11 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo,
             case CL_MEM_OBJECT_IMAGE2D_ARRAY:
                 log_error( "ERROR: Unable to create 2D image array of size %d x %d x %d (%s)", (int)imageInfo->width, (int)imageInfo->height, (int)imageInfo->arraySize, IGetErrorString( error ) );
                 break;
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+                log_error(
+                    "ERROR: Unable to create 1D image buffer of size %d (%s)",
+                    (int)imageInfo->width, IGetErrorString(error));
+                break;
         }
         return -1;
     }
@@ -148,6 +172,7 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo,
   switch (imageInfo->type)
   {
     case CL_MEM_OBJECT_IMAGE1D:
+    case CL_MEM_OBJECT_IMAGE1D_BUFFER:
     case CL_MEM_OBJECT_IMAGE1D_ARRAY:
       required_height = 0;
       break;
@@ -175,8 +200,7 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo,
     case CL_MEM_OBJECT_IMAGE2D:
     case CL_MEM_OBJECT_IMAGE1D_ARRAY:
     case CL_MEM_OBJECT_IMAGE2D_ARRAY:
-      required_depth = 0;
+    case CL_MEM_OBJECT_IMAGE1D_BUFFER: required_depth = 0; break;
-      break;
     case CL_MEM_OBJECT_IMAGE3D:
       required_depth = imageInfo->depth;
       break;
@@ -198,8 +222,7 @@ int test_get_image_info_single( cl_context context, image_descriptor *imageInfo,
     case CL_MEM_OBJECT_IMAGE1D:
     case CL_MEM_OBJECT_IMAGE2D:
     case CL_MEM_OBJECT_IMAGE3D:
-      required_array_size = 0;
+    case CL_MEM_OBJECT_IMAGE1D_BUFFER: required_array_size = 0; break;
-      break;
     case CL_MEM_OBJECT_IMAGE1D_ARRAY:
     case CL_MEM_OBJECT_IMAGE2D_ARRAY:
       required_array_size = imageInfo->arraySize;

test_conformance/images/clGetInfo/test_loops.cpp

@@ -21,6 +21,10 @@ extern int test_get_image_info_2D( cl_device_id device, cl_context context, cl_i
 extern int test_get_image_info_3D( cl_device_id device, cl_context context, cl_image_format *format, cl_mem_flags flags );
 extern int test_get_image_info_1D_array( cl_device_id device, cl_context context, cl_image_format *format, cl_mem_flags flags );
 extern int test_get_image_info_2D_array( cl_device_id device, cl_context context, cl_image_format *format, cl_mem_flags flags );
+extern int test_get_image_info_1D_buffer(cl_device_id device,
+                                         cl_context context,
+                                         cl_image_format *format,
+                                         cl_mem_flags flags);
 
 int test_image_type( cl_device_id device, cl_context context, cl_mem_object_type image_type, cl_mem_flags flags )
 {
@@ -66,6 +70,10 @@ int test_image_type( cl_device_id device, cl_context context, cl_mem_object_type
           case CL_MEM_OBJECT_IMAGE2D_ARRAY:
             test_return = test_get_image_info_2D_array( device, context, &formatList[ i ], flags );
             break;
+          case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+              test_return = test_get_image_info_1D_buffer(
+                  device, context, &formatList[i], flags);
+              break;
         }
 
         if (test_return) {

test_conformance/images/clReadWriteImage/CMakeLists.txt

@@ -3,6 +3,7 @@ set(MODULE_NAME CL_READ_WRITE_IMAGES)
 set(${MODULE_NAME}_SOURCES
     main.cpp
     test_read_1D.cpp
+    test_read_1D_buffer.cpp
     test_read_1D_array.cpp
     test_read_2D.cpp
     test_read_2D_array.cpp

test_conformance/images/clReadWriteImage/main.cpp

@@ -51,13 +51,15 @@ int test_2Darray(cl_device_id device, cl_context context, cl_command_queue queue
 {
     return test_image_set( device, context, queue, CL_MEM_OBJECT_IMAGE2D_ARRAY );
 }
+int test_1Dbuffer(cl_device_id device, cl_context context,
+                  cl_command_queue queue, int num_elements)
+{
+    return test_image_set(device, context, queue, CL_MEM_OBJECT_IMAGE1D_BUFFER);
+}
 
 test_definition test_list[] = {
-    ADD_TEST( 1D ),
+    ADD_TEST(1D),      ADD_TEST(2D),      ADD_TEST(3D),
-    ADD_TEST( 2D ),
+    ADD_TEST(1Darray), ADD_TEST(2Darray), ADD_TEST(1Dbuffer),
-    ADD_TEST( 3D ),
-    ADD_TEST( 1Darray ),
-    ADD_TEST( 2Darray ),
 };
 
 const int test_num = ARRAY_SIZE( test_list );

test_conformance/images/clReadWriteImage/test_loops.cpp

@@ -33,6 +33,11 @@ extern int test_read_image_set_2D_array(cl_device_id device, cl_context context,
                                         cl_command_queue queue,
                                         cl_image_format *format,
                                         cl_mem_flags flags);
+extern int test_read_image_set_1D_buffer(cl_device_id device,
+                                         cl_context context,
+                                         cl_command_queue queue,
+                                         cl_image_format *format,
+                                         cl_mem_flags flags);
 
 int test_image_type( cl_device_id device, cl_context context, cl_command_queue queue, cl_mem_object_type imageType, cl_mem_flags flags )
 {
@@ -97,6 +102,10 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q
                 test_return = test_read_image_set_2D_array(
                     device, context, queue, &formatList[i], flags);
                 break;
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+                test_return = test_read_image_set_1D_buffer(
+                    device, context, queue, &formatList[i], flags);
+                break;
         }
 
         if (test_return)

test_conformance/images/clReadWriteImage/test_read_1D_buffer.cpp

@@ -0,0 +1,256 @@
+//
+// Copyright (c) 2023 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "../testBase.h"
+#include <CL/cl.h>
+
+int test_read_image_1D_buffer(cl_context context, cl_command_queue queue,
+                              image_descriptor *imageInfo, MTdata d,
+                              cl_mem_flags flags)
+{
+    int error;
+
+    clMemWrapper image;
+    clMemWrapper buffer;
+
+    // Generate some data to test against
+    BufferOwningPtr<char> imageValues;
+    generate_random_image_data(imageInfo, imageValues, d);
+
+    if (gDebugTrace)
+    {
+        log_info(" - Creating 1D image %d...\n", (int)imageInfo->width);
+        log_info(" with %llu mip levels\n",
+                 (unsigned long long)imageInfo->num_mip_levels);
+    }
+
+    buffer = clCreateBuffer(context, flags, imageInfo->rowPitch, NULL, &error);
+    if (error != CL_SUCCESS)
+    {
+        log_error("ERROR: Unable to create buffer for 1D image buffer of size "
+                  "%d (%s)",
+                  (int)imageInfo->rowPitch, IGetErrorString(error));
+    }
+
+    image = create_image_1d(context, flags, imageInfo->format, imageInfo->width,
+                            0, NULL, buffer, &error);
+    if (image == NULL)
+    {
+        log_error("ERROR: Unable to create 1D image buffer of size %d (%s)",
+                  (int)imageInfo->width, IGetErrorString(error));
+        return -1;
+    }
+
+    if (gDebugTrace) log_info(" - Writing image...\n");
+
+    size_t origin[3] = { 0, 0, 0 };
+    size_t region[3] = { imageInfo->width, 1, 1 };
+    size_t fullImageSize = imageInfo->rowPitch;
+
+    BufferOwningPtr<char> resultValues(malloc(fullImageSize));
+    size_t imgValMipLevelOffset = 0;
+
+    error = clEnqueueWriteImage(queue, image, CL_FALSE, origin, region,
+                                (gEnablePitch ? imageInfo->rowPitch : 0), 0,
+                                (char *)imageValues + imgValMipLevelOffset, 0,
+                                NULL, NULL);
+    if (error != CL_SUCCESS)
+    {
+        log_error("ERROR: Unable to write to 1D image of size %d \n",
+                  (int)imageInfo->width);
+        return -1;
+    }
+
+    // To verify, we just read the results right back and see whether they
+    // match the input
+    if (gDebugTrace)
+    {
+        log_info(" - Initing result array...\n");
+    }
+
+    // Note: we read back without any pitch, to verify pitch actually WORKED
+    size_t scanlineSize = imageInfo->width * get_pixel_size(imageInfo->format);
+    size_t imageSize = scanlineSize;
+    memset(resultValues, 0xff, imageSize);
+
+    if (gDebugTrace) log_info(" - Reading results...\n");
+
+    error = clEnqueueReadImage(queue, image, CL_TRUE, origin, region, 0, 0,
+                               resultValues, 0, NULL, NULL);
+    test_error(error, "Unable to read image values");
+
+    // Verify scanline by scanline, since the pitches are different
+    char *sourcePtr = (char *)imageValues + imgValMipLevelOffset;
+    char *destPtr = resultValues;
+
+    if (memcmp(sourcePtr, destPtr, scanlineSize) != 0)
+    {
+        log_error("ERROR: Scanline did not verify for image size %d pitch "
+                  "%d (extra %d bytes)\n",
+                  (int)imageInfo->width, (int)imageInfo->rowPitch,
+                  (int)imageInfo->rowPitch
+                      - (int)imageInfo->width
+                          * (int)get_pixel_size(imageInfo->format));
+
+        log_error("First few values: \n");
+        log_error(" Input: ");
+        uint32_t *s = (uint32_t *)sourcePtr;
+        uint32_t *d = (uint32_t *)destPtr;
+        for (int q = 0; q < 12; q++) log_error("%08x ", s[q]);
+        log_error("\nOutput: ");
+        for (int q = 0; q < 12; q++) log_error("%08x ", d[q]);
+        log_error("\n");
+
+        int outX;
+        int offset = (int)get_pixel_size(imageInfo->format)
+            * (int)(imageInfo->width - 16);
+        if (offset < 0) offset = 0;
+        int foundCount = debug_find_vector_in_image(
+            (char *)imageValues + imgValMipLevelOffset, imageInfo,
+            destPtr + offset, get_pixel_size(imageInfo->format), &outX, NULL,
+            NULL);
+        if (foundCount > 0)
+        {
+            int returnedOffset =
+                (offset / (int)get_pixel_size(imageInfo->format)) - outX;
+
+            if (memcmp(sourcePtr
+                           + returnedOffset * get_pixel_size(imageInfo->format),
+                       destPtr, get_pixel_size(imageInfo->format) * 8)
+                == 0)
+                log_error("       Values appear to be offsetted by %d\n",
+                          returnedOffset);
+            else
+                log_error(
+                    "       Calculated offset is %d but unable to verify\n",
+                    returnedOffset);
+        }
+        else
+        {
+            log_error("      Unable to determine offset\n");
+        }
+        return -1;
+    }
+    imgValMipLevelOffset +=
+        imageInfo->width * get_pixel_size(imageInfo->format);
+    return 0;
+}
+
+int test_read_image_set_1D_buffer(cl_device_id device, cl_context context,
+                                  cl_command_queue queue,
+                                  cl_image_format *format, cl_mem_flags flags)
+{
+    size_t maxWidth;
+    cl_ulong maxAllocSize, memSize;
+    image_descriptor imageInfo = { 0 };
+    RandomSeed seed(gRandomSeed);
+    size_t pixelSize;
+
+    if (gTestMipmaps)
+    {
+        // 1D image buffers don't support mipmaps
+        // https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html#cl_khr_mipmap_image
+        return 0;
+    }
+
+    imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+    imageInfo.format = format;
+    imageInfo.height = imageInfo.depth = imageInfo.slicePitch = 0;
+    pixelSize = get_pixel_size(imageInfo.format);
+
+    int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE,
+                                sizeof(maxWidth), &maxWidth, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
+                             sizeof(maxAllocSize), &maxAllocSize, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize),
+                             &memSize, NULL);
+    test_error(error, "Unable to get max image 2D size from device");
+
+    if (memSize > (cl_ulong)SIZE_MAX)
+    {
+        memSize = (cl_ulong)SIZE_MAX;
+        maxAllocSize = (cl_ulong)SIZE_MAX;
+    }
+
+    if (gTestSmallImages)
+    {
+        for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++)
+        {
+            imageInfo.rowPitch = imageInfo.width * pixelSize;
+
+            if (gDebugTrace) log_info("   at size %d\n", (int)imageInfo.width);
+
+            int ret = test_read_image_1D_buffer(context, queue, &imageInfo,
+                                                seed, flags);
+            if (ret) return -1;
+        }
+    }
+    else if (gTestMaxImages)
+    {
+        // Try a specific set of maximum sizes
+        size_t numbeOfSizes;
+        size_t sizes[100][3];
+
+        get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1,
+                      maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER,
+                      imageInfo.format);
+
+        for (size_t idx = 0; idx < numbeOfSizes; idx++)
+        {
+            imageInfo.width = sizes[idx][0];
+            imageInfo.rowPitch = imageInfo.width * pixelSize;
+
+            log_info("Testing %d\n", (int)imageInfo.width);
+            if (gDebugTrace) log_info("   at max size %d\n", (int)maxWidth);
+            if (test_read_image_1D_buffer(context, queue, &imageInfo, seed,
+                                          flags))
+                return -1;
+        }
+    }
+    else
+    {
+        for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++)
+        {
+            cl_ulong size;
+            // Loop until we get a size that a) will fit in the max alloc size
+            // and b) that an allocation of that image, the result array, plus
+            // offset arrays, will fit in the global ram space
+            do
+            {
+                imageInfo.width =
+                    (size_t)random_log_in_range(16, (int)maxWidth / 32, seed);
+
+                imageInfo.rowPitch = imageInfo.width * pixelSize;
+                if (gEnablePitch)
+                {
+                    size_t extraWidth = (int)random_log_in_range(0, 64, seed);
+                    imageInfo.rowPitch += extraWidth * pixelSize;
+                }
+
+                size = (size_t)imageInfo.rowPitch * 4;
+            } while (size > maxAllocSize || (size / 3) > memSize);
+
+            if (gDebugTrace)
+                log_info("   at size %d (row pitch %d) out of %d\n",
+                         (int)imageInfo.width, (int)imageInfo.rowPitch,
+                         (int)maxWidth);
+            int ret = test_read_image_1D_buffer(context, queue, &imageInfo,
+                                                seed, flags);
+            if (ret) return -1;
+        }
+    }
+
+    return 0;
+}

test_conformance/images/kernel_image_methods/CMakeLists.txt

@@ -3,6 +3,7 @@ set(MODULE_NAME KERNEL_IMAGE_METHODS)
 set(${MODULE_NAME}_SOURCES
     main.cpp
     test_1D.cpp
+    test_1D_buffer.cpp
     test_1D_array.cpp
     test_2D.cpp
     test_2D_array.cpp

test_conformance/images/kernel_image_methods/main.cpp

@@ -51,13 +51,15 @@ int test_2Darray(cl_device_id device, cl_context context, cl_command_queue queue
 {
     return test_image_set( device, context, queue, CL_MEM_OBJECT_IMAGE2D_ARRAY );
 }
+int test_1Dbuffer(cl_device_id device, cl_context context,
+                  cl_command_queue queue, int num_elements)
+{
+    return test_image_set(device, context, queue, CL_MEM_OBJECT_IMAGE1D_BUFFER);
+}
 
 test_definition test_list[] = {
-    ADD_TEST( 1D ),
+    ADD_TEST(1D),      ADD_TEST(2D),      ADD_TEST(3D),
-    ADD_TEST( 2D ),
+    ADD_TEST(1Darray), ADD_TEST(2Darray), ADD_TEST(1Dbuffer),
-    ADD_TEST( 3D ),
-    ADD_TEST( 1Darray ),
-    ADD_TEST( 2Darray ),
 };
 
 const int test_num = ARRAY_SIZE( test_list );

test_conformance/images/kernel_image_methods/test_1D_buffer.cpp

@@ -0,0 +1,282 @@
+//
+// Copyright (c) 2023 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "../testBase.h"
+#include <CL/cl.h>
+
+
+struct image_kernel_data
+{
+    cl_int width;
+    cl_int channelType;
+    cl_int channelOrder;
+    cl_int expectedChannelType;
+    cl_int expectedChannelOrder;
+};
+
+static const char *methodTest1DImageKernelPattern =
+    "typedef struct {\n"
+    "    int width;\n"
+    "    int channelType;\n"
+    "    int channelOrder;\n"
+    "    int expectedChannelType;\n"
+    "    int expectedChannelOrder;\n"
+    " } image_kernel_data;\n"
+    "__kernel void sample_kernel( %s image1d_buffer_t input, __global "
+    "image_kernel_data *outData )\n"
+    "{\n"
+    "   outData->width = get_image_width( input );\n"
+    "   outData->channelType = get_image_channel_data_type( input );\n"
+    "   outData->channelOrder = get_image_channel_order( input );\n"
+    "\n"
+    "   outData->expectedChannelType = %s;\n"
+    "   outData->expectedChannelOrder = %s;\n"
+    "}";
+
+static int test_get_1Dimage_buffer_info_single(cl_context context,
+                                               cl_command_queue queue,
+                                               image_descriptor *imageInfo,
+                                               MTdata d, cl_mem_flags flags)
+{
+    int error = 0;
+
+    clProgramWrapper program;
+    clKernelWrapper kernel;
+    clMemWrapper image, outDataBuffer, buffer;
+    char programSrc[10240];
+
+    image_kernel_data outKernelData;
+
+    // Generate some data to test against
+    BufferOwningPtr<char> imageValues;
+    generate_random_image_data(imageInfo, imageValues, d);
+
+    // Construct testing source
+    if (gDebugTrace)
+        log_info(" - Creating 1D image %d ...\n", (int)imageInfo->width);
+
+    buffer = clCreateBuffer(
+        context, flags, imageInfo->width * get_pixel_size(imageInfo->format),
+        NULL, &error);
+    if (error != CL_SUCCESS)
+    {
+        log_error("ERROR: Unable to create buffer for 1D image buffer of size "
+                  "%d (%s)",
+                  (int)imageInfo->rowPitch, IGetErrorString(error));
+    }
+
+    image = create_image_1d(context, flags, imageInfo->format, imageInfo->width,
+                            imageInfo->rowPitch, NULL, buffer, &error);
+    if (image == NULL)
+    {
+        log_error("ERROR: Unable to create 1D image of size %d (%s)",
+                  (int)imageInfo->width, IGetErrorString(error));
+        return -1;
+    }
+
+    char channelTypeConstantString[256] = { 0 };
+    char channelOrderConstantString[256] = { 0 };
+
+    const char *channelTypeName =
+        GetChannelTypeName(imageInfo->format->image_channel_data_type);
+    const char *channelOrderName =
+        GetChannelOrderName(imageInfo->format->image_channel_order);
+    const char *image_access_qualifier =
+        (flags == CL_MEM_READ_ONLY) ? "read_only" : "write_only";
+
+    if (channelTypeName && strlen(channelTypeName))
+        sprintf(channelTypeConstantString, "CLK_%s",
+                &channelTypeName[3]); // replace CL_* with CLK_*
+
+    if (channelOrderName && strlen(channelOrderName))
+        sprintf(channelOrderConstantString, "CLK_%s",
+                &channelOrderName[3]); // replace CL_* with CLK_*
+
+    // Create a program to run against
+    sprintf(programSrc, methodTest1DImageKernelPattern, image_access_qualifier,
+            channelTypeConstantString, channelOrderConstantString);
+
+    // log_info("-----------------------------------\n%s\n", programSrc);
+    error = clFinish(queue);
+    if (error) print_error(error, "clFinish failed.\n");
+    const char *ptr = programSrc;
+    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
+                                        "sample_kernel");
+    test_error(error, "Unable to create kernel to test against");
+
+    // Create an output buffer
+    outDataBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                   sizeof(outKernelData), NULL, &error);
+    test_error(error, "Unable to create output buffer");
+
+    // Set up arguments and run
+    error = clSetKernelArg(kernel, 0, sizeof(image), &image);
+    test_error(error, "Unable to set kernel argument");
+    error = clSetKernelArg(kernel, 1, sizeof(outDataBuffer), &outDataBuffer);
+    test_error(error, "Unable to set kernel argument");
+
+    size_t threads[1] = { 1 }, localThreads[1] = { 1 };
+
+    error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads,
+                                   localThreads, 0, NULL, NULL);
+    test_error(error, "Unable to run kernel");
+
+    error = clEnqueueReadBuffer(queue, outDataBuffer, CL_TRUE, 0,
+                                sizeof(outKernelData), &outKernelData, 0, NULL,
+                                NULL);
+    test_error(error, "Unable to read data buffer");
+
+
+    // Verify the results now
+    if (outKernelData.width != (cl_int)imageInfo->width)
+    {
+        log_error(
+            "ERROR: Returned width did not validate (expected %d, got %d)\n",
+            (int)imageInfo->width, (int)outKernelData.width);
+        error = -1;
+    }
+    if (outKernelData.channelType != (cl_int)outKernelData.expectedChannelType)
+    {
+        log_error(
+            "ERROR: Returned channel type did not validate (expected %s (%d), "
+            "got %d)\n",
+            GetChannelTypeName(imageInfo->format->image_channel_data_type),
+            (int)outKernelData.expectedChannelType,
+            (int)outKernelData.channelType);
+        error = -1;
+    }
+    if (outKernelData.channelOrder
+        != (cl_int)outKernelData.expectedChannelOrder)
+    {
+        log_error("ERROR: Returned channel order did not validate (expected %s "
+                  "(%d), got %d)\n",
+                  GetChannelOrderName(imageInfo->format->image_channel_order),
+                  (int)outKernelData.expectedChannelOrder,
+                  (int)outKernelData.channelOrder);
+        error = -1;
+    }
+
+    if (clFinish(queue) != CL_SUCCESS)
+    {
+        log_error("ERROR: CL Finished failed in %s \n", __FUNCTION__);
+        error = -1;
+    }
+
+    return error;
+}
+
+int test_get_image_info_1D_buffer(cl_device_id device, cl_context context,
+                                  cl_command_queue queue,
+                                  cl_image_format *format, cl_mem_flags flags)
+{
+    size_t maxWidth;
+    cl_ulong maxAllocSize, memSize;
+    image_descriptor imageInfo = { 0 };
+    RandomSeed seed(gRandomSeed);
+    size_t pixelSize;
+
+    imageInfo.type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+    imageInfo.format = format;
+    imageInfo.height = imageInfo.depth = imageInfo.slicePitch = 0;
+    pixelSize = get_pixel_size(imageInfo.format);
+
+    int error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_MAX_BUFFER_SIZE,
+                                sizeof(maxWidth), &maxWidth, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
+                             sizeof(maxAllocSize), &maxAllocSize, NULL);
+    error |= clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize),
+                             &memSize, NULL);
+    test_error(error, "Unable to get max image 1D size from device");
+
+    if (memSize > (cl_ulong)SIZE_MAX)
+    {
+        memSize = (cl_ulong)SIZE_MAX;
+        maxAllocSize = (cl_ulong)SIZE_MAX;
+    }
+
+    if (gTestSmallImages)
+    {
+        for (imageInfo.width = 1; imageInfo.width < 13; imageInfo.width++)
+        {
+            imageInfo.rowPitch = imageInfo.width * pixelSize;
+            if (gDebugTrace) log_info("   at size %d\n", (int)imageInfo.width);
+
+            int ret = test_get_1Dimage_buffer_info_single(
+                context, queue, &imageInfo, seed, flags);
+            if (ret) return -1;
+        }
+    }
+    else if (gTestMaxImages)
+    {
+        // Try a specific set of maximum sizes
+        size_t numbeOfSizes;
+        size_t sizes[100][3];
+
+        get_max_sizes(&numbeOfSizes, 100, sizes, maxWidth, 1, 1, 1,
+                      maxAllocSize, memSize, CL_MEM_OBJECT_IMAGE1D_BUFFER,
+                      imageInfo.format);
+
+        for (size_t idx = 0; idx < numbeOfSizes; idx++)
+        {
+            imageInfo.width = sizes[idx][0];
+            imageInfo.rowPitch = imageInfo.width * pixelSize;
+
+            log_info("Testing %d\n", (int)sizes[idx][0]);
+            if (gDebugTrace)
+                log_info("   at max size %d\n", (int)sizes[idx][0]);
+            if (test_get_1Dimage_buffer_info_single(context, queue, &imageInfo,
+                                                    seed, flags))
+                return -1;
+        }
+    }
+    else
+    {
+        for (int i = 0; i < NUM_IMAGE_ITERATIONS; i++)
+        {
+            cl_ulong size;
+            // Loop until we get a size that a) will fit in the max alloc size
+            // and b) that an allocation of that image, the result array, plus
+            // offset arrays, will fit in the global ram space
+            do
+            {
+                imageInfo.width =
+                    (size_t)random_log_in_range(16, (int)maxWidth / 32, seed);
+
+                imageInfo.rowPitch = imageInfo.width * pixelSize;
+                size_t extraWidth = (int)random_log_in_range(0, 64, seed);
+                imageInfo.rowPitch += extraWidth;
+
+                do
+                {
+                    extraWidth++;
+                    imageInfo.rowPitch += extraWidth;
+                } while ((imageInfo.rowPitch % pixelSize) != 0);
+
+                size = (cl_ulong)imageInfo.rowPitch * (cl_ulong)imageInfo.height
+                    * 4;
+            } while (size > maxAllocSize || (size * 3) > memSize);
+
+            if (gDebugTrace)
+                log_info("   at size %d (row pitch %d) out of %d\n",
+                         (int)imageInfo.width, (int)imageInfo.rowPitch,
+                         (int)maxWidth);
+            int ret = test_get_1Dimage_buffer_info_single(
+                context, queue, &imageInfo, seed, flags);
+            if (ret) return -1;
+        }
+    }
+
+    return 0;
+}

test_conformance/images/kernel_image_methods/test_loops.cpp

@@ -34,6 +34,11 @@ extern int test_get_image_info_2D_array(cl_device_id device, cl_context context,
                                         cl_command_queue queue,
                                         cl_image_format *format,
                                         cl_mem_flags flags);
+extern int test_get_image_info_1D_buffer(cl_device_id device,
+                                         cl_context context,
+                                         cl_command_queue queue,
+                                         cl_image_format *format,
+                                         cl_mem_flags flags);
 
 int test_image_type( cl_device_id device, cl_context context, cl_command_queue queue, cl_mem_object_type imageType, cl_mem_flags flags )
 {
@@ -84,6 +89,10 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q
                 test_return = test_get_image_info_2D_array(
                     device, context, queue, &formatList[i], flags);
                 break;
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+                test_return = test_get_image_info_1D_buffer(
+                    device, context, queue, &formatList[i], flags);
+                break;
         }
 
         if (test_return) {
@@ -109,6 +118,8 @@ int test_image_set( cl_device_id device, cl_context context, cl_command_queue qu
         switch (imageType) {
             case CL_MEM_OBJECT_IMAGE1D:
             test_missing_feature(version_check, "image_1D");
+            case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+                test_missing_feature(version_check, "image_1D_buffer");
             case CL_MEM_OBJECT_IMAGE1D_ARRAY:
             test_missing_feature(version_check, "image_1D_array");
             case CL_MEM_OBJECT_IMAGE2D_ARRAY:

test_conformance/images/testBase.h

@@ -53,6 +53,9 @@ enum MethodsToTest
     k2DTo2DArray = (1 << 8),
     k2DArrayTo3D = (1 << 9),
     k3DTo2DArray = (1 << 10),
+    k1DBuffer = (1 << 11),
+    k1DTo1DBuffer = (1 << 12),
+    k1DBufferTo1D = (1 << 13),
 };