diff --git a/test_conformance/api/test_kernels.cpp b/test_conformance/api/test_kernels.cpp index 30452caa..a8d02558 100644 --- a/test_conformance/api/test_kernels.cpp +++ b/test_conformance/api/test_kernels.cpp @@ -87,6 +87,19 @@ const char *sample_two_kernel_program[] = { "\n" "}\n" }; +const char *sample_read_only_image_test_kernel = R"( + __kernel void read_only_image_test(__write_only image2d_t img, __global uint4 *src) + { + write_imageui(img, (int2)(get_global_id(0), get_global_id(1)), src[0]); + } +)"; + +const char *sample_write_only_image_test_kernel = R"( + __kernel void write_only_image_test(__read_only image2d_t src, __global uint4 *dst) + { + dst[0]=read_imageui(src, (int2)(get_global_id(0), get_global_id(1))); + } +)"; REGISTER_TEST(get_kernel_info) { @@ -704,3 +717,58 @@ REGISTER_TEST(negative_set_immutable_memory_to_writeable_kernel_arg) return TEST_PASS; } + +REGISTER_TEST(negative_set_read_write_image_arg) +{ + cl_int error = CL_SUCCESS; + clProgramWrapper program; + clKernelWrapper write_image_kernel, read_image_kernel; + clMemWrapper write_only_image, read_only_image; + const char *test_kernels[2] = { sample_read_only_image_test_kernel, + sample_write_only_image_test_kernel }; + constexpr cl_image_format format = { CL_RGBA, CL_UNSIGNED_INT8 }; + const int 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"); + + read_image_kernel = clCreateKernel(program, "read_only_image_test", &error); + test_error(error, + "Unable to get read_only_image_test kernel for built program"); + + write_image_kernel = + clCreateKernel(program, "write_only_image_test", &error); + test_error(error, + "Unable to get write_only_image_test kernel for built program"); + + read_only_image = create_image_2d(context, CL_MEM_READ_ONLY, &format, + size_dim, size_dim, 0, nullptr, &error); + test_error(error, "create_image_2d failed"); + + write_only_image = create_image_2d(context, CL_MEM_WRITE_ONLY, &format, + size_dim, size_dim, 0, nullptr, &error); + test_error(error, "create_image_2d failed"); + + // Run the test + error = clSetKernelArg(read_image_kernel, 0, sizeof(read_only_image), + &read_only_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_READ_ONLY is " + "passed to a write_only kernel argument", + TEST_FAIL); + + error = clSetKernelArg(write_image_kernel, 0, sizeof(write_only_image), + &write_only_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_WRITE_ONLY is " + "passed to a read_only kernel argument", + TEST_FAIL); + + return TEST_PASS; +}