-
Notifications
You must be signed in to change notification settings - Fork 208
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
base: main
Are you sure you want to change the base?
Changes from all commits
669e01b
5ae78e7
aba1567
e0cb8e5
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
Original file line number | Diff line number | Diff line change |
---|---|---|
|
@@ -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
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Minor: do we need to add a new macro here? I think |
||
|
||
/* Prints out the standard device header for all tests given the device to print | ||
* for */ | ||
extern int printDeviceHeader(cl_device_id device); | ||
|
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")) | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe 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 |
||
{ | ||
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; | ||
} |
Original file line number | Diff line number | Diff line change | ||||||||||||||||||||||||
---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
|
@@ -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" | ||||||||||||||||||||||||||
|
@@ -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) | ||||||||||||||||||||||||||
{ | ||||||||||||||||||||||||||
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
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. I think this should be
Suggested change
|
||||||||||||||||||||||||||
|
||||||||||||||||||||||||||
return TEST_PASS; | ||||||||||||||||||||||||||
} |
There was a problem hiding this comment.
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 write
test_assert_error_ret
andtest_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.