Skip to content
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.

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

Add tests for cl_ext_immutable_memory_objects #2286

Open
wants to merge 4 commits into
base: main
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from all commits
Commits
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
6 changes: 6 additions & 0 deletions test_common/harness/errorHelpers.h
Original file line number Diff line number Diff line change
Expand Up @@ -120,6 +120,12 @@ static int vlog_win32(const char *format, ...);
return retValue; \
} \
}
#define test_object_failure_ret(object, errCode, expectedErrCode, msg, \
retValue) \
{ \
test_assert_error_ret(object == nullptr, msg, retValue); \
test_failure_error_ret(errCode, expectedErrCode, msg, retValue); \
}
Comment on lines +123 to +128
Copy link
Contributor

Choose a reason for hiding this comment

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

I feel like we already have too many of these error macros, and this one seems like it will be used only in a few narrow cases (it's only used twice in this PR). Could we just writetest_assert_error_ret and test_failure_error_ret instead?

Note that "object" is really just a "pointer" in actual use, so if we decide to keep it, I'd prefer to give it a different name.

#define print_failure_error(errCode, expectedErrCode, msg) \
log_error("ERROR: %s! (Got %s, expected %s from %s:%d)\n", msg, \
IGetErrorString(errCode), IGetErrorString(expectedErrCode), \
Expand Down
8 changes: 8 additions & 0 deletions test_common/harness/kernelHelpers.h
Original file line number Diff line number Diff line change
Expand Up @@ -162,6 +162,14 @@ get_default_rounding_mode(const cl_device_id device,
return TEST_SKIPPED_ITSELF; \
}

#define PASSIVE_REQUIRE_IMMUTABLE_MEMORY_OBJECTS(device) \
if (!is_extension_available(device, "cl_ext_immutable_memory_objects")) \
{ \
log_info("\n\tNote: device does not support " \
"'cl_ext_immutable_memory_objects'. Skipping test...\n"); \
return TEST_SKIPPED_ITSELF; \
}
Comment on lines +165 to +171
Copy link
Contributor

Choose a reason for hiding this comment

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

Minor: do we need to add a new macro here? I think REQUIRE_EXTENSION should be sufficient and preferred.


/* Prints out the standard device header for all tests given the device to print
* for */
extern int printDeviceHeader(cl_device_id device);
Expand Down
1 change: 1 addition & 0 deletions test_conformance/api/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
198 changes: 198 additions & 0 deletions test_conformance/api/negative_enqueue_map_image.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,198 @@
//
// 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 <array>
#include <vector>
#include <memory>

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;

if (is_extension_available(device, "cl_ext_immutable_memory_objects"))
Copy link
Contributor

Choose a reason for hiding this comment

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

Minor: I think it would be preferable to skip this test if immutable memory objects are not supported. Can we use REQUIRE_EXTENSION to do this?

{
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<cl_uchar, decltype(&free)>;

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<cl_image_format> 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<cl_uchar *>(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<cl_uchar *>(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<cl_uchar *>(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<cl_uchar *>(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<cl_uchar *>(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);
test_object_failure_ret(
map, error, CL_INVALID_OPERATION,
"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_FAIL);

map = clEnqueueMapImage(queue, image, CL_TRUE,
CL_MAP_WRITE_INVALIDATE_REGION,
origin, region, nullptr, nullptr, 0,
nullptr, nullptr, &error);
test_object_failure_ret(
map, error, CL_INVALID_OPERATION,
"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_FAIL);
}
}
}
}

return TEST_PASS;
}
73 changes: 73 additions & 0 deletions test_conformance/api/test_kernels.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,7 @@
#include "testBase.h"
#include "harness/typeWrappers.h"
#include "harness/conversions.h"
#include <vector>

const char *sample_single_test_kernel[] = {
"__kernel void sample_test(__global float *src, __global int *dst)\n"
Expand Down Expand Up @@ -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"
Expand Down Expand Up @@ -631,3 +643,64 @@ REGISTER_TEST(kernel_global_constant)

return 0;
}

REGISTER_TEST(negative_set_immutable_memory_to_writeable_kernel_arg)
{
PASSIVE_REQUIRE_IMMUTABLE_MEMORY_OBJECTS(device);

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<cl_uchar> 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 read_only kernel argument",
TEST_FAIL);
Comment on lines +698 to +703
Copy link
Contributor

Choose a reason for hiding this comment

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

I think this should be write_only instead of read_only?

Suggested change
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 read_only kernel argument",
TEST_FAIL);
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;
}
Loading
Loading