From b681d4f2c89e015aac0aa107a78e6abc153e966b Mon Sep 17 00:00:00 2001 From: Michael Rizkalla Date: Tue, 13 Jan 2026 17:46:02 +0000 Subject: [PATCH] Add `cl_ext_immutable_memory_objects` tests writing to and from buffer (#2432) This change extends the test coverage for https://github.com/KhronosGroup/OpenCL-Docs/pull/1280 The change tests: 1. Writing to immutable buffers. 2. Writing to buffer/image from immutable buffers. 3. Reading from immutable buffers. This change adds the following tests: 1. `test_negative_imagearraycopy` 2. `test_negative_imagearraycopy3d` 3. `test_immutable_bufferreadwriterect` 4. `test_immutable_arrayreadwrite` 5. `test_write_from_immutable_buffer_to_buffer` 6. `test_immutable_buffer_map_*` and extends the following tests: 1. `test_arrayimagecopy3d` 2. `test_arrayimagecopy` 3. `test_imagearraycopy3d` 4. `test_imagearraycopy` 5. `test_buffer_copy` 6. `test_buffer_partial_copy` Signed-off-by: Michael Rizkalla --- .../basic/test_arrayimagecopy.cpp | 28 +- .../basic/test_arrayreadwrite.cpp | 32 +- .../basic/test_bufferreadwriterect.cpp | 77 +++++ .../basic/test_imagearraycopy.cpp | 128 +++++++- test_conformance/buffers/main.cpp | 21 +- test_conformance/buffers/testBase.h | 2 +- test_conformance/buffers/test_buffer_copy.cpp | 155 ++++++++-- test_conformance/buffers/test_buffer_fill.cpp | 12 + test_conformance/buffers/test_buffer_map.cpp | 123 ++++++++ test_conformance/buffers/test_buffer_read.cpp | 18 ++ .../buffers/test_buffer_write.cpp | 284 ++++++++++++++++++ 11 files changed, 834 insertions(+), 46 deletions(-) diff --git a/test_conformance/basic/test_arrayimagecopy.cpp b/test_conformance/basic/test_arrayimagecopy.cpp index bb44abff..16b24390 100644 --- a/test_conformance/basic/test_arrayimagecopy.cpp +++ b/test_conformance/basic/test_arrayimagecopy.cpp @@ -188,9 +188,17 @@ REGISTER_TEST(arrayimagecopy) { PASSIVE_REQUIRE_IMAGE_SUPPORT(device) - return test_arrayimagecommon(device, context, queue, CL_MEM_READ_WRITE, - CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, - test_arrayimagecopy_single_format); + int error = test_arrayimagecommon(device, context, queue, CL_MEM_READ_WRITE, + CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, + test_arrayimagecopy_single_format); + if (is_extension_available(device, "cl_ext_immutable_memory_objects")) + { + error |= test_arrayimagecommon( + device, context, queue, CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR, + CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, + test_arrayimagecopy_single_format); + } + return error; } @@ -198,7 +206,15 @@ REGISTER_TEST(arrayimagecopy3d) { PASSIVE_REQUIRE_3D_IMAGE_SUPPORT(device) - return test_arrayimagecommon(device, context, queue, CL_MEM_READ_WRITE, - CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE3D, - test_arrayimagecopy_single_format); + int error = test_arrayimagecommon(device, context, queue, CL_MEM_READ_WRITE, + CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE3D, + test_arrayimagecopy_single_format); + if (is_extension_available(device, "cl_ext_immutable_memory_objects")) + { + error |= test_arrayimagecommon( + device, context, queue, CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR, + CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE3D, + test_arrayimagecopy_single_format); + } + return error; } diff --git a/test_conformance/basic/test_arrayreadwrite.cpp b/test_conformance/basic/test_arrayreadwrite.cpp index fe4bb995..4b0555fd 100644 --- a/test_conformance/basic/test_arrayreadwrite.cpp +++ b/test_conformance/basic/test_arrayreadwrite.cpp @@ -72,16 +72,36 @@ static int test_arrayreadwrite_impl(cl_device_id device, cl_context context, err = clEnqueueWriteBuffer( queue, buffer, CL_TRUE, offset * sizeof(cl_uint), sizeof(cl_uint) * cb, &reference_vals[offset], 0, nullptr, nullptr); - test_error(err, "clEnqueueWriteBuffer failed"); + if (flags & CL_MEM_IMMUTABLE_EXT) + { + test_failure_error_ret(err, CL_INVALID_OPERATION, + "clEnqueueWriteBuffer is expected to fail " + "with CL_INVALID_OPERATION when the buffer " + "is created with CL_MEM_IMMUTABLE_EXT", + TEST_FAIL); + } + else + { + test_error(err, "clEnqueueWriteBuffer failed"); + } err = clEnqueueReadBuffer( queue, buffer, CL_TRUE, offset * sizeof(cl_uint), cb * sizeof(cl_uint), &outptr[offset], 0, nullptr, nullptr); test_error(err, "clEnqueueReadBuffer failed"); + const cl_uint* expected_buffer_values = nullptr; + if (flags & CL_MEM_IMMUTABLE_EXT) + { + expected_buffer_values = inptr.data(); + } + else + { + expected_buffer_values = reference_vals.data(); + } for (int j = offset; j < offset + cb; j++) { - if (reference_vals[j] != outptr[j]) + if (expected_buffer_values[j] != outptr[j]) { log_error("ARRAY read, write test failed\n"); err = -1; @@ -105,3 +125,11 @@ REGISTER_TEST(arrayreadwrite) return test_arrayreadwrite_impl(device, context, queue, num_elements, CL_MEM_READ_WRITE); } + +REGISTER_TEST(immutable_arrayreadwrite) +{ + REQUIRE_EXTENSION("cl_ext_immutable_memory_objects"); + + return test_arrayreadwrite_impl(device, context, queue, num_elements, + CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR); +} diff --git a/test_conformance/basic/test_bufferreadwriterect.cpp b/test_conformance/basic/test_bufferreadwriterect.cpp index 883bff7c..03ba2706 100644 --- a/test_conformance/basic/test_bufferreadwriterect.cpp +++ b/test_conformance/basic/test_bufferreadwriterect.cpp @@ -14,6 +14,7 @@ // limitations under the License. // #include "harness/compat.h" +#include "errorHelpers.h" #include #include @@ -194,6 +195,43 @@ int copy_region(size_t src, size_t soffset[3], size_t sregion[3], size_t dst, si return 0; } +int immutable_copy_region(size_t src, size_t soffset[3], size_t sregion[3], + size_t dst, size_t doffset[3], size_t dregion[3]) +{ + + // Copy between cl buffers. + size_t src_slice_pitch = + (width[src] * height[src] != 1) ? width[src] * height[src] : 0; + size_t dst_slice_pitch = + (width[dst] * height[dst] != 1) ? width[dst] * height[dst] : 0; + size_t src_row_pitch = width[src]; + + cl_int err; + if (check_overlap_rect(soffset, doffset, sregion, src_row_pitch, + src_slice_pitch)) + { + log_info("Copy overlap reported, skipping copy buffer rect\n"); + return CL_SUCCESS; + } + else + { + err = clEnqueueCopyBufferRect(gQueue, buffer[src], buffer[dst], soffset, + doffset, sregion, /*dregion,*/ + width[src], src_slice_pitch, width[dst], + dst_slice_pitch, 0, nullptr, nullptr); + if (err != CL_INVALID_OPERATION) + { + log_error( + "clEnqueueCopyBufferRect should return " + "CL_INVALID_OPERATION but returned %s between %zu and %zu", + IGetErrorString(err), src, dst); + return TEST_FAIL; + } + } + + return TEST_PASS; +} + // This function compares the destination region in the buffer pointed // to by device, to the source region of the specified verify buffer. int verify_region(BufferType* device, size_t src, size_t soffset[3], size_t sregion[3], size_t dst, size_t doffset[3]) { @@ -337,6 +375,32 @@ int write_region(size_t src, size_t soffset[3], size_t sregion[3], size_t dst, s return 0; } +int immutable_write_region(size_t src, size_t soffset[3], size_t sregion[3], + size_t dst, size_t doffset[3], size_t dregion[3]) +{ + initialize_image(tmp_buffer, tmp_buffer_size, 1, 1, mt); + + size_t src_slice_pitch = + (width[src] * height[src] != 1) ? width[src] * height[src] : 0; + size_t dst_slice_pitch = + (width[dst] * height[dst] != 1) ? width[dst] * height[dst] : 0; + + cl_int error = clEnqueueWriteBufferRect( + gQueue, buffer[dst], CL_TRUE, doffset, soffset, dregion, width[dst], + dst_slice_pitch, width[src], src_slice_pitch, tmp_buffer, 0, nullptr, + nullptr); + + if (error != CL_INVALID_OPERATION) + { + log_error("clEnqueueWriteBufferRect should return CL_INVALID_OPERATION " + "but retured %s between %zu and %zu", + IGetErrorString(error), src, dst); + return TEST_FAIL; + } + + return TEST_PASS; +} + void CL_CALLBACK mem_obj_destructor_callback( cl_mem, void *data ) { free( data ); @@ -591,3 +655,16 @@ REGISTER_TEST(bufferreadwriterect) device, context, queue, num_elements, CL_MEM_USE_HOST_PTR | CL_MEM_READ_WRITE, test_functions); } + +REGISTER_TEST(immutable_bufferreadwriterect) +{ + REQUIRE_EXTENSION("cl_ext_immutable_memory_objects"); + + TestFunctions test_functions; + test_functions.copy = immutable_copy_region; + test_functions.read = read_verify_region; + test_functions.write = immutable_write_region; + return test_bufferreadwriterect_impl( + device, context, queue, num_elements, + CL_MEM_USE_HOST_PTR | CL_MEM_IMMUTABLE_EXT, test_functions); +} diff --git a/test_conformance/basic/test_imagearraycopy.cpp b/test_conformance/basic/test_imagearraycopy.cpp index a400c460..d0ce67a1 100644 --- a/test_conformance/basic/test_imagearraycopy.cpp +++ b/test_conformance/basic/test_imagearraycopy.cpp @@ -27,6 +27,82 @@ using test_function_t = int (*)(cl_device_id, cl_context, cl_command_queue, cl_mem_flags, cl_mem_flags, cl_mem_object_type, const cl_image_format *); +static int test_negative_imagearraycopy_single_format( + cl_device_id device, cl_context context, cl_command_queue queue, + cl_mem_flags image_flags, cl_mem_flags buffer_flags, + cl_mem_object_type image_type, const cl_image_format *format) +{ + std::unique_ptr bufptr{ nullptr, free }, + imgptr{ nullptr, free }; + clMemWrapper image; + clMemWrapper buffer; + const size_t img_width = 512; + const size_t img_height = 512; + const size_t img_depth = (image_type == CL_MEM_OBJECT_IMAGE3D) ? 32 : 1; + size_t elem_size; + size_t buffer_size; + cl_int err; + RandomSeed seed(gRandomSeed); + + const size_t origin[3] = { 0, 0, 0 }, + region[3] = { img_width, img_height, img_depth }; + + log_info("Testing %s %s\n", + GetChannelOrderName(format->image_channel_order), + GetChannelTypeName(format->image_channel_data_type)); + + elem_size = get_pixel_size(format); + buffer_size = + sizeof(cl_uchar) * elem_size * img_width * img_height * img_depth; + + if (image_flags & CL_MEM_USE_HOST_PTR || image_flags & CL_MEM_COPY_HOST_PTR) + { + imgptr.reset(static_cast( + create_random_data(kUChar, seed, buffer_size))); + } + + bufptr.reset( + static_cast(create_random_data(kUChar, seed, buffer_size))); + + if (CL_MEM_OBJECT_IMAGE2D == image_type) + { + image = create_image_2d(context, image_flags, format, img_width, + img_height, 0, imgptr.get(), &err); + } + else + { + image = + create_image_3d(context, image_flags, format, img_width, img_height, + img_depth, 0, 0, imgptr.get(), &err); + } + test_error(err, "create_image_xd failed"); + + if (!(image_flags & CL_MEM_USE_HOST_PTR + || image_flags & CL_MEM_COPY_HOST_PTR)) + { + imgptr.reset(static_cast( + create_random_data(kUChar, seed, buffer_size))); + + err = clEnqueueWriteImage(queue, image, CL_TRUE, origin, region, 0, 0, + imgptr.get(), 0, nullptr, nullptr); + test_error(err, "clEnqueueWriteImage failed"); + } + + buffer = + clCreateBuffer(context, buffer_flags, buffer_size, bufptr.get(), &err); + test_error(err, "clCreateBuffer failed"); + + err = clEnqueueCopyImageToBuffer(queue, image, buffer, origin, region, 0, 0, + nullptr, nullptr); + test_failure_error_ret( + err, CL_INVALID_OPERATION, + "clEnqueueCopyImageToBuffer should return CL_INVALID_OPERATION when: " + "\" dst_buffer is created with CL_MEM_IMMUTABLE_EXT flag\"", + TEST_FAIL); + + return TEST_PASS; +} + static int test_imagearraycopy_single_format( cl_device_id device, cl_context context, cl_command_queue queue, cl_mem_flags image_flags, cl_mem_flags buffer_flags, @@ -188,9 +264,18 @@ REGISTER_TEST(imagearraycopy) { PASSIVE_REQUIRE_IMAGE_SUPPORT(device) - return test_imagearraycommon(device, context, queue, CL_MEM_READ_WRITE, - CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, - test_imagearraycopy_single_format); + int error = test_imagearraycommon(device, context, queue, CL_MEM_READ_WRITE, + CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, + test_imagearraycopy_single_format); + + if (is_extension_available(device, "cl_ext_immutable_memory_objects")) + { + error |= test_imagearraycommon( + device, context, queue, CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR, + CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, + test_imagearraycopy_single_format); + } + return error; } @@ -198,7 +283,38 @@ REGISTER_TEST(imagearraycopy3d) { PASSIVE_REQUIRE_3D_IMAGE_SUPPORT(device) - return test_imagearraycommon(device, context, queue, CL_MEM_READ_ONLY, - CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE3D, - test_imagearraycopy_single_format); + int error = test_imagearraycommon(device, context, queue, CL_MEM_READ_ONLY, + CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE3D, + test_imagearraycopy_single_format); + + if (is_extension_available(device, "cl_ext_immutable_memory_objects")) + { + error |= test_imagearraycommon( + device, context, queue, CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR, + CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE3D, + test_imagearraycopy_single_format); + } + return error; +} + +REGISTER_TEST(negative_imagearraycopy) +{ + PASSIVE_REQUIRE_IMAGE_SUPPORT(device); + REQUIRE_EXTENSION("cl_ext_immutable_memory_objects"); + + return test_imagearraycommon(device, context, queue, CL_MEM_READ_WRITE, + CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR, + CL_MEM_OBJECT_IMAGE2D, + test_negative_imagearraycopy_single_format); +} + +REGISTER_TEST(negative_imagearraycopy3d) +{ + PASSIVE_REQUIRE_3D_IMAGE_SUPPORT(device); + REQUIRE_EXTENSION("cl_ext_immutable_memory_objects"); + + return test_imagearraycommon(device, context, queue, CL_MEM_READ_ONLY, + CL_MEM_IMMUTABLE_EXT | CL_MEM_USE_HOST_PTR, + CL_MEM_OBJECT_IMAGE3D, + test_negative_imagearraycopy_single_format); } diff --git a/test_conformance/buffers/main.cpp b/test_conformance/buffers/main.cpp index f2a8c2a3..496d3b7d 100644 --- a/test_conformance/buffers/main.cpp +++ b/test_conformance/buffers/main.cpp @@ -19,19 +19,24 @@ #include "testBase.h" -const cl_mem_flags flag_set[] = { - CL_MEM_ALLOC_HOST_PTR, - CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR, - CL_MEM_USE_HOST_PTR, - CL_MEM_COPY_HOST_PTR, - 0 -}; +const cl_mem_flags flag_set[] = { CL_MEM_ALLOC_HOST_PTR, + CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR, + CL_MEM_USE_HOST_PTR, + CL_MEM_COPY_HOST_PTR, + 0, + 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 }; const char* flag_set_names[] = { "CL_MEM_ALLOC_HOST_PTR", "CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR", "CL_MEM_USE_HOST_PTR", "CL_MEM_COPY_HOST_PTR", - "0" + "0", + "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", }; int main( int argc, const char *argv[] ) diff --git a/test_conformance/buffers/testBase.h b/test_conformance/buffers/testBase.h index 8c5bb0e4..4cd17155 100644 --- a/test_conformance/buffers/testBase.h +++ b/test_conformance/buffers/testBase.h @@ -25,6 +25,6 @@ extern const cl_mem_flags flag_set[]; extern const char* flag_set_names[]; -#define NUM_FLAGS 5 +#define NUM_FLAGS 8 #endif // _testBase_h diff --git a/test_conformance/buffers/test_buffer_copy.cpp b/test_conformance/buffers/test_buffer_copy.cpp index 81dbd5cf..cba2c626 100644 --- a/test_conformance/buffers/test_buffer_copy.cpp +++ b/test_conformance/buffers/test_buffer_copy.cpp @@ -39,7 +39,8 @@ static int verify_copy_buffer(int *inptr, int *outptr, int n) using alignedOwningPtr = std::unique_ptr; -static int test_copy( cl_command_queue queue, cl_context context, int num_elements, MTdata d ) +static int test_copy(cl_device_id device, cl_command_queue queue, + cl_context context, int num_elements, MTdata d) { clMemWrapper buffers[2]; cl_int err = CL_SUCCESS; @@ -76,10 +77,19 @@ static int test_copy( cl_command_queue queue, cl_context context, int num_elemen return TEST_FAIL; } + const bool has_immutable_memory_extension = + is_extension_available(device, "cl_ext_immutable_memory_objects"); + for (int src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) { for (int dst_flag_id = 0; dst_flag_id < NUM_FLAGS; dst_flag_id++) { + if (((flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + || (flag_set[dst_flag_id] & CL_MEM_IMMUTABLE_EXT)) + && !has_immutable_memory_extension) + { + continue; + } log_info("Testing with cl_mem_flags src: %s dst: %s\n", flag_set_names[src_flag_id], flag_set_names[dst_flag_id]); for (int i = 0; i < num_elements; i++) @@ -89,7 +99,6 @@ static int test_copy( cl_command_queue queue, cl_context context, int num_elemen reference_ptr[i] = (int)genrand_int32(d); } - if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)) buffers[0] = clCreateBuffer(context, flag_set[src_flag_id], sizeof(cl_int) * num_elements, @@ -116,7 +125,9 @@ static int test_copy( cl_command_queue queue, cl_context context, int num_elemen return TEST_FAIL; } - if (!(flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) && !(flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)) { + if (!(flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) + && !(flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)) + { err = clEnqueueWriteBuffer(queue, buffers[0], CL_TRUE, 0, sizeof(cl_int) * num_elements, reference_ptr.get(), 0, nullptr, @@ -130,11 +141,44 @@ static int test_copy( cl_command_queue queue, cl_context context, int num_elemen err = clEnqueueCopyBuffer(queue, buffers[0], buffers[1], 0, 0, sizeof(cl_int) * num_elements, 0, nullptr, nullptr); - if ( err != CL_SUCCESS ){ + if ((flag_set[dst_flag_id] & CL_MEM_IMMUTABLE_EXT)) + { + if (err != CL_INVALID_OPERATION) + { + test_failure_error_ret(err, CL_INVALID_OPERATION, + "clEnqueueCopyBuffer should return " + "CL_INVALID_OPERATION when: " + "\"dst_buffer is created with " + "CL_MEM_IMMUTABLE_EXT flag\"", + TEST_FAIL); + return TEST_FAIL; + } + } + else if (err != CL_SUCCESS) + { print_error(err, "clCopyArray failed\n"); return TEST_FAIL; } + err = clEnqueueReadBuffer(queue, buffers[0], true, 0, + sizeof(int) * num_elements, out_ptr.get(), + 0, nullptr, nullptr); + if (verify_copy_buffer(reference_ptr.get(), out_ptr.get(), + num_elements)) + { + log_error("test failed\n"); + return TEST_FAIL; + } + else + { + log_info("test passed\n"); + } + + // Reset out_ptr + for (int i = 0; i < num_elements; i++) + { + out_ptr[i] = (int)0xdeadbeef; // seed with incorrect data + } err = clEnqueueReadBuffer(queue, buffers[1], true, 0, sizeof(int) * num_elements, out_ptr.get(), 0, nullptr, nullptr); @@ -143,14 +187,20 @@ static int test_copy( cl_command_queue queue, cl_context context, int num_elemen return TEST_FAIL; } - if (verify_copy_buffer(reference_ptr.get(), out_ptr.get(), - num_elements)) + int *target_buffer = reference_ptr.get(); + if (flag_set[dst_flag_id] & CL_MEM_IMMUTABLE_EXT) { - log_error( " test failed\n" ); + target_buffer = invalid_ptr.get(); + } + + if (verify_copy_buffer(target_buffer, out_ptr.get(), num_elements)) + { + log_error("test failed\n"); return TEST_FAIL; } - else{ - log_info( " test passed\n" ); + else + { + log_info("test passed\n"); } } // dst flags } // src flags @@ -160,7 +210,10 @@ static int test_copy( cl_command_queue queue, cl_context context, int num_elemen } // end test_copy() -static int testPartialCopy( cl_command_queue queue, cl_context context, int num_elements, cl_uint srcStart, cl_uint dstStart, int size, MTdata d ) +static int testPartialCopy(cl_device_id device, cl_command_queue queue, + cl_context context, int num_elements, + cl_uint srcStart, cl_uint dstStart, int size, + MTdata d) { clMemWrapper buffers[2]; cl_int err = CL_SUCCESS; @@ -197,10 +250,19 @@ static int testPartialCopy( cl_command_queue queue, cl_context context, int num_ return TEST_FAIL; } + const bool has_immutable_memory_extension = + is_extension_available(device, "cl_ext_immutable_memory_objects"); + for (int src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) { for (int dst_flag_id = 0; dst_flag_id < NUM_FLAGS; dst_flag_id++) { + if (((flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + || (flag_set[dst_flag_id] & CL_MEM_IMMUTABLE_EXT)) + && !has_immutable_memory_extension) + { + continue; + } log_info("Testing with cl_mem_flags src: %s dst: %s\n", flag_set_names[src_flag_id], flag_set_names[dst_flag_id]); for (int i = 0; i < num_elements; i++) @@ -236,7 +298,9 @@ static int testPartialCopy( cl_command_queue queue, cl_context context, int num_ return TEST_FAIL; } - if (!(flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) && !(flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)){ + if (!(flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) + && !(flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)) + { err = clEnqueueWriteBuffer(queue, buffers[0], CL_TRUE, 0, sizeof(cl_int) * num_elements, reference_ptr.get(), 0, nullptr, @@ -251,27 +315,72 @@ static int testPartialCopy( cl_command_queue queue, cl_context context, int num_ queue, buffers[0], buffers[1], srcStart * sizeof(cl_int), dstStart * sizeof(cl_int), sizeof(cl_int) * size, 0, nullptr, nullptr); - if ( err != CL_SUCCESS){ - print_error(err, "clEnqueueCopyBuffer failed\n"); + if ((flag_set[dst_flag_id] & CL_MEM_IMMUTABLE_EXT)) + { + if (err != CL_INVALID_OPERATION) + { + test_failure_error_ret(err, CL_INVALID_OPERATION, + "clEnqueueCopyBuffer should return " + "CL_INVALID_OPERATION when: " + "\"dst_buffer is created with " + "CL_MEM_IMMUTABLE_EXT flag\"", + TEST_FAIL); + } + } + else if (err != CL_SUCCESS) + { + print_error(err, "clCopyArray failed\n"); return TEST_FAIL; } + err = clEnqueueReadBuffer(queue, buffers[0], true, 0, + sizeof(int) * num_elements, out_ptr.get(), + 0, nullptr, nullptr); + if (err != CL_SUCCESS) + { + print_error(err, "clEnqueueReadBuffer failed\n"); + return TEST_FAIL; + } + if (verify_copy_buffer(reference_ptr.get(), out_ptr.get(), + num_elements)) + { + log_error("test failed\n"); + return TEST_FAIL; + } + else + { + log_info("test passed\n"); + } + + // Reset out_ptr + for (int i = 0; i < num_elements; i++) + { + out_ptr[i] = (int)0xdeadbeef; // seed with incorrect data + } err = clEnqueueReadBuffer(queue, buffers[1], true, 0, sizeof(int) * num_elements, out_ptr.get(), 0, nullptr, nullptr); - if ( err != CL_SUCCESS){ + if (err != CL_SUCCESS) + { print_error(err, "clEnqueueReadBuffer failed\n"); return TEST_FAIL; } - if (verify_copy_buffer(reference_ptr.get() + srcStart, - out_ptr.get() + dstStart, size)) + cl_int *target_buffer = reference_ptr.get() + srcStart; + if (flag_set[dst_flag_id] & CL_MEM_IMMUTABLE_EXT) { - log_error("buffer_COPY test failed\n"); + target_buffer = invalid_ptr.get(); + } + + if (verify_copy_buffer(target_buffer, out_ptr.get() + dstStart, + size)) + { + log_error("test failed\n"); return TEST_FAIL; } - else{ - log_info("buffer_COPY test passed\n"); + else + { + log_info("test passed\n"); } } // dst mem flags } // src mem flags @@ -289,7 +398,7 @@ REGISTER_TEST(buffer_copy) // test the preset size log_info( "set size: %d: ", num_elements ); - if (test_copy(queue, context, num_elements, d) != TEST_PASS) + if (test_copy(device, queue, context, num_elements, d) != TEST_PASS) { err++; } @@ -298,7 +407,7 @@ REGISTER_TEST(buffer_copy) for ( i = 0; i < 8; i++ ){ size = (int)get_random_float(2.f,131072.f, d); log_info( "random size: %d: ", size ); - if (test_copy(queue, context, size, d) != TEST_PASS) + if (test_copy(device, queue, context, size, d) != TEST_PASS) { err++; } @@ -324,8 +433,8 @@ REGISTER_TEST(buffer_partial_copy) size = (int)get_random_float( 8.f, (float)(num_elements - srcStart), d ); dstStart = (cl_uint)get_random_float( 0.f, (float)(num_elements - size), d ); log_info( "random partial copy from %d to %d, size: %d: ", (int)srcStart, (int)dstStart, size ); - if (testPartialCopy(queue, context, num_elements, srcStart, dstStart, - size, d) + if (testPartialCopy(device, queue, context, num_elements, srcStart, + dstStart, size, d) != TEST_PASS) { err++; diff --git a/test_conformance/buffers/test_buffer_fill.cpp b/test_conformance/buffers/test_buffer_fill.cpp index 2e7a22de..d8fa7654 100644 --- a/test_conformance/buffers/test_buffer_fill.cpp +++ b/test_conformance/buffers/test_buffer_fill.cpp @@ -598,6 +598,12 @@ static int test_buffer_fill(cl_device_id deviceID, cl_context context, for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) { + // Skip immutable memory flags + if (flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } + clEventWrapper event[2]; clMemWrapper buffers[2]; if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)) @@ -721,6 +727,12 @@ REGISTER_TEST(buffer_fill_struct) for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) { + // Skip immutable memory flags + if (flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } + clProgramWrapper program; clKernelWrapper kernel; log_info("Testing with cl_mem_flags: %s\n", diff --git a/test_conformance/buffers/test_buffer_map.cpp b/test_conformance/buffers/test_buffer_map.cpp index 5cac90ab..3299902c 100644 --- a/test_conformance/buffers/test_buffer_map.cpp +++ b/test_conformance/buffers/test_buffer_map.cpp @@ -592,6 +592,12 @@ static int test_buffer_map_read( cl_device_id deviceID, cl_context context, cl_c for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) { + // Skip immutable memory flags + if (flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } + clMemWrapper buffer; outptr[i] = align_malloc( ptrSizes[i] * num_elements, min_alignment); if ( ! outptr[i] ){ @@ -671,6 +677,101 @@ static int test_buffer_map_read( cl_device_id deviceID, cl_context context, cl_c } // end test_buffer_map_read() +int test_immutable_buffer_map(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements, + size_t size, const char *type, int loops) +{ + REQUIRE_EXTENSION("cl_ext_immutable_memory_objects"); + + void *outptr[5]; + cl_int err; + int i; + size_t ptrSizes[5]; + int total_errors = 0; + MTdataHolder mtdata(gRandomSeed); + + size_t min_alignment = get_min_alignment(context); + + ptrSizes[0] = size; + ptrSizes[1] = ptrSizes[0] << 1; + ptrSizes[2] = ptrSizes[1] << 1; + ptrSizes[3] = ptrSizes[2] << 1; + ptrSizes[4] = ptrSizes[3] << 1; + + // embedded devices don't support long/ulong so skip over + if (!gHasLong && strstr(type, "long")) return TEST_SKIPPED_ITSELF; + + for (i = 0; i < loops; i++) + { + for (int src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) + { + // Testing writing from immutable flags + if (!(flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT)) + { + continue; + } + + clMemWrapper buffer; + outptr[i] = align_malloc(ptrSizes[i] * num_elements, min_alignment); + if (!outptr[i]) + { + log_error(" unable to allocate %d bytes of memory\n", + (int)ptrSizes[i] * num_elements); + return TEST_FAIL; + } + generate_random_data(kUChar, ptrSizes[i] * num_elements, mtdata, + outptr[i]); + + buffer = + clCreateBuffer(context, flag_set[src_flag_id], + ptrSizes[i] * num_elements, outptr[i], &err); + + if (nullptr == buffer || CL_SUCCESS != err) + { + print_error(err, "clCreateBuffer failed\n"); + align_free(outptr[i]); + return TEST_FAIL; + } + + void *mappedPtr = clEnqueueMapBuffer( + queue, buffer, CL_TRUE, CL_MAP_READ, 0, + ptrSizes[i] * num_elements, 0, nullptr, nullptr, &err); + if (err != CL_SUCCESS) + { + print_error(err, "clEnqueueMapBuffer failed"); + align_free(outptr[i]); + return TEST_FAIL; + } + + if (memcmp(mappedPtr, outptr[i], ptrSizes[i] * num_elements) != 0) + { + log_error(" %s%d test failed. cl_mem_flags src: %s\n", type, + 1 << i, flag_set_names[src_flag_id]); + total_errors++; + } + else + { + log_info(" %s%d test passed. cl_mem_flags src: %s\n", type, + 1 << i, flag_set_names[src_flag_id]); + } + + err = clEnqueueUnmapMemObject(queue, buffer, mappedPtr, 0, nullptr, + nullptr); + test_error(err, "clEnqueueUnmapMemObject failed"); + + // If we are using the outptr[i] as backing via USE_HOST_PTR we need + // to make sure we are done before freeing. + if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR)) + { + err = clFinish(queue); + test_error(err, "clFinish failed"); + } + align_free(outptr[i]); + } + } // cl_mem_flags + + return total_errors > 0 ? TEST_FAIL : TEST_PASS; +} #define DECLARE_LOCK_TEST(type, realType) \ REGISTER_TEST(buffer_map_read_##type) \ @@ -691,6 +792,28 @@ DECLARE_LOCK_TEST(char, cl_char) DECLARE_LOCK_TEST(uchar, cl_uchar) DECLARE_LOCK_TEST(float, cl_float) +#undef DECLARE_LOCK_TEST + +#define DECLARE_LOCK_TEST(type, realType) \ + REGISTER_TEST(immutable_buffer_map_##type) \ + { \ + return test_immutable_buffer_map(device, context, queue, num_elements, \ + sizeof(realType), #type, 5); \ + } + +DECLARE_LOCK_TEST(int, cl_int) +DECLARE_LOCK_TEST(uint, cl_uint) +DECLARE_LOCK_TEST(long, cl_long) +DECLARE_LOCK_TEST(ulong, cl_ulong) +DECLARE_LOCK_TEST(short, cl_short) +DECLARE_LOCK_TEST(ushort, cl_ushort) +DECLARE_LOCK_TEST(char, cl_char) +DECLARE_LOCK_TEST(uchar, cl_uchar) +DECLARE_LOCK_TEST(float, cl_float) + +#undef DECLARE_LOCK_TEST + + REGISTER_TEST(buffer_map_read_struct) { int (*foo)(void *,int); diff --git a/test_conformance/buffers/test_buffer_read.cpp b/test_conformance/buffers/test_buffer_read.cpp index dbf39ab4..814dee45 100644 --- a/test_conformance/buffers/test_buffer_read.cpp +++ b/test_conformance/buffers/test_buffer_read.cpp @@ -666,6 +666,12 @@ static int test_buffer_read(cl_device_id deviceID, cl_context context, for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) { + // Skip immutable memory flags + if (flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } + clMemWrapper buffer; outptr[i] = align_malloc( ptrSizes[i] * num_elements, min_alignment); if ( ! outptr[i] ){ @@ -809,6 +815,12 @@ static int test_buffer_read_async(cl_device_id deviceID, cl_context context, for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) { + // Skip immutable memory flags + if (flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } + clMemWrapper buffer; clEventWrapper event; outptr[i] = align_malloc(ptrSizes[i] * num_elements, min_alignment); @@ -946,6 +958,12 @@ static int test_buffer_read_array_barrier( for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) { + // Skip immutable memory flags + if (flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } + clMemWrapper buffer; clEventWrapper event; outptr[i] = align_malloc(ptrSizes[i] * num_elements, min_alignment); diff --git a/test_conformance/buffers/test_buffer_write.cpp b/test_conformance/buffers/test_buffer_write.cpp index 36dcc963..7c92dfd9 100644 --- a/test_conformance/buffers/test_buffer_write.cpp +++ b/test_conformance/buffers/test_buffer_write.cpp @@ -660,8 +660,18 @@ static int test_buffer_write(cl_device_id deviceID, cl_context context, for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) { + // Skip immutable memory flags + if (flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } for (dst_flag_id = 0; dst_flag_id < NUM_FLAGS; dst_flag_id++) { + // Skip immutable memory flags + if (flag_set[dst_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } clMemWrapper buffers[2]; if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)) @@ -834,8 +844,19 @@ REGISTER_TEST(buffer_write_struct) for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) { + // Skip immutable memory flags + if (flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } for (dst_flag_id = 0; dst_flag_id < NUM_FLAGS; dst_flag_id++) { + // Skip immutable memory flags + if (flag_set[dst_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } + clMemWrapper buffers[2]; inptr[i] = (TestStruct *)align_malloc(ptrSizes[i] * num_elements, min_alignment); @@ -996,7 +1017,17 @@ static int test_buffer_write_array_async( ptrSizes[4] = ptrSizes[3] << 1; for (src_flag_id=0; src_flag_id < NUM_FLAGS; src_flag_id++) { + // Skip immutable memory flags + if (flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } for (dst_flag_id=0; dst_flag_id < NUM_FLAGS; dst_flag_id++) { + // Skip immutable memory flags + if (flag_set[dst_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } log_info("Testing with cl_mem_flags src: %s dst: %s\n", flag_set_names[src_flag_id], flag_set_names[dst_flag_id]); loops = ( loops < 5 ? loops : 5 ); @@ -1974,3 +2005,256 @@ REGISTER_TEST(buffer_write_async_ulong) } // end test_buffer_ulong_write_array_async() + +int immutable_test_buffer_write(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements, + size_t size, const char *type, int loops, + void *inptr[5], const char *kernelCode[], + const char *kernelName[], + int (*fn)(void *, void *, int), MTdataHolder &d) +{ + REQUIRE_EXTENSION("cl_ext_immutable_memory_objects"); + + void *outptr[5]; + clProgramWrapper program[5]; + clKernelWrapper kernel[5]; + size_t ptrSizes[5]; + size_t global_work_size[3]; + cl_int err; + int i; + int src_flag_id, dst_flag_id; + int total_errors = 0; + + size_t min_alignment = get_min_alignment(context); + + global_work_size[0] = (size_t)num_elements; + + ptrSizes[0] = size; + ptrSizes[1] = ptrSizes[0] << 1; + ptrSizes[2] = ptrSizes[1] << 1; + ptrSizes[3] = ptrSizes[2] << 1; + ptrSizes[4] = ptrSizes[3] << 1; + + loops = (loops < 5 ? loops : 5); + for (i = 0; i < loops; i++) + { + err = create_single_kernel_helper(context, &program[i], &kernel[i], 1, + &kernelCode[i], kernelName[i]); + if (err) + { + log_error(" Error creating program for %s\n", type); + return TEST_FAIL; + } + + for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++) + { + // Testing writing from immutable flags + if (!(flag_set[src_flag_id] & CL_MEM_IMMUTABLE_EXT)) + { + continue; + } + for (dst_flag_id = 0; dst_flag_id < NUM_FLAGS; dst_flag_id++) + { + // Skip immutable memory flags + if (flag_set[dst_flag_id] & CL_MEM_IMMUTABLE_EXT) + { + continue; + } + + cl_mem_flags src_mem_flags = flag_set[src_flag_id]; + cl_mem_flags dst_mem_flags = flag_set[dst_flag_id]; + clMemWrapper buffers[2]; + + buffers[0] = + clCreateBuffer(context, src_mem_flags, + ptrSizes[i] * num_elements, inptr[i], &err); + + if (nullptr == buffers[0] || CL_SUCCESS != err) + { + align_free(outptr[i]); + print_error(err, " clCreateBuffer failed\n"); + return TEST_FAIL; + } + if (!strcmp(type, "half")) + { + outptr[i] = align_malloc(ptrSizes[i] * (num_elements * 2), + min_alignment); + buffers[1] = clCreateBuffer(context, dst_mem_flags, + ptrSizes[i] * 2 * num_elements, + outptr[i], &err); + } + else + { + outptr[i] = + align_malloc(ptrSizes[i] * num_elements, min_alignment); + if ((dst_mem_flags & CL_MEM_USE_HOST_PTR) + || (dst_mem_flags & CL_MEM_COPY_HOST_PTR)) + buffers[1] = clCreateBuffer(context, dst_mem_flags, + ptrSizes[i] * num_elements, + outptr[i], &err); + else + buffers[1] = clCreateBuffer(context, dst_mem_flags, + ptrSizes[i] * num_elements, + nullptr, &err); + } + if (err) + { + align_free(outptr[i]); + print_error(err, " clCreateBuffer failed\n"); + return TEST_FAIL; + } + + err = clSetKernelArg(kernel[i], 0, sizeof(cl_mem), + (void *)&buffers[0]); + err |= clSetKernelArg(kernel[i], 1, sizeof(cl_mem), + (void *)&buffers[1]); + if (err != CL_SUCCESS) + { + align_free(outptr[i]); + print_error(err, " clSetKernelArg failed"); + return TEST_FAIL; + } + + err = clEnqueueNDRangeKernel(queue, kernel[i], 1, nullptr, + global_work_size, nullptr, 0, + nullptr, nullptr); + if (err != CL_SUCCESS) + { + print_error(err, " clEnqueueNDRangeKernel failed"); + align_free(outptr[i]); + return TEST_FAIL; + } + + err = clEnqueueReadBuffer(queue, buffers[1], true, 0, + ptrSizes[i] * num_elements, outptr[i], + 0, nullptr, nullptr); + + if (err != CL_SUCCESS) + { + align_free(outptr[i]); + print_error(err, " clEnqueueReadBuffer failed"); + return TEST_FAIL; + } + + if (fn(inptr[i], outptr[i], + (int)(ptrSizes[i] * (size_t)num_elements / ptrSizes[0]))) + { + log_error( + " %s%d test failed. cl_mem_flags src: %s, dst: %s\n", + type, 1 << i, flag_set_names[src_flag_id], + flag_set_names[dst_flag_id]); + total_errors++; + } + else + { + log_info( + " %s%d test passed. cl_mem_flags src: %s, dst: %s\n", + type, 1 << i, flag_set_names[src_flag_id], + flag_set_names[dst_flag_id]); + } + // cleanup + align_free(outptr[i]); + } + } // dst cl_mem_flag + } // src cl_mem_flag + + return total_errors; + +} // end test_buffer_write() + +REGISTER_TEST(write_from_immutable_buffer_to_buffer) +{ + REQUIRE_EXTENSION("cl_ext_immutable_memory_objects"); + + static const char *immutable_buffer_write_int_kernel_code[] = { + R"( + __kernel void test_buffer_write_int(constant int *src, __global int *dst) + { + int tid = get_global_id(0); + + dst[tid] = src[tid]; + })", + + R"( + __kernel void test_buffer_write_int2(constant int2 *src, __global int2 *dst) + { + int tid = get_global_id(0); + + dst[tid] = src[tid]; + })", + + R"( + __kernel void test_buffer_write_int4(constant int4 *src, __global int4 *dst) + { + int tid = get_global_id(0); + + dst[tid] = src[tid]; + })", + + R"( + __kernel void test_buffer_write_int8(constant int8 *src, __global int8 *dst) + { + int tid = get_global_id(0); + + dst[tid] = src[tid]; + })", + + R"( + __kernel void test_buffer_write_int16(constant int16 *src, __global int16 *dst) + { + int tid = get_global_id(0); + + dst[tid] = src[tid]; + })" + }; + + static const char *immutable_int_kernel_name[] = { + "test_buffer_write_int", "test_buffer_write_int2", + "test_buffer_write_int4", "test_buffer_write_int8", + "test_buffer_write_int16" + }; + + if (gTestMap) + { + log_error("Immutable buffers cannot be mapped with CL_MEM_WRITE\n"); + return TEST_SKIPPED_ITSELF; + } + + int *inptr[5]; + size_t ptrSizes[5]; + int i, err; + cl_uint j; + int (*foo)(void *, void *, int); + MTdataHolder d(gRandomSeed); + + size_t min_alignment = get_min_alignment(context); + + foo = verify_write_int; + + ptrSizes[0] = sizeof(cl_int); + ptrSizes[1] = ptrSizes[0] << 1; + ptrSizes[2] = ptrSizes[1] << 1; + ptrSizes[3] = ptrSizes[2] << 1; + ptrSizes[4] = ptrSizes[3] << 1; + + for (i = 0; i < 5; i++) + { + inptr[i] = + (int *)align_malloc(ptrSizes[i] * num_elements, min_alignment); + + for (j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++) + inptr[i][j] = (int)genrand_int32(d); + } + + err = immutable_test_buffer_write(device, context, queue, num_elements, + sizeof(cl_int), "int", 5, (void **)inptr, + immutable_buffer_write_int_kernel_code, + immutable_int_kernel_name, foo, d); + + for (i = 0; i < 5; i++) + { + align_free((void *)inptr[i]); + } + + return err; +}