Skip to content

Commit

Permalink
a few fixes (thread safety & cl_khr_command_buffer UB) (#1840)
Browse files Browse the repository at this point in the history 
some fixes we've been carrying in our CTS fork:

* fix UB in `command_buffer_event_sync.cpp`: enqueue of two commands in
two separate queues, with both using the same buffer argument, and no
synchronization between the commands.
* fix UB in `command_buffer_test_barrier.cpp`: missing synchronization
between zeroing command and command-buffer using two separate queues
* make `test_thread_dimensions.cpp` thread-safe to avoid spurious
errors.
  • Loading branch information
@franz
franz authored Oct 8, 2024
1 parent 617e7cb commit c40c8d5
Show file tree
Hide file tree
Showing 3 changed files with 94 additions and 79 deletions.
7 changes: 4 additions & 3 deletions test_conformance/extensions/cl_khr_command_buffer/command_buffer_event_sync.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -583,7 +583,7 @@ struct CommandBufferEventSync : public BasicCommandBufferTest

// process secondary queue
error =
clEnqueueFillBuffer(queue_sec, in_mem, &pattern_pri, sizeof(cl_int),
clEnqueueFillBuffer(queue_sec, in_mem, &pattern_sec, sizeof(cl_int),
0, data_size(), 0, nullptr, nullptr);
test_error(error, "clEnqueueFillBuffer failed");

Expand All @@ -593,8 +593,9 @@ struct CommandBufferEventSync : public BasicCommandBufferTest
"clEnqueueCommandBufferKHR in secondary queue failed");

// process primary queue
error = clEnqueueFillBuffer(queue, in_mem, &pattern_pri, sizeof(cl_int),
0, data_size(), 0, nullptr, event_ptrs[0]);
error =
clEnqueueFillBuffer(queue, in_mem, &pattern_pri, sizeof(cl_int), 0,
data_size(), 1, &test_event, event_ptrs[0]);
test_error(error, "clEnqueueFillBuffer failed");

cl_event wait_list[] = { test_event,
Expand Down
3 changes: 3 additions & 0 deletions test_conformance/extensions/cl_khr_command_buffer/command_buffer_test_barrier.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -93,6 +93,9 @@ struct BarrierWithWaitListKHR : public BasicCommandBufferTest
0, data_size(), 0, nullptr, nullptr);
test_error(error, "clEnqueueFillBufferKHR failed");

error = clFinish(queue);
test_error(error, "clFinish");

error = clEnqueueCommandBufferKHR(
0, nullptr, out_of_order_command_buffer, 0, nullptr, &event);
test_error(error, "clEnqueueCommandBufferKHR failed");
Expand Down
163 changes: 87 additions & 76 deletions test_conformance/thread_dimensions/test_thread_dimensions.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -245,10 +245,8 @@ static const char *thread_dimension_kernel_code_not_atomic_not_long =
"\n"
"}\n";

char dim_str[128];
char *print_dimensions(size_t x, size_t y, size_t z, cl_uint dim)
char *print_dimensions(char *dim_str, size_t x, size_t y, size_t z, cl_uint dim)
{
// Not thread safe...
if (dim == 1)
{
snprintf(dim_str, 128, "[%d]", (int)x);
Expand All @@ -268,10 +266,9 @@ char *print_dimensions(size_t x, size_t y, size_t z, cl_uint dim)
return dim_str;
}

char dim_str2[128];
char *print_dimensions2(size_t x, size_t y, size_t z, cl_uint dim)
char *print_dimensions2(char *dim_str2, size_t x, size_t y, size_t z,
cl_uint dim)
{
// Not thread safe...
if (dim == 1)
{
snprintf(dim_str2, 128, "[%d]", (int)x);
Expand Down Expand Up @@ -315,6 +312,9 @@ int run_test(cl_context context, cl_command_queue queue, cl_kernel kernel,
global_size[2] = final_z_size;
local_size[2] = local_z_size;

char dim_str[128];
char dim_str2[128];

cl_ulong start_valid_memory_address = 0;
cl_ulong end_valid_memory_address = memory_size;
cl_ulong last_memory_address = (cl_ulong)final_x_size
Expand Down Expand Up @@ -387,8 +387,9 @@ int run_test(cl_context context, cl_command_queue queue, cl_kernel kernel,
if (DEBUG)
log_info("\t\t\tExecuting kernel with global %s, NULL local, "
"%d dim, start address %llu, end address %llu.\n",
print_dimensions(global_size[0], global_size[1],
global_size[2], dimensions),
print_dimensions(dim_str, global_size[0],
global_size[1], global_size[2],
dimensions),
dimensions, start_valid_memory_address,
end_valid_memory_address);
}
Expand All @@ -398,14 +399,15 @@ int run_test(cl_context context, cl_command_queue queue, cl_kernel kernel,
clEnqueueNDRangeKernel(queue, kernel, dimensions, NULL,
global_size, local_size, 0, NULL, NULL);
if (DEBUG)
log_info("\t\t\tExecuting kernel with global %s, local %s, %d "
"dim, start address %llu, end address %llu.\n",
print_dimensions(global_size[0], global_size[1],
global_size[2], dimensions),
print_dimensions2(local_size[0], local_size[1],
local_size[2], dimensions),
dimensions, start_valid_memory_address,
end_valid_memory_address);
log_info(
"\t\t\tExecuting kernel with global %s, local %s, %d "
"dim, start address %llu, end address %llu.\n",
print_dimensions(dim_str, global_size[0], global_size[1],
global_size[2], dimensions),
print_dimensions2(dim_str2, local_size[0], local_size[1],
local_size[2], dimensions),
dimensions, start_valid_memory_address,
end_valid_memory_address);
}
if (err == CL_OUT_OF_RESOURCES)
{
Expand Down Expand Up @@ -482,18 +484,15 @@ int run_test(cl_context context, cl_command_queue queue, cl_kernel kernel,
}


static cl_uint max_x_size = 1, min_x_size = 1, max_y_size = 1, min_y_size = 1,
max_z_size = 1, min_z_size = 1;

static void set_min(cl_uint *x, cl_uint *y, cl_uint *z)
{
if (*x < min_x_size) *x = min_x_size;
if (*y < min_y_size) *y = min_y_size;
if (*z < min_z_size) *z = min_z_size;
if (*x > max_x_size) *x = max_x_size;
if (*y > max_y_size) *y = max_y_size;
if (*z > max_z_size) *z = max_z_size;
}
#define set_min(x, y, z) \
{ \
if (x < min_x_size) x = min_x_size; \
if (y < min_y_size) y = min_y_size; \
if (z < min_z_size) z = min_z_size; \
if (x > max_x_size) x = max_x_size; \
if (y > max_y_size) y = max_y_size; \
if (z > max_z_size) z = max_z_size; \
}


int test_thread_dimensions(cl_device_id device, cl_context context,
Expand All @@ -512,6 +511,12 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
int use_atomics = 1;
MTdata d;

char dim_str[128];
char dim_str2[128];

cl_uint max_x_size = 1, min_x_size = 1, max_y_size = 1, min_y_size = 1,
max_z_size = 1, min_z_size = 1;

if (getenv("CL_WIMPY_MODE") && !quick_test)
{
log_info("CL_WIMPY_MODE enabled, skipping test\n");
Expand Down Expand Up @@ -678,9 +683,6 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
cl_uint local_tests_per_size = 1 + dimensions + 2;
if (explicit_local == 0) local_tests_per_size = 1;

max_x_size = 1, min_x_size = 1, max_y_size = 1, min_y_size = 1,
max_z_size = 1, min_z_size = 1;

if (dimensions > 3)
{
log_error("Invalid dimensions: %d\n", dimensions);
Expand All @@ -700,7 +702,8 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
}

log_info("Testing with dimensions up to %s.\n",
print_dimensions(max_x_size, max_y_size, max_z_size, dimensions));
print_dimensions(dim_str, max_x_size, max_y_size, max_z_size,
dimensions));
if (bufferSize)
{
log_info("Testing with buffer size %d.\n", bufferSize);
Expand All @@ -723,7 +726,8 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
{

log_info("Base test size %s:\n",
print_dimensions(x_size, y_size, z_size, dimensions));
print_dimensions(dim_str, x_size, y_size, z_size,
dimensions));

cl_uint sub_test;
cl_uint final_x_size, final_y_size, final_z_size;
Expand All @@ -736,10 +740,10 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
if (sub_test == 0)
{
if (DEBUG)
log_info(
"\tTesting with base dimensions %s.\n",
print_dimensions(final_x_size, final_y_size,
final_z_size, dimensions));
log_info("\tTesting with base dimensions %s.\n",
print_dimensions(
dim_str, final_x_size, final_y_size,
final_z_size, dimensions));
}
else if (quick_test)
{
Expand All @@ -749,12 +753,13 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
final_x_size--;
final_y_size--;
final_z_size--;
set_min(&final_x_size, &final_y_size, &final_z_size);
set_min(final_x_size, final_y_size, final_z_size);
if (DEBUG)
log_info(
"\tTesting with all base dimensions - 1 %s.\n",
print_dimensions(final_x_size, final_y_size,
final_z_size, dimensions));
print_dimensions(dim_str, final_x_size,
final_y_size, final_z_size,
dimensions));
}
else if (sub_test <= dimensions * 2)
{
Expand All @@ -781,38 +786,41 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
dim_to_change);
return -1;
}
set_min(&final_x_size, &final_y_size, &final_z_size);
set_min(final_x_size, final_y_size, final_z_size);
if (DEBUG)
log_info(
"\tTesting with one base dimension +/- 1 %s.\n",
print_dimensions(final_x_size, final_y_size,
final_z_size, dimensions));
print_dimensions(dim_str, final_x_size,
final_y_size, final_z_size,
dimensions));
}
else if (sub_test == (dimensions * 2 + 1))
{
if (dimensions == 1) continue;
final_x_size--;
final_y_size--;
final_z_size--;
set_min(&final_x_size, &final_y_size, &final_z_size);
set_min(final_x_size, final_y_size, final_z_size);
if (DEBUG)
log_info(
"\tTesting with all base dimensions - 1 %s.\n",
print_dimensions(final_x_size, final_y_size,
final_z_size, dimensions));
print_dimensions(dim_str, final_x_size,
final_y_size, final_z_size,
dimensions));
}
else if (sub_test == (dimensions * 2 + 2))
{
if (dimensions == 1) continue;
final_x_size++;
final_y_size++;
final_z_size++;
set_min(&final_x_size, &final_y_size, &final_z_size);
set_min(final_x_size, final_y_size, final_z_size);
if (DEBUG)
log_info(
"\tTesting with all base dimensions + 1 %s.\n",
print_dimensions(final_x_size, final_y_size,
final_z_size, dimensions));
print_dimensions(dim_str, final_x_size,
final_y_size, final_z_size,
dimensions));
}
else
{
Expand All @@ -828,12 +836,12 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
(int)get_random_float(
0, (z_size / size_increase_per_iteration), d)
+ z_size / size_increase_per_iteration;
set_min(&final_x_size, &final_y_size, &final_z_size);
set_min(final_x_size, final_y_size, final_z_size);
if (DEBUG)
log_info(
"\tTesting with random dimensions %s.\n",
print_dimensions(final_x_size, final_y_size,
final_z_size, dimensions));
log_info("\tTesting with random dimensions %s.\n",
print_dimensions(
dim_str, final_x_size, final_y_size,
final_z_size, dimensions));
}

if (limit_size
Expand All @@ -842,8 +850,9 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
{
log_info("Skipping size %s as it exceeds max test "
"threads of %d.\n",
print_dimensions(final_x_size, final_y_size,
final_z_size, dimensions),
print_dimensions(dim_str, final_x_size,
final_y_size, final_z_size,
dimensions),
MAX_TOTAL_GLOBAL_THREADS_FOR_TEST);
continue;
}
Expand Down Expand Up @@ -993,26 +1002,27 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
}

if (DEBUG)
log_info(
"\t\tTesting local size %s.\n",
print_dimensions(local_x_size, local_y_size,
local_z_size, dimensions));
log_info("\t\tTesting local size %s.\n",
print_dimensions(
dim_str, local_x_size, local_y_size,
local_z_size, dimensions));

if (explicit_local == 0)
{
log_info(
"\tTesting global %s local [NULL]...\n",
print_dimensions(final_x_size, final_y_size,
final_z_size, dimensions));
log_info("\tTesting global %s local [NULL]...\n",
print_dimensions(
dim_str, final_x_size, final_y_size,
final_z_size, dimensions));
}
else
{
log_info(
"\tTesting global %s local %s...\n",
print_dimensions(final_x_size, final_y_size,
final_z_size, dimensions),
print_dimensions2(local_x_size, local_y_size,
local_z_size, dimensions));
log_info("\tTesting global %s local %s...\n",
print_dimensions(dim_str, final_x_size,
final_y_size,
final_z_size, dimensions),
print_dimensions2(
dim_str2, local_x_size, local_y_size,
local_z_size, dimensions));
}

// Avoid running with very small local sizes on very
Expand Down Expand Up @@ -1052,12 +1062,13 @@ int test_thread_dimensions(cl_device_id device, cl_context context,
// Otherwise, if we had errors add them up.
if (err)
{
log_error(
"Test global %s local %s failed.\n",
print_dimensions(final_x_size, final_y_size,
final_z_size, dimensions),
print_dimensions2(local_x_size, local_y_size,
local_z_size, dimensions));
log_error("Test global %s local %s failed.\n",
print_dimensions(
dim_str, final_x_size, final_y_size,
final_z_size, dimensions),
print_dimensions2(
dim_str2, local_x_size, local_y_size,
local_z_size, dimensions));
errors++;
clReleaseMemObject(array);
clReleaseKernel(kernel);
Expand Down

0 comments on commit c40c8d5

Please sign in to comment.