Skip to content

Added test to verify negative result of clSetKernelArg with CL_INVALID_MEM_OBJECT and MSAA image object argument #2472

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
shajder wants to merge 4 commits into
base: main
Choose a base branch
from
+62 −0
Open

Added test to verify negative result of clSetKernelArg with CL_INVALID_MEM_OBJECT and MSAA image object argument #2472

Changes from 2 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
cb7308b
Added test to verify negative result of clSetKernelArg
shajder Aug 5, 2025
2e9ae20
Merge branch 'main' into clSetKernelArg_negative_invalid_image_msaa
shajder Aug 7, 2025
82e829f
Merge branch 'main' into clSetKernelArg_negative_invalid_image_msaa
shajder Dec 16, 2025
18ca59b
Correction due to code review
shajder Dec 16, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
67 changes: 67 additions & 0 deletions test_conformance/api/test_kernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,8 @@
#include "testBase.h"
#include "harness/typeWrappers.h"
#include "harness/conversions.h"
#include "harness/featureHelpers.h"
#include "harness/stringHelpers.h"
#include <vector>

const char *sample_single_test_kernel[] = {
Expand Down Expand Up @@ -87,6 +89,15 @@ const char *sample_two_kernel_program[] = {
"\n"
"}\n" };

const char *sample_arg_image_msaa_test_kernel = R"(
#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable
__kernel void arg_image_test(%s src, __global float4 *dst)
{
int%s coord = (int%s)(get_global_id(0));
dst[0]=read_imagef(src, coord, 0);
}
)";

const char *sample_read_only_image_test_kernel = R"(
__kernel void read_only_image_test(__write_only image2d_t img, __global uint4 *src)
{
Expand Down Expand Up @@ -718,6 +729,62 @@ REGISTER_TEST(negative_set_immutable_memory_to_writeable_kernel_arg)
return TEST_PASS;
}

REGISTER_TEST(negative_set_kernel_arg_invalid_image_msaa)
{
PASSIVE_REQUIRE_IMAGE_SUPPORT(device)

if (!is_extension_available(device, "cl_khr_gl_msaa_sharing"))
{
log_info("Test not run because 'cl_khr_gl_msaa_sharing' extension is "
"not supported by the tested device\n");
return TEST_SKIPPED_ITSELF;
}
Copy link
Contributor

@bashbaug bashbaug Dec 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
if (!is_extension_available(device, "cl_khr_gl_msaa_sharing"))
{
log_info("Test not run because 'cl_khr_gl_msaa_sharing' extension is "
"not supported by the tested device\n");
return TEST_SKIPPED_ITSELF;
}
REQUIRE_EXTENSION("cl_khr_gl_msaa_sharing");

Copy link
Contributor Author

@shajder shajder Dec 16, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done


cl_int error = CL_SUCCESS;
clProgramWrapper program;
clKernelWrapper kernel;

std::vector<std::pair<std::string, std::string>> image_types = {
{ "image2d_msaa_t", "2" },
{ "image2d_array_msaa_t", "4" },
{ "image2d_msaa_depth_t", "4" },
{ "image2d_array_msaa_depth_t", "4" }
Copy link
Contributor

@bashbaug bashbaug Dec 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm not convinced that all of these cases should return an error.

I think image2d_msaa_t should work. Why isn't the 2D image created below trivially an MSAA image, just one with a single sample?

It's probably reasonable to expect that image2d_array_msaa_t and image2d_array_msaa_depth_t will fail, since the image that is created is a 2D image and not a 2D image array. Do we have similar testing for non-MSAA images?

I think that image2d_msaa_depth_t is an especially grey area, since the image is a 2D image, just not one with a CL_DEPTH image channel order. Do we expect implementations to check the channel order?

Note that the error description in clSetKernelArg was clarified recently in KhronosGroup/OpenCL-Docs#1493. The text in the currently published specs is:

If the argument is a multi-sample 2D depth image, the arg_value entry must be a pointer to a multisample depth image object.

This is somewhat ambiguous though - what is considered to be "a multisample depth image object"? Also, note that the spec does not require the image to be a 2D image, which seems like an oversight.

The updated text in the GitHub spec sources but not yet in a published specification is:

If the kernel argument is a 2D MSAA depth image, then the image memory object must be of image type CL_MEM_OBJECT_IMAGE2D and image channel order CL_DEPTH.

This is more precise, but maybe we got the CL_DEPTH image channel order part wrong.

Copy link
Contributor Author

@shajder shajder Dec 16, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think image2d_msaa_t should work. Why isn't the 2D image created below trivially an MSAA image, just one with a single sample?

While mathematically a single-sample image is a subset of MSAA, I believe image2d_msaa_t and image2d_t are treated as mutually exclusive types by the API. Specifically, we cannot pass an image2d_msaa_t to a read_imageX function that uses a sampler, nor can we use an image2d_t with the sampler-less read functions. It’s a similar situation to OpenGL, where you cannot use a sampler2D for a texture created with glTexImage2DMultisample (and vice-versa with sampler2DMS).

If the kernel argument is a 2D MSAA depth image, then the image memory object must be of image type CL_MEM_OBJECT_IMAGE2D and image channel order CL_DEPTH.

My interpretation relies on the "must" keyword here. It implies that if the kernel expects image2d_msaa_depth_t, the implementation is required to verify that the passed object actually has the CL_DEPTH channel order. If a user passes a non-depth image, that fails to meet the CL_DEPTH channel order requirement, which is why we expect clSetKernelArg to return an error. Does that make sense?

};

constexpr cl_image_format format = { CL_RGBA, CL_UNSIGNED_INT8 };
constexpr size_t size_dim = 32;
clMemWrapper image =
create_image_2d(context, CL_MEM_READ_ONLY, &format, size_dim, size_dim,
0, nullptr, &error);
test_error(error, "clCreateBuffer failed");

for (auto &el : image_types)
{
std::string program_source =
str_sprintf(std::string(sample_arg_image_msaa_test_kernel),
el.first.c_str(), el.second.c_str(), el.second.c_str());

const char *ptr = program_source.c_str();
cl_int error = create_single_kernel_helper(context, &program, &kernel,
1, &ptr, "arg_image_test");
test_error(error, "Unable to build test program");

kernel = clCreateKernel(program, "arg_image_test", &error);
test_error(error,
"Unable to get arg_image_test kernel for built program");

error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &image);
test_failure_error_ret(
error, CL_INVALID_MEM_OBJECT,
"clSetKernelArg is supposed to fail with CL_INVALID_MEM_OBJECT "
"when arg_value does not follow the rules described in spec for "
"msaa images",
TEST_FAIL);
}

return TEST_PASS;
}

REGISTER_TEST(negative_set_read_write_image_arg)
{
cl_int error = CL_SUCCESS;
Expand Down
Loading