From 8701acfa90b5e95772782bbdfb1997a1623d43c1 Mon Sep 17 00:00:00 2001 From: Michael Rizkalla Date: Tue, 17 Jun 2025 18:19:11 +0100 Subject: [PATCH] Add tests for cl_ext_immutable_memory_objects (#2286) This change provides partial test coverage for KhronosGroup/OpenCL-Docs#1280 Adding CTS tests for: 1. clEnqueueMapBuffer, clEnqueueMapImage. 2. Command buffer negative tests. 3. clSetKernelArgs negative tests. The bulk of the tests is to make sure that the CL driver does not allow writing to a memory object that is created with `CL_MEM_IMMUTABLE_EXT` flag when used with the above APIs. --------- Signed-off-by: Michael Rizkalla --- test_conformance/api/CMakeLists.txt | 1 + .../api/negative_enqueue_map_image.cpp | 191 ++++++++++++ test_conformance/api/test_kernels.cpp | 73 +++++ test_conformance/basic/test_enqueue_map.cpp | 78 ++++- .../command_buffer_with_immutable_memory.h | 36 +++ .../negative_command_buffer_copy.cpp | 279 ++++++++++++++++++ .../negative_command_buffer_fill.cpp | 97 ++++++ 7 files changed, 745 insertions(+), 10 deletions(-) create mode 100644 test_conformance/api/negative_enqueue_map_image.cpp create mode 100644 test_conformance/extensions/cl_khr_command_buffer/command_buffer_with_immutable_memory.h diff --git a/test_conformance/api/CMakeLists.txt b/test_conformance/api/CMakeLists.txt index 3df9a81f..f2bfac35 100644 --- a/test_conformance/api/CMakeLists.txt +++ b/test_conformance/api/CMakeLists.txt @@ -4,6 +4,7 @@ set(${MODULE_NAME}_SOURCES main.cpp negative_platform.cpp negative_queue.cpp + negative_enqueue_map_image.cpp test_api_consistency.cpp test_bool.cpp test_retain.cpp diff --git a/test_conformance/api/negative_enqueue_map_image.cpp b/test_conformance/api/negative_enqueue_map_image.cpp new file mode 100644 index 00000000..95542163 --- /dev/null +++ b/test_conformance/api/negative_enqueue_map_image.cpp @@ -0,0 +1,191 @@ +// +// Copyright (c) 2024 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 "harness/clImageHelper.h" + +#include +#include +#include + +static constexpr cl_mem_object_type image_types[] = { + CL_MEM_OBJECT_IMAGE2D, CL_MEM_OBJECT_IMAGE3D, CL_MEM_OBJECT_IMAGE2D_ARRAY, + CL_MEM_OBJECT_IMAGE1D, CL_MEM_OBJECT_IMAGE1D_ARRAY +}; + +REGISTER_TEST(negative_enqueue_map_image) +{ + constexpr size_t image_dim = 32; + + REQUIRE_EXTENSION("cl_ext_immutable_memory_objects"); + + static constexpr cl_mem_flags mem_flags[]{ + CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR, + CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR, + CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR + }; + + static constexpr const char *mem_flags_string[]{ + "CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR", + "CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR", + "CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR | " + "CL_MEM_ALLOC_HOST_PTR" + }; + + static_assert(ARRAY_SIZE(mem_flags) == ARRAY_SIZE(mem_flags_string), + "mem_flags and mem_flags_string must be of the same size"); + + using CLUCharPtr = std::unique_ptr; + + for (size_t index = 0; index < ARRAY_SIZE(mem_flags); ++index) + { + cl_mem_flags mem_flag = mem_flags[index]; + + log_info("Testing memory flag: %s\n", mem_flags_string[index]); + for (cl_mem_object_type image_type : image_types) + { + // find supported image formats + cl_uint num_formats = 0; + + cl_int error = clGetSupportedImageFormats( + context, mem_flag, image_type, 0, nullptr, &num_formats); + test_error(error, + "clGetSupportedImageFormats failed to return supported " + "formats"); + + std::vector formats(num_formats); + error = clGetSupportedImageFormats(context, mem_flag, image_type, + num_formats, formats.data(), + nullptr); + test_error(error, + "clGetSupportedImageFormats failed to return supported " + "formats"); + + clMemWrapper image; + for (cl_image_format &fmt : formats) + { + log_info("Testing %s %s\n", + GetChannelOrderName(fmt.image_channel_order), + GetChannelTypeName(fmt.image_channel_data_type)); + + RandomSeed seed(gRandomSeed); + size_t origin[3] = { 0, 0, 0 }; + size_t region[3] = { image_dim, image_dim, image_dim }; + switch (image_type) + { + case CL_MEM_OBJECT_IMAGE1D: { + const size_t pixel_size = get_pixel_size(&fmt); + const size_t image_size = + image_dim * pixel_size * sizeof(cl_uchar); + CLUCharPtr imgptr{ static_cast( + create_random_data(kUChar, seed, + image_size)), + free }; + image = + create_image_1d(context, mem_flag, &fmt, image_dim, + 0, imgptr.get(), nullptr, &error); + region[1] = 1; + region[2] = 1; + break; + } + case CL_MEM_OBJECT_IMAGE2D: { + const size_t pixel_size = get_pixel_size(&fmt); + const size_t image_size = image_dim * image_dim + * pixel_size * sizeof(cl_uchar); + CLUCharPtr imgptr{ static_cast( + create_random_data(kUChar, seed, + image_size)), + free }; + image = + create_image_2d(context, mem_flag, &fmt, image_dim, + image_dim, 0, imgptr.get(), &error); + region[2] = 1; + break; + } + case CL_MEM_OBJECT_IMAGE3D: { + const size_t pixel_size = get_pixel_size(&fmt); + const size_t image_size = image_dim * image_dim + * image_dim * pixel_size * sizeof(cl_uchar); + CLUCharPtr imgptr{ static_cast( + create_random_data(kUChar, seed, + image_size)), + free }; + image = create_image_3d(context, mem_flag, &fmt, + image_dim, image_dim, image_dim, + 0, 0, imgptr.get(), &error); + break; + } + case CL_MEM_OBJECT_IMAGE1D_ARRAY: { + const size_t pixel_size = get_pixel_size(&fmt); + const size_t image_size = image_dim * image_dim + * pixel_size * sizeof(cl_uchar); + CLUCharPtr imgptr{ static_cast( + create_random_data(kUChar, seed, + image_size)), + free }; + image = create_image_1d_array(context, mem_flag, &fmt, + image_dim, image_dim, 0, + 0, imgptr.get(), &error); + region[1] = 1; + region[2] = 1; + break; + } + case CL_MEM_OBJECT_IMAGE2D_ARRAY: { + const size_t pixel_size = get_pixel_size(&fmt); + const size_t image_size = image_dim * image_dim + * image_dim * pixel_size * sizeof(cl_uchar); + CLUCharPtr imgptr{ static_cast( + create_random_data(kUChar, seed, + image_size)), + free }; + image = create_image_2d_array( + context, mem_flag, &fmt, image_dim, image_dim, + image_dim, 0, 0, imgptr.get(), &error); + region[2] = 1; + break; + } + } + test_error(error, "Failed to create image"); + + void *map = clEnqueueMapImage( + queue, image, CL_TRUE, CL_MAP_WRITE, origin, region, + nullptr, nullptr, 0, nullptr, nullptr, &error); + + constexpr const char *write_err_msg = + "clEnqueueMapImage should return CL_INVALID_OPERATION " + "when: \"image has been created with CL_MEM_IMMUTABLE_EXT " + "and CL_MAP_WRITE is set in map_flags\""; + test_assert_error(map == nullptr, write_err_msg); + test_failure_error_ret(error, CL_INVALID_OPERATION, + write_err_msg, TEST_FAIL); + + map = clEnqueueMapImage(queue, image, CL_TRUE, + CL_MAP_WRITE_INVALIDATE_REGION, origin, + region, nullptr, nullptr, 0, nullptr, + nullptr, &error); + + constexpr const char *write_invalidate_err_msg = + "clEnqueueMapImage should return CL_INVALID_OPERATION " + "when: \"image has been created with CL_MEM_IMMUTABLE_EXT " + "and CL_MAP_WRITE_INVALIDATE_REGION is set in map_flags\""; + test_assert_error(map == nullptr, write_invalidate_err_msg); + test_failure_error_ret(error, CL_INVALID_OPERATION, + write_invalidate_err_msg, TEST_FAIL); + } + } + } + + return TEST_PASS; +} diff --git a/test_conformance/api/test_kernels.cpp b/test_conformance/api/test_kernels.cpp index 2b5e9c56..30452caa 100644 --- a/test_conformance/api/test_kernels.cpp +++ b/test_conformance/api/test_kernels.cpp @@ -16,6 +16,7 @@ #include "testBase.h" #include "harness/typeWrappers.h" #include "harness/conversions.h" +#include const char *sample_single_test_kernel[] = { "__kernel void sample_test(__global float *src, __global int *dst)\n" @@ -49,6 +50,17 @@ const char *sample_const_test_kernel[] = { "\n" "}\n" }; +const char *sample_image_test_kernel[] = { + "__kernel void sample_image_test(__read_only image2d_t src, __write_only " + "image2d_t dst)\n" + "{\n" + " int2 coord = (int2)(get_global_id(0), get_global_id(1));\n" + " uint4 value = read_imageui(src, coord);\n" + " write_imageui(dst, coord, value);\n" + "\n" + "}\n" +}; + const char *sample_const_global_test_kernel[] = { "__constant int addFactor = 1024;\n" "__kernel void sample_test(__global int *src1, __global int *dst)\n" @@ -631,3 +643,64 @@ REGISTER_TEST(kernel_global_constant) return 0; } + +REGISTER_TEST(negative_set_immutable_memory_to_writeable_kernel_arg) +{ + REQUIRE_EXTENSION("cl_ext_immutable_memory_objects"); + + cl_int error = CL_SUCCESS; + clProgramWrapper program; + clKernelWrapper kernels[2]; + clMemWrapper image, buffer; + const char *test_kernels[2] = { sample_const_test_kernel[0], + sample_image_test_kernel[0] }; + constexpr cl_image_format formats = { CL_RGBA, CL_UNSIGNED_INT8 }; + constexpr size_t size_dim = 128; + + // Setup the test + error = create_single_kernel_helper(context, &program, nullptr, 2, + test_kernels, nullptr); + test_error(error, "Unable to build test program"); + + kernels[0] = clCreateKernel(program, "sample_test", &error); + test_error(error, "Unable to get sample_test kernel for built program"); + + kernels[1] = clCreateKernel(program, "sample_image_test", &error); + test_error(error, + "Unable to get sample_image_test kernel for built program"); + + std::vector mem_data(size_dim * size_dim); + buffer = clCreateBuffer(context, CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR, + sizeof(cl_int) * size_dim, mem_data.data(), &error); + test_error(error, "clCreateBuffer failed"); + + image = create_image_2d(context, CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR, + &formats, size_dim, size_dim, 0, mem_data.data(), + &error); + test_error(error, "create_image_2d failed"); + + // Run the test + error = clSetKernelArg(kernels[0], 0, sizeof(buffer), &buffer); + test_error(error, "clSetKernelArg failed"); + + error = clSetKernelArg(kernels[0], 2, sizeof(buffer), &buffer); + test_failure_error_ret(error, CL_INVALID_ARG_VALUE, + "clSetKernelArg is supposed to fail " + "with CL_INVALID_ARG_VALUE when a buffer is " + "created with CL_MEM_IMMUTABLE_EXT is " + "passed to a non-constant kernel argument", + TEST_FAIL); + + error = clSetKernelArg(kernels[1], 0, sizeof(image), &image); + test_error(error, "clSetKernelArg failed"); + + error = clSetKernelArg(kernels[1], 1, sizeof(image), &image); + test_failure_error_ret(error, CL_INVALID_ARG_VALUE, + "clSetKernelArg is supposed to fail " + "with CL_INVALID_ARG_VALUE when an image is " + "created with CL_MEM_IMMUTABLE_EXT is " + "passed to a write_only kernel argument", + TEST_FAIL); + + return TEST_PASS; +} diff --git a/test_conformance/basic/test_enqueue_map.cpp b/test_conformance/basic/test_enqueue_map.cpp index c13cebe4..90062bc2 100644 --- a/test_conformance/basic/test_enqueue_map.cpp +++ b/test_conformance/basic/test_enqueue_map.cpp @@ -29,6 +29,9 @@ const cl_mem_flags flag_set[] = { CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR, CL_MEM_USE_HOST_PTR, CL_MEM_COPY_HOST_PTR, + CL_MEM_USE_HOST_PTR | CL_MEM_IMMUTABLE_EXT, + CL_MEM_COPY_HOST_PTR | CL_MEM_IMMUTABLE_EXT, + CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR | CL_MEM_IMMUTABLE_EXT, 0 }; @@ -37,6 +40,9 @@ const char *flag_set_names[] = { "CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR", "CL_MEM_USE_HOST_PTR", "CL_MEM_COPY_HOST_PTR", + "CL_MEM_USE_HOST_PTR | CL_MEM_IMMUTABLE_EXT", + "CL_MEM_COPY_HOST_PTR | CL_MEM_IMMUTABLE_EXT", + "CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR | CL_MEM_IMMUTABLE_EXT", "0" }; // clang-format on @@ -44,7 +50,7 @@ const char *flag_set_names[] = { REGISTER_TEST(enqueue_map_buffer) { int error; - const size_t bufferSize = 256 * 256; + constexpr size_t bufferSize = 256 * 256; MTdataHolder d{ gRandomSeed }; BufferOwningPtr hostPtrData{ malloc(bufferSize) }; BufferOwningPtr referenceData{ malloc(bufferSize) }; @@ -57,18 +63,28 @@ REGISTER_TEST(enqueue_map_buffer) log_info("Testing with cl_mem_flags src: %s\n", flag_set_names[src_flag_id]); + if ((flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + && !is_extension_available(device, + "cl_ext_immutable_memory_objects")) + { + log_info("Device does not support CL_MEM_IMMUTABLE_EXT. " + "Skipping the memory flag.\n"); + continue; + } + generate_random_data(kChar, (unsigned int)bufferSize, d, hostPtrData); memcpy(referenceData, hostPtrData, bufferSize); void *hostPtr = nullptr; cl_mem_flags flags = flag_set[src_flag_id]; + const bool is_immutable_buffer = flags & CL_MEM_IMMUTABLE_EXT; bool hasHostPtr = (flags & CL_MEM_USE_HOST_PTR) || (flags & CL_MEM_COPY_HOST_PTR); if (hasHostPtr) hostPtr = hostPtrData; memObject = clCreateBuffer(context, flags, bufferSize, hostPtr, &error); test_error(error, "Unable to create testing buffer"); - if (!hasHostPtr) + if (!hasHostPtr && !is_immutable_buffer) { error = clEnqueueWriteBuffer(queue, memObject, CL_TRUE, 0, bufferSize, @@ -86,7 +102,18 @@ REGISTER_TEST(enqueue_map_buffer) cl_char *mappedRegion = (cl_char *)clEnqueueMapBuffer( queue, memObject, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, offset, length, 0, NULL, NULL, &error); - if (error != CL_SUCCESS) + + // Mapping should fail if the buffer is immutable + if (is_immutable_buffer) + { + test_failure_error_ret( + error, CL_INVALID_OPERATION, + "clEnqueueMapBuffer call was expected to fail " + "with CL_INVALID_OPERATION", + TEST_FAIL); + continue; + } + else if (error != CL_SUCCESS) { print_error(error, "clEnqueueMapBuffer call failed"); log_error("\tOffset: %d Length: %d\n", (int)offset, @@ -122,6 +149,11 @@ REGISTER_TEST(enqueue_map_buffer) finalData, 0, NULL, NULL); test_error(error, "Unable to read results"); + if (is_immutable_buffer && !hasHostPtr) + { + continue; + } + for (size_t q = 0; q < bufferSize; q++) { if (referenceData[q] != finalData[q]) @@ -140,9 +172,10 @@ REGISTER_TEST(enqueue_map_buffer) REGISTER_TEST(enqueue_map_image) { int error; - cl_image_format format = { CL_RGBA, CL_UNSIGNED_INT32 }; - const size_t imageSize = 256; - const size_t imageDataSize = imageSize * imageSize * 4 * sizeof(cl_uint); + constexpr cl_image_format format = { CL_RGBA, CL_UNSIGNED_INT32 }; + constexpr size_t imageSize = 256; + constexpr size_t imageDataSize = + imageSize * imageSize * 4 * sizeof(cl_uint); PASSIVE_REQUIRE_IMAGE_SUPPORT(device) @@ -158,20 +191,30 @@ REGISTER_TEST(enqueue_map_image) log_info("Testing with cl_mem_flags src: %s\n", flag_set_names[src_flag_id]); + if ((flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + && !is_extension_available(device, + "cl_ext_immutable_memory_objects")) + { + log_info("Device does not support CL_MEM_IMMUTABLE_EXT. " + "Skipping the memory flag.\n"); + continue; + } + generate_random_data(kUInt, (unsigned int)(imageSize * imageSize * 4), d, hostPtrData); memcpy(referenceData, hostPtrData, imageDataSize); cl_mem_flags flags = flag_set[src_flag_id]; + bool is_immutable_image = flags & CL_MEM_IMMUTABLE_EXT; bool hasHostPtr = (flags & CL_MEM_USE_HOST_PTR) || (flags & CL_MEM_COPY_HOST_PTR); void *hostPtr = nullptr; if (hasHostPtr) hostPtr = hostPtrData; - memObject = create_image_2d(context, CL_MEM_READ_WRITE | flags, &format, - imageSize, imageSize, 0, hostPtr, &error); + memObject = create_image_2d(context, flags, &format, imageSize, + imageSize, 0, hostPtr, &error); test_error(error, "Unable to create testing buffer"); - if (!hasHostPtr) + if (!hasHostPtr && !is_immutable_image) { size_t write_origin[3] = { 0, 0, 0 }, write_region[3] = { imageSize, imageSize, 1 }; @@ -198,7 +241,17 @@ REGISTER_TEST(enqueue_map_image) cl_uint *mappedRegion = (cl_uint *)clEnqueueMapImage( queue, memObject, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, offset, region, &rowPitch, NULL, 0, NULL, NULL, &error); - if (error != CL_SUCCESS) + + if (is_immutable_image) + { + test_failure_error_ret( + error, CL_INVALID_OPERATION, + "clEnqueueMapImage call was expected to fail " + "with CL_INVALID_OPERATION", + TEST_FAIL); + continue; + } + else if (error != CL_SUCCESS) { print_error(error, "clEnqueueMapImage call failed"); log_error("\tOffset: %d,%d Region: %d,%d\n", (int)offset[0], @@ -245,6 +298,11 @@ REGISTER_TEST(enqueue_map_image) finalRegion, 0, 0, finalData, 0, NULL, NULL); test_error(error, "Unable to read results"); + if (is_immutable_image && !hasHostPtr) + { + continue; + } + for (size_t q = 0; q < imageSize * imageSize * 4; q++) { if (referenceData[q] != finalData[q]) diff --git a/test_conformance/extensions/cl_khr_command_buffer/command_buffer_with_immutable_memory.h b/test_conformance/extensions/cl_khr_command_buffer/command_buffer_with_immutable_memory.h new file mode 100644 index 00000000..bf6b238f --- /dev/null +++ b/test_conformance/extensions/cl_khr_command_buffer/command_buffer_with_immutable_memory.h @@ -0,0 +1,36 @@ +// +// Copyright (c) 2024 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. + +#pragma once + +#include "basic_command_buffer.h" +#include + +template +struct CommandBufferWithImmutableMemoryObjectsTest : public TBase +{ + using TBase::TBase; + + static_assert(std::is_base_of::value, + "TBase must be BasicCommandBufferTest or a derived class"); + + bool Skip() override + { + bool is_immutable_memory_objects_supported = is_extension_available( + BasicCommandBufferTest::device, "cl_ext_immutable_memory_objects"); + + return !is_immutable_memory_objects_supported || TBase::Skip(); + } +}; diff --git a/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_copy.cpp b/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_copy.cpp index 5a93518c..7b7ec603 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_copy.cpp +++ b/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_copy.cpp @@ -14,6 +14,9 @@ // limitations under the License. // #include "basic_command_buffer.h" +#include "command_buffer_with_immutable_memory.h" +#include "imageHelpers.h" +#include //-------------------------------------------------------------------------- template @@ -577,6 +580,252 @@ struct CommandBufferCopyImageMutableHandleNotNull return CL_SUCCESS; } }; + +struct CommandBufferCopyToImmutableImage + : public CommandBufferWithImmutableMemoryObjectsTest< + CommandBufferCopyBaseTest> +{ + using CommandBufferWithImmutableMemoryObjectsTest:: + CommandBufferWithImmutableMemoryObjectsTest; + + cl_int Run() override + { + cl_int error = clCommandFillImageKHR( + command_buffer, nullptr, nullptr, src_image, fill_color_1, origin, + region, 0, nullptr, nullptr, nullptr); + + test_error(error, "clCommandFillImageKHR failed"); + + error = clCommandCopyImageKHR(command_buffer, nullptr, nullptr, + src_image, dst_image, origin, origin, + region, 0, 0, nullptr, nullptr); + + test_failure_error_ret(error, CL_INVALID_OPERATION, + "clCommandCopyImageKHR is supposed to fail " + "with CL_INVALID_OPERATION when dst_image is " + "created with CL_MEM_IMMUTABLE_EXT", + TEST_FAIL); + + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = BasicCommandBufferTest::SetUp(elements); + test_error(error, "BasicCommandBufferTest::SetUp failed"); + + src_image = create_image_2d(context, CL_MEM_READ_ONLY, &format, + img_width, img_height, 0, nullptr, &error); + test_error(error, "create_image_2d failed"); + + size_t pixel_size = get_pixel_size(&format); + size_t image_size = + pixel_size * sizeof(cl_uchar) * img_width * img_height; + + std::vector imgptr(image_size); + + dst_image = create_image_2d( + context, CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR, &format, + img_width, img_height, 0, imgptr.data(), &error); + test_error(error, "create_image_2d failed"); + + return CL_SUCCESS; + } + + clMemWrapper dst_image; + clMemWrapper src_image; + static constexpr cl_uint pattern_1 = 0x05; + const cl_uint fill_color_1[4] = { pattern_1, pattern_1, pattern_1, + pattern_1 }; +}; + +struct CommandBufferCopyToImmutableBuffer + : public CommandBufferWithImmutableMemoryObjectsTest< + CommandBufferCopyBaseTest> +{ + using CommandBufferWithImmutableMemoryObjectsTest:: + CommandBufferWithImmutableMemoryObjectsTest; + + cl_int Run() override + { + cl_int error = clCommandCopyBufferKHR(command_buffer, nullptr, nullptr, + in_mem, buffer, 0, 0, data_size, + 0, nullptr, nullptr, nullptr); + test_failure_error_ret(error, CL_INVALID_OPERATION, + "clCommandCopyBufferKHR is supposed to fail " + "with CL_INVALID_OPERATION when dst_buffer is " + "created with CL_MEM_IMMUTABLE_EXT", + TEST_FAIL); + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = BasicCommandBufferTest::SetUp(elements); + test_error(error, "BasicCommandBufferTest::SetUp failed"); + + in_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, data_size, nullptr, + &error); + test_error(error, "clCreateBuffer failed"); + + std::vector data(data_size); + + buffer = + clCreateBuffer(context, CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR, + data_size, data.data(), &error); + test_error(error, "clCreateBuffer failed"); + + return CL_SUCCESS; + } +}; + +struct CommandBufferCopyBufferToImmutableImage + : public CommandBufferWithImmutableMemoryObjectsTest< + CommandBufferCopyBaseTest> +{ + using CommandBufferWithImmutableMemoryObjectsTest:: + CommandBufferWithImmutableMemoryObjectsTest; + + cl_int Run() override + { + cl_int error = clCommandFillBufferKHR( + command_buffer, nullptr, nullptr, buffer, &pattern_1, + sizeof(pattern_1), 0, data_size, 0, nullptr, nullptr, nullptr); + + test_error(error, "clCommandFillBufferKHR failed"); + + error = clCommandCopyBufferToImageKHR(command_buffer, nullptr, nullptr, + buffer, image, 0, origin, region, + 0, 0, nullptr, nullptr); + + test_failure_error_ret( + error, CL_INVALID_OPERATION, + "clCommandCopyBufferToImageKHR is supposed to fail " + "with CL_INVALID_OPERATION when dst_image is " + "created with CL_MEM_IMMUTABLE_EXT", + TEST_FAIL); + + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = BasicCommandBufferTest::SetUp(elements); + test_error(error, "BasicCommandBufferTest::SetUp failed"); + + buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, data_size, nullptr, + &error); + test_error(error, "Unable to create buffer"); + + size_t pixel_size = get_pixel_size(&format); + size_t image_size = + pixel_size * sizeof(cl_uchar) * img_width * img_height; + + std::vector imgptr(image_size); + + image = create_image_2d( + context, CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR, &format, + img_width, img_height, 0, imgptr.data(), &error); + test_error(error, "create_image_2d failed"); + + return CL_SUCCESS; + } + + const uint8_t pattern_1 = 0x05; +}; + +struct CommandBufferCopyImageToImmutableBuffer + : public CommandBufferWithImmutableMemoryObjectsTest< + CommandBufferCopyBaseTest> +{ + using CommandBufferWithImmutableMemoryObjectsTest:: + CommandBufferWithImmutableMemoryObjectsTest; + + cl_int Run() override + { + cl_int error = clCommandFillImageKHR( + command_buffer, nullptr, nullptr, image, fill_color_1, origin, + region, 0, nullptr, nullptr, nullptr); + + test_error(error, "clCommandFillImageKHR failed"); + + error = clCommandCopyImageToBufferKHR(command_buffer, nullptr, nullptr, + image, buffer, origin, region, 0, + 0, nullptr, nullptr, nullptr); + + test_failure_error_ret( + error, CL_INVALID_OPERATION, + "clCommandCopyImageToBufferKHR is supposed to fail " + "with CL_INVALID_OPERATION when dst_buffer is " + "created with CL_MEM_IMMUTABLE_EXT", + TEST_FAIL); + + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = BasicCommandBufferTest::SetUp(elements); + test_error(error, "BasicCommandBufferTest::SetUp failed"); + + image = create_image_2d(context, CL_MEM_READ_WRITE, &format, img_width, + img_height, 0, NULL, &error); + test_error(error, "create_image_2d failed"); + + std::vector data(data_size); + + buffer = + clCreateBuffer(context, CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR, + data_size, data.data(), &error); + test_error(error, "Unable to create buffer"); + + return CL_SUCCESS; + } + + static constexpr cl_uint pattern_1 = 0x12; + const cl_uint fill_color_1[4] = { pattern_1, pattern_1, pattern_1, + pattern_1 }; +}; + +struct CommandBufferCopyToImmutableBufferRect + : public CommandBufferWithImmutableMemoryObjectsTest< + CommandBufferCopyBaseTest> +{ + using CommandBufferWithImmutableMemoryObjectsTest:: + CommandBufferWithImmutableMemoryObjectsTest; + + cl_int Run() override + { + cl_int error = clCommandCopyBufferRectKHR( + command_buffer, nullptr, nullptr, in_mem, buffer, origin, origin, + region, 0, 0, 0, 0, 0, nullptr, nullptr, nullptr); + test_failure_error_ret(error, CL_INVALID_OPERATION, + "clCommandCopyBufferRectKHR is supposed to fail " + "with CL_INVALID_OPERATION when dst_buffer is " + "created with CL_MEM_IMMUTABLE_EXT", + TEST_FAIL); + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = BasicCommandBufferTest::SetUp(elements); + test_error(error, "BasicCommandBufferTest::SetUp failed"); + + in_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, data_size, nullptr, + &error); + test_error(error, "clCreateBuffer failed"); + + std::vector data(data_size); + + buffer = + clCreateBuffer(context, CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR, + data_size, data.data(), &error); + test_error(error, "clCreateBuffer failed"); + + return CL_SUCCESS; + } +}; } REGISTER_TEST(negative_command_buffer_command_copy_buffer_queue_not_null) @@ -657,3 +906,33 @@ REGISTER_TEST( return MakeAndRunTest( device, context, queue, num_elements); } + +REGISTER_TEST(negative_copy_to_immutable_buffer) +{ + return MakeAndRunTest( + device, context, queue, num_elements); +} + +REGISTER_TEST(negative_copy_to_immutable_buffer_rect) +{ + return MakeAndRunTest( + device, context, queue, num_elements); +} + +REGISTER_TEST(negative_copy_image_to_immutable_buffer) +{ + return MakeAndRunTest( + device, context, queue, num_elements); +} + +REGISTER_TEST(negative_copy_to_immutable_image) +{ + return MakeAndRunTest( + device, context, queue, num_elements); +} + +REGISTER_TEST(negative_copy_buffer_to_immutable_image) +{ + return MakeAndRunTest( + device, context, queue, num_elements); +} diff --git a/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_fill.cpp b/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_fill.cpp index ae2067ba..ef14ade3 100644 --- a/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_fill.cpp +++ b/test_conformance/extensions/cl_khr_command_buffer/negative_command_buffer_fill.cpp @@ -14,6 +14,8 @@ // limitations under the License. // #include "basic_command_buffer.h" +#include "command_buffer_with_immutable_memory.h" +#include "imageHelpers.h" #include //-------------------------------------------------------------------------- @@ -456,6 +458,89 @@ struct CommandBufferCommandFillImageMutableHandleNotNull } }; +// CL_INVALID_OPERATION if destination buffer is immutable memory +struct CommandBufferCommandFillImmutableBuffer + : CommandBufferWithImmutableMemoryObjectsTest> +{ + using CommandBufferWithImmutableMemoryObjectsTest:: + CommandBufferWithImmutableMemoryObjectsTest; + + cl_int Run() override + { + cl_int error = clCommandFillBufferKHR( + command_buffer, nullptr, nullptr, buffer, &pattern_1, + sizeof(pattern_1), 0, buffer_size, 0, nullptr, nullptr, nullptr); + + test_failure_error_ret(error, CL_INVALID_OPERATION, + "clCommandFillBufferKHR is supposed to fail " + "with CL_INVALID_OPERATION when buffer is " + "created with CL_MEM_IMMUTABLE_EXT", + TEST_FAIL); + + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = BasicCommandBufferTest::SetUp(elements); + test_error(error, "BasicCommandBufferTest::SetUp failed"); + + std::vector data(buffer_size); + + buffer = + clCreateBuffer(context, CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR, + buffer_size, data.data(), &error); + test_error(error, "clCreateBuffer failed"); + + return CL_SUCCESS; + } + + clMemWrapper buffer; + const size_t buffer_size = 512; + const uint8_t pattern_1 = 0x0f; +}; + +struct CommandBufferCommandFillImmutableImage + : CommandBufferWithImmutableMemoryObjectsTest> +{ + using CommandBufferWithImmutableMemoryObjectsTest:: + CommandBufferWithImmutableMemoryObjectsTest; + + cl_int Run() override + { + cl_int error = clCommandFillImageKHR( + command_buffer, nullptr, nullptr, image, fill_color_1, origin, + region, 0, nullptr, nullptr, nullptr); + + test_failure_error_ret(error, CL_INVALID_OPERATION, + "clCommandFillImageKHR is supposed to fail " + "with CL_INVALID_OPERATION when image is " + "created with CL_MEM_IMMUTABLE_EXT", + TEST_FAIL); + + return CL_SUCCESS; + } + + cl_int SetUp(int elements) override + { + cl_int error = BasicCommandBufferTest::SetUp(elements); + test_error(error, "BasicCommandBufferTest::SetUp failed"); + + size_t pixel_size = get_pixel_size(&formats); + size_t image_size = pixel_size * sizeof(cl_uchar) * 512 * 512; + + std::vector imgptr(image_size); + + image = create_image_2d(context, + CL_MEM_IMMUTABLE_EXT | CL_MEM_COPY_HOST_PTR, + &formats, 512, 512, 0, imgptr.data(), &error); + test_error(error, "create_image_2d failed"); + + return CL_SUCCESS; + } + + clMemWrapper image; +}; } REGISTER_TEST(negative_command_buffer_command_fill_buffer_queue_not_null) @@ -537,3 +622,15 @@ REGISTER_TEST( return MakeAndRunTest( device, context, queue, num_elements); } + +REGISTER_TEST(negative_fill_immutable_image) +{ + return MakeAndRunTest( + device, context, queue, num_elements); +} + +REGISTER_TEST(negative_fill_immutable_buffer) +{ + return MakeAndRunTest( + device, context, queue, num_elements); +}