mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-22 07:09:03 +00:00
[NFC] Reformat code in events test (#1497)
Signed-off-by: Stuart Brady <stuart.brady@arm.com>
This commit is contained in:
Stuart Brady
@@ -1,6 +1,6 @@
|
||||
//
|
||||
// Copyright (c) 2017 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
|
||||
@@ -39,61 +39,79 @@ const char *write_kernels[] = {
|
||||
/*
|
||||
Tests event dependencies by running two kernels that use the same buffer.
|
||||
If two_queues is set they are run in separate queues.
|
||||
If test_enqueue_wait_for_events is set then clEnqueueWaitForEvent is called between them.
|
||||
If test_barrier is set then clEnqueueBarrier is called between them (only for single queue).
|
||||
If neither are set, nothing is done to prevent them from executing in the wrong order. This can be used for verification.
|
||||
If test_enqueue_wait_for_events is set then clEnqueueWaitForEvent is called
|
||||
between them. If test_barrier is set then clEnqueueBarrier is called between
|
||||
them (only for single queue). If neither are set, nothing is done to prevent
|
||||
them from executing in the wrong order. This can be used for verification.
|
||||
*/
|
||||
int test_event_enqueue_wait_for_events_run_test( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, int two_queues, int two_devices,
|
||||
int test_enqueue_wait_for_events, int test_barrier, int use_waitlist, int use_marker)
|
||||
int test_event_enqueue_wait_for_events_run_test(
|
||||
cl_device_id deviceID, cl_context context, cl_command_queue queue,
|
||||
int num_elements, int two_queues, int two_devices,
|
||||
int test_enqueue_wait_for_events, int test_barrier, int use_waitlist,
|
||||
int use_marker)
|
||||
{
|
||||
cl_int error = CL_SUCCESS;
|
||||
size_t threads[3] = {TEST_SIZE,0,0};
|
||||
size_t threads[3] = { TEST_SIZE, 0, 0 };
|
||||
int i, loop_count, event_count, expected_value, failed;
|
||||
int expected_if_only_queue[2];
|
||||
int max_count = TEST_SIZE;
|
||||
|
||||
cl_platform_id platform;
|
||||
cl_command_queue queues[2]; // Not a wrapper so we don't autorelease if they are the same
|
||||
clCommandQueueWrapper queueWrappers[2]; // If they are different, we use the wrapper so it will auto release
|
||||
cl_command_queue
|
||||
queues[2]; // Not a wrapper so we don't autorelease if they are the same
|
||||
clCommandQueueWrapper queueWrappers[2]; // If they are different, we use the
|
||||
// wrapper so it will auto release
|
||||
clContextWrapper context_to_use;
|
||||
clMemWrapper data;
|
||||
clProgramWrapper program;
|
||||
clKernelWrapper kernel1[TEST_COUNT], kernel2[TEST_COUNT];
|
||||
clEventWrapper event[TEST_COUNT*4+2]; // If we usemarkers we get 2 more events per iteration
|
||||
clEventWrapper event[TEST_COUNT * 4 + 2]; // If we usemarkers we get 2 more
|
||||
// events per iteration
|
||||
|
||||
if (test_enqueue_wait_for_events)
|
||||
log_info("\tTesting with clEnqueueBarrierWithWaitList as barrier function.\n");
|
||||
log_info("\tTesting with clEnqueueBarrierWithWaitList as barrier "
|
||||
"function.\n");
|
||||
if (test_barrier)
|
||||
log_info("\tTesting with clEnqueueBarrierWithWaitList as barrier function.\n");
|
||||
log_info("\tTesting with clEnqueueBarrierWithWaitList as barrier "
|
||||
"function.\n");
|
||||
if (use_waitlist)
|
||||
log_info("\tTesting with waitlist-based depenednecies between kernels.\n");
|
||||
log_info(
|
||||
"\tTesting with waitlist-based depenednecies between kernels.\n");
|
||||
if (use_marker)
|
||||
log_info("\tTesting with clEnqueueMarker as a barrier function.\n");
|
||||
if (test_barrier && (two_queues || two_devices)) {
|
||||
log_error("\tTest requested with clEnqueueBarrier across two queues. This is not a valid combination.\n");
|
||||
if (test_barrier && (two_queues || two_devices))
|
||||
{
|
||||
log_error("\tTest requested with clEnqueueBarrier across two queues. "
|
||||
"This is not a valid combination.\n");
|
||||
return -1;
|
||||
}
|
||||
|
||||
error = clGetPlatformIDs(1, &platform, NULL);
|
||||
test_error(error, "clGetPlatformIDs failed.");
|
||||
|
||||
// If we are to use two devices, then get them and create a context with both.
|
||||
// If we are to use two devices, then get them and create a context with
|
||||
// both.
|
||||
cl_device_id *two_device_ids;
|
||||
if (two_devices) {
|
||||
two_device_ids = (cl_device_id*)malloc(sizeof(cl_device_id)*2);
|
||||
if (two_devices)
|
||||
{
|
||||
two_device_ids = (cl_device_id *)malloc(sizeof(cl_device_id) * 2);
|
||||
cl_uint number_returned;
|
||||
error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 2, two_device_ids, &number_returned);
|
||||
test_error( error, "clGetDeviceIDs for CL_DEVICE_TYPE_ALL failed.");
|
||||
if (number_returned != 2) {
|
||||
error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 2, two_device_ids,
|
||||
&number_returned);
|
||||
test_error(error, "clGetDeviceIDs for CL_DEVICE_TYPE_ALL failed.");
|
||||
if (number_returned != 2)
|
||||
{
|
||||
log_info("Failed to obtain two devices. Test can not run.\n");
|
||||
free(two_device_ids);
|
||||
return 0;
|
||||
}
|
||||
|
||||
for (i=0; i<2; i++) {
|
||||
for (i = 0; i < 2; i++)
|
||||
{
|
||||
cl_device_type type;
|
||||
error = clGetDeviceInfo(two_device_ids[i], CL_DEVICE_TYPE, sizeof(cl_device_type), &type, NULL);
|
||||
test_error( error, "clGetDeviceInfo failed.");
|
||||
error = clGetDeviceInfo(two_device_ids[i], CL_DEVICE_TYPE,
|
||||
sizeof(cl_device_type), &type, NULL);
|
||||
test_error(error, "clGetDeviceInfo failed.");
|
||||
if (type & CL_DEVICE_TYPE_CPU)
|
||||
log_info("\tDevice %d is CL_DEVICE_TYPE_CPU.\n", i);
|
||||
if (type & CL_DEVICE_TYPE_GPU)
|
||||
@@ -104,12 +122,16 @@ int test_event_enqueue_wait_for_events_run_test( cl_device_id deviceID, cl_conte
|
||||
log_info("\tDevice %d is CL_DEVICE_TYPE_DEFAULT.\n", i);
|
||||
}
|
||||
|
||||
context_to_use = clCreateContext(NULL, 2, two_device_ids, notify_callback, NULL, &error);
|
||||
context_to_use = clCreateContext(NULL, 2, two_device_ids,
|
||||
notify_callback, NULL, &error);
|
||||
test_error(error, "clCreateContext failed for two devices.");
|
||||
|
||||
log_info("\tTesting with two devices.\n");
|
||||
} else {
|
||||
context_to_use = clCreateContext(NULL, 1, &deviceID, NULL, NULL, &error);
|
||||
}
|
||||
else
|
||||
{
|
||||
context_to_use =
|
||||
clCreateContext(NULL, 1, &deviceID, NULL, NULL, &error);
|
||||
test_error(error, "clCreateContext failed for one device.");
|
||||
|
||||
log_info("\tTesting with one device.\n");
|
||||
@@ -117,41 +139,55 @@ int test_event_enqueue_wait_for_events_run_test( cl_device_id deviceID, cl_conte
|
||||
|
||||
// If we are using two queues then create them
|
||||
cl_command_queue_properties props = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
|
||||
if (two_queues) {
|
||||
if (two_queues)
|
||||
{
|
||||
// Get a second queue
|
||||
if (two_devices)
|
||||
{
|
||||
if( !checkDeviceForQueueSupport( two_device_ids[ 0 ], props ) ||
|
||||
!checkDeviceForQueueSupport( two_device_ids[ 1 ], props ) )
|
||||
if (!checkDeviceForQueueSupport(two_device_ids[0], props)
|
||||
|| !checkDeviceForQueueSupport(two_device_ids[1], props))
|
||||
{
|
||||
log_info( "WARNING: One or more device for multi-device test does not support out-of-order exec mode; skipping test.\n" );
|
||||
log_info(
|
||||
"WARNING: One or more device for multi-device test does "
|
||||
"not support out-of-order exec mode; skipping test.\n");
|
||||
return -1942;
|
||||
}
|
||||
|
||||
queueWrappers[0] = clCreateCommandQueue(context_to_use, two_device_ids[0], props, &error);
|
||||
test_error(error, "clCreateCommandQueue for first queue on first device failed.");
|
||||
queueWrappers[1] = clCreateCommandQueue(context_to_use, two_device_ids[1], props, &error);
|
||||
test_error(error, "clCreateCommandQueue for second queue on second device failed.");
|
||||
|
||||
queueWrappers[0] = clCreateCommandQueue(
|
||||
context_to_use, two_device_ids[0], props, &error);
|
||||
test_error(
|
||||
error,
|
||||
"clCreateCommandQueue for first queue on first device failed.");
|
||||
queueWrappers[1] = clCreateCommandQueue(
|
||||
context_to_use, two_device_ids[1], props, &error);
|
||||
test_error(error,
|
||||
"clCreateCommandQueue for second queue on second device "
|
||||
"failed.");
|
||||
}
|
||||
else
|
||||
{
|
||||
// Single device has already been checked for out-of-order exec support
|
||||
queueWrappers[0] = clCreateCommandQueue(context_to_use, deviceID, props, &error);
|
||||
// Single device has already been checked for out-of-order exec
|
||||
// support
|
||||
queueWrappers[0] =
|
||||
clCreateCommandQueue(context_to_use, deviceID, props, &error);
|
||||
test_error(error, "clCreateCommandQueue for first queue failed.");
|
||||
queueWrappers[1] = clCreateCommandQueue(context_to_use, deviceID, props, &error);
|
||||
queueWrappers[1] =
|
||||
clCreateCommandQueue(context_to_use, deviceID, props, &error);
|
||||
test_error(error, "clCreateCommandQueue for second queue failed.");
|
||||
}
|
||||
// Ugly hack to make sure we only have the wrapper auto-release if they are different queues
|
||||
// Ugly hack to make sure we only have the wrapper auto-release if they
|
||||
// are different queues
|
||||
queues[0] = queueWrappers[0];
|
||||
queues[1] = queueWrappers[1];
|
||||
log_info("\tTesting with two queues.\n");
|
||||
}
|
||||
else
|
||||
{
|
||||
// (Note: single device has already been checked for out-of-order exec support)
|
||||
// Otherwise create one queue and have the second one be the same
|
||||
queueWrappers[0] = clCreateCommandQueue(context_to_use, deviceID, props, &error);
|
||||
// (Note: single device has already been checked for out-of-order exec
|
||||
// support) Otherwise create one queue and have the second one be the
|
||||
// same
|
||||
queueWrappers[0] =
|
||||
clCreateCommandQueue(context_to_use, deviceID, props, &error);
|
||||
test_error(error, "clCreateCommandQueue for first queue failed.");
|
||||
queues[0] = queueWrappers[0];
|
||||
queues[1] = (cl_command_queue)queues[0];
|
||||
@@ -160,236 +196,346 @@ int test_event_enqueue_wait_for_events_run_test( cl_device_id deviceID, cl_conte
|
||||
|
||||
|
||||
// Setup - create a buffer and the two kernels
|
||||
data = clCreateBuffer(context_to_use, CL_MEM_READ_WRITE, TEST_SIZE*sizeof(cl_int), NULL, &error);
|
||||
test_error( error, "clCreateBuffer failed");
|
||||
data = clCreateBuffer(context_to_use, CL_MEM_READ_WRITE,
|
||||
TEST_SIZE * sizeof(cl_int), NULL, &error);
|
||||
test_error(error, "clCreateBuffer failed");
|
||||
|
||||
|
||||
// Initialize the values to zero
|
||||
cl_int *values = (cl_int*)malloc(TEST_SIZE*sizeof(cl_int));
|
||||
for (i=0; i<(int)TEST_SIZE; i++)
|
||||
values[i] = 0;
|
||||
error = clEnqueueWriteBuffer(queues[0], data, CL_TRUE, 0, TEST_SIZE*sizeof(cl_int), values, 0, NULL, NULL);
|
||||
test_error( error, "clEnqueueWriteBuffer failed");
|
||||
cl_int *values = (cl_int *)malloc(TEST_SIZE * sizeof(cl_int));
|
||||
for (i = 0; i < (int)TEST_SIZE; i++) values[i] = 0;
|
||||
error =
|
||||
clEnqueueWriteBuffer(queues[0], data, CL_TRUE, 0,
|
||||
TEST_SIZE * sizeof(cl_int), values, 0, NULL, NULL);
|
||||
test_error(error, "clEnqueueWriteBuffer failed");
|
||||
expected_value = 0;
|
||||
|
||||
// Build the kernels
|
||||
if (create_single_kernel_helper( context_to_use, &program, &kernel1[0], 1, write_kernels, "write_up" ))
|
||||
if (create_single_kernel_helper(context_to_use, &program, &kernel1[0], 1,
|
||||
write_kernels, "write_up"))
|
||||
return -1;
|
||||
|
||||
error = clSetKernelArg(kernel1[0], 0, sizeof(data), &data);
|
||||
error |= clSetKernelArg(kernel1[0], 1, sizeof(max_count), &max_count);
|
||||
test_error( error, "clSetKernelArg 1 failed");
|
||||
test_error(error, "clSetKernelArg 1 failed");
|
||||
|
||||
for (i=1; i<TEST_COUNT; i++) {
|
||||
for (i = 1; i < TEST_COUNT; i++)
|
||||
{
|
||||
kernel1[i] = clCreateKernel(program, "write_up", &error);
|
||||
test_error( error, "clCreateKernel 1 failed");
|
||||
test_error(error, "clCreateKernel 1 failed");
|
||||
|
||||
error = clSetKernelArg(kernel1[i], 0, sizeof(data), &data);
|
||||
error |= clSetKernelArg(kernel1[i], 1, sizeof(max_count), &max_count);
|
||||
test_error( error, "clSetKernelArg 1 failed");
|
||||
test_error(error, "clSetKernelArg 1 failed");
|
||||
}
|
||||
|
||||
for (i=0; i<TEST_COUNT; i++) {
|
||||
for (i = 0; i < TEST_COUNT; i++)
|
||||
{
|
||||
kernel2[i] = clCreateKernel(program, "write_down", &error);
|
||||
test_error( error, "clCreateKernel 2 failed");
|
||||
test_error(error, "clCreateKernel 2 failed");
|
||||
|
||||
error = clSetKernelArg(kernel2[i], 0, sizeof(data), &data);
|
||||
error |= clSetKernelArg(kernel2[i], 1, sizeof(max_count), &max_count);
|
||||
test_error( error, "clSetKernelArg 2 failed");
|
||||
test_error(error, "clSetKernelArg 2 failed");
|
||||
}
|
||||
|
||||
// Execution - run the first kernel, then enqueue the wait on the events, then the second kernel
|
||||
// If clEnqueueBarrierWithWaitList works, the buffer will be filled with 1s, then multiplied by 4s,
|
||||
// then incremented to 5s, repeatedly. Otherwise the values may be 2s (if the first one doesn't work) or 8s
|
||||
// (if the second one doesn't work).
|
||||
// Execution - run the first kernel, then enqueue the wait on the events,
|
||||
// then the second kernel If clEnqueueBarrierWithWaitList works, the buffer
|
||||
// will be filled with 1s, then multiplied by 4s, then incremented to 5s,
|
||||
// repeatedly. Otherwise the values may be 2s (if the first one doesn't
|
||||
// work) or 8s (if the second one doesn't work).
|
||||
if (RANDOMIZE)
|
||||
log_info("Queues chosen randomly for each kernel execution.\n");
|
||||
else
|
||||
log_info("Queues chosen alternatily for each kernel execution.\n");
|
||||
|
||||
event_count = 0;
|
||||
for (i=0; i<(int)TEST_SIZE; i++)
|
||||
values[i] = 1;
|
||||
error = clEnqueueWriteBuffer(queues[0], data, CL_FALSE, 0, TEST_SIZE*sizeof(cl_int), values, 0, NULL, &event[event_count]);
|
||||
test_error( error, "clEnqueueWriteBuffer 2 failed");
|
||||
for (i = 0; i < (int)TEST_SIZE; i++) values[i] = 1;
|
||||
error = clEnqueueWriteBuffer(queues[0], data, CL_FALSE, 0,
|
||||
TEST_SIZE * sizeof(cl_int), values, 0, NULL,
|
||||
&event[event_count]);
|
||||
test_error(error, "clEnqueueWriteBuffer 2 failed");
|
||||
expected_value = 1;
|
||||
expected_if_only_queue[0] = 1;
|
||||
expected_if_only_queue[1] = 1;
|
||||
|
||||
int queue_to_use = 1;
|
||||
if (test_enqueue_wait_for_events) {
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 1, &event[event_count], NULL );
|
||||
test_error( error, "Unable to queue wait for events" );
|
||||
} else if (test_barrier) {
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 0, NULL, NULL);
|
||||
test_error( error, "Unable to queue barrier" );
|
||||
if (test_enqueue_wait_for_events)
|
||||
{
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 1,
|
||||
&event[event_count], NULL);
|
||||
test_error(error, "Unable to queue wait for events");
|
||||
}
|
||||
else if (test_barrier)
|
||||
{
|
||||
error =
|
||||
clEnqueueBarrierWithWaitList(queues[queue_to_use], 0, NULL, NULL);
|
||||
test_error(error, "Unable to queue barrier");
|
||||
}
|
||||
|
||||
for (loop_count=0; loop_count<TEST_COUNT; loop_count++) {
|
||||
for (loop_count = 0; loop_count < TEST_COUNT; loop_count++)
|
||||
{
|
||||
// Execute kernel 1
|
||||
event_count++;
|
||||
if (use_waitlist | use_marker) {
|
||||
if (DEBUG_OUT) log_info("clEnqueueNDRangeKernel(queues[%d], kernel1[%d], 1, NULL, threads, NULL, 1, &event[%d], &event[%d])\n", queue_to_use, loop_count, event_count-1, event_count);
|
||||
error = clEnqueueNDRangeKernel(queues[queue_to_use], kernel1[loop_count], 1, NULL, threads, NULL, 1, &event[event_count-1], &event[event_count]);
|
||||
} else {
|
||||
if (DEBUG_OUT) log_info("clEnqueueNDRangeKernel(queues[%d], kernel1[%d], 1, NULL, threads, NULL, 0, NULL, &event[%d])\n", queue_to_use, loop_count, event_count);
|
||||
error = clEnqueueNDRangeKernel(queues[queue_to_use], kernel1[loop_count], 1, NULL, threads, NULL, 0, NULL, &event[event_count]);
|
||||
if (use_waitlist | use_marker)
|
||||
{
|
||||
if (DEBUG_OUT)
|
||||
log_info("clEnqueueNDRangeKernel(queues[%d], kernel1[%d], 1, "
|
||||
"NULL, threads, NULL, 1, &event[%d], &event[%d])\n",
|
||||
queue_to_use, loop_count, event_count - 1,
|
||||
event_count);
|
||||
error = clEnqueueNDRangeKernel(
|
||||
queues[queue_to_use], kernel1[loop_count], 1, NULL, threads,
|
||||
NULL, 1, &event[event_count - 1], &event[event_count]);
|
||||
}
|
||||
if (error) {
|
||||
else
|
||||
{
|
||||
if (DEBUG_OUT)
|
||||
log_info("clEnqueueNDRangeKernel(queues[%d], kernel1[%d], 1, "
|
||||
"NULL, threads, NULL, 0, NULL, &event[%d])\n",
|
||||
queue_to_use, loop_count, event_count);
|
||||
error = clEnqueueNDRangeKernel(
|
||||
queues[queue_to_use], kernel1[loop_count], 1, NULL, threads,
|
||||
NULL, 0, NULL, &event[event_count]);
|
||||
}
|
||||
if (error)
|
||||
{
|
||||
log_info("\tLoop count %d\n", loop_count);
|
||||
print_error( error, "clEnqueueNDRangeKernel for kernel 1 failed");
|
||||
print_error(error, "clEnqueueNDRangeKernel for kernel 1 failed");
|
||||
return error;
|
||||
}
|
||||
expected_value *= 2;
|
||||
expected_if_only_queue[queue_to_use] *= 2;
|
||||
|
||||
// If we are using a marker, it needs to go in the same queue
|
||||
if (use_marker) {
|
||||
if (use_marker)
|
||||
{
|
||||
event_count++;
|
||||
if (DEBUG_OUT) log_info("clEnqueueMarker(queues[%d], event[%d])\n", queue_to_use, event_count);
|
||||
|
||||
#ifdef CL_VERSION_1_2
|
||||
error = clEnqueueMarkerWithWaitList(queues[queue_to_use], 0, NULL, &event[event_count]);
|
||||
#else
|
||||
error = clEnqueueMarker(queues[queue_to_use], &event[event_count]);
|
||||
#endif
|
||||
if (DEBUG_OUT)
|
||||
log_info("clEnqueueMarker(queues[%d], event[%d])\n",
|
||||
queue_to_use, event_count);
|
||||
|
||||
#ifdef CL_VERSION_1_2
|
||||
error = clEnqueueMarkerWithWaitList(queues[queue_to_use], 0, NULL,
|
||||
&event[event_count]);
|
||||
#else
|
||||
error = clEnqueueMarker(queues[queue_to_use], &event[event_count]);
|
||||
#endif
|
||||
}
|
||||
|
||||
// Pick the next queue to run
|
||||
if (RANDOMIZE)
|
||||
queue_to_use = rand()%2;
|
||||
queue_to_use = rand() % 2;
|
||||
else
|
||||
queue_to_use = (queue_to_use + 1)%2;
|
||||
queue_to_use = (queue_to_use + 1) % 2;
|
||||
|
||||
// Put in a barrier if requested
|
||||
if (test_enqueue_wait_for_events) {
|
||||
if (DEBUG_OUT) log_info("clEnqueueBarrierWithWaitList(queues[%d], 1, &event[%d], NULL)\n", queue_to_use, event_count);
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 1, &event[event_count], NULL);
|
||||
test_error( error, "Unable to queue wait for events" );
|
||||
} else if (test_barrier) {
|
||||
if (DEBUG_OUT) log_info("clEnqueueBarrierWithWaitList(queues[%d])\n", queue_to_use);
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 0, NULL, NULL);
|
||||
test_error( error, "Unable to queue barrier" );
|
||||
if (test_enqueue_wait_for_events)
|
||||
{
|
||||
if (DEBUG_OUT)
|
||||
log_info("clEnqueueBarrierWithWaitList(queues[%d], 1, "
|
||||
"&event[%d], NULL)\n",
|
||||
queue_to_use, event_count);
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 1,
|
||||
&event[event_count], NULL);
|
||||
test_error(error, "Unable to queue wait for events");
|
||||
}
|
||||
else if (test_barrier)
|
||||
{
|
||||
if (DEBUG_OUT)
|
||||
log_info("clEnqueueBarrierWithWaitList(queues[%d])\n",
|
||||
queue_to_use);
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 0, NULL,
|
||||
NULL);
|
||||
test_error(error, "Unable to queue barrier");
|
||||
}
|
||||
|
||||
// Execute Kernel 2
|
||||
event_count++;
|
||||
if (use_waitlist | use_marker) {
|
||||
if (DEBUG_OUT) log_info("clEnqueueNDRangeKernel(queues[%d], kernel2[%d], 1, NULL, threads, NULL, 1, &event[%d], &event[%d])\n", queue_to_use, loop_count, event_count-1, event_count);
|
||||
error = clEnqueueNDRangeKernel(queues[queue_to_use], kernel2[loop_count], 1, NULL, threads, NULL, 1, &event[event_count-1], &event[event_count]);
|
||||
} else {
|
||||
if (DEBUG_OUT) log_info("clEnqueueNDRangeKernel(queues[%d], kernel2[%d], 1, NULL, threads, NULL, 0, NULL, &event[%d])\n", queue_to_use, loop_count, event_count);
|
||||
error = clEnqueueNDRangeKernel(queues[queue_to_use], kernel2[loop_count], 1, NULL, threads, NULL, 0, NULL, &event[event_count]);
|
||||
if (use_waitlist | use_marker)
|
||||
{
|
||||
if (DEBUG_OUT)
|
||||
log_info("clEnqueueNDRangeKernel(queues[%d], kernel2[%d], 1, "
|
||||
"NULL, threads, NULL, 1, &event[%d], &event[%d])\n",
|
||||
queue_to_use, loop_count, event_count - 1,
|
||||
event_count);
|
||||
error = clEnqueueNDRangeKernel(
|
||||
queues[queue_to_use], kernel2[loop_count], 1, NULL, threads,
|
||||
NULL, 1, &event[event_count - 1], &event[event_count]);
|
||||
}
|
||||
if (error) {
|
||||
else
|
||||
{
|
||||
if (DEBUG_OUT)
|
||||
log_info("clEnqueueNDRangeKernel(queues[%d], kernel2[%d], 1, "
|
||||
"NULL, threads, NULL, 0, NULL, &event[%d])\n",
|
||||
queue_to_use, loop_count, event_count);
|
||||
error = clEnqueueNDRangeKernel(
|
||||
queues[queue_to_use], kernel2[loop_count], 1, NULL, threads,
|
||||
NULL, 0, NULL, &event[event_count]);
|
||||
}
|
||||
if (error)
|
||||
{
|
||||
log_info("\tLoop count %d\n", loop_count);
|
||||
print_error( error, "clEnqueueNDRangeKernel for kernel 2 failed");
|
||||
print_error(error, "clEnqueueNDRangeKernel for kernel 2 failed");
|
||||
return error;
|
||||
}
|
||||
expected_value--;
|
||||
expected_if_only_queue[queue_to_use]--;
|
||||
|
||||
// If we are using a marker, it needs to go in the same queue
|
||||
if (use_marker) {
|
||||
if (use_marker)
|
||||
{
|
||||
event_count++;
|
||||
if (DEBUG_OUT) log_info("clEnqueueMarker(queues[%d], event[%d])\n", queue_to_use, event_count);
|
||||
if (DEBUG_OUT)
|
||||
log_info("clEnqueueMarker(queues[%d], event[%d])\n",
|
||||
queue_to_use, event_count);
|
||||
|
||||
#ifdef CL_VERSION_1_2
|
||||
error = clEnqueueMarkerWithWaitList(queues[queue_to_use], 0, NULL, &event[event_count]);
|
||||
#else
|
||||
#ifdef CL_VERSION_1_2
|
||||
error = clEnqueueMarkerWithWaitList(queues[queue_to_use], 0, NULL,
|
||||
&event[event_count]);
|
||||
#else
|
||||
error = clEnqueueMarker(queues[queue_to_use], &event[event_count]);
|
||||
#endif
|
||||
#endif
|
||||
}
|
||||
|
||||
// Pick the next queue to run
|
||||
if (RANDOMIZE)
|
||||
queue_to_use = rand()%2;
|
||||
queue_to_use = rand() % 2;
|
||||
else
|
||||
queue_to_use = (queue_to_use + 1)%2;
|
||||
queue_to_use = (queue_to_use + 1) % 2;
|
||||
|
||||
// Put in a barrier if requested
|
||||
if (test_enqueue_wait_for_events) {
|
||||
if (DEBUG_OUT) log_info("clEnqueueBarrierWithWaitList(queues[%d], 1, &event[%d], NULL)\n", queue_to_use, event_count);
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 1, &event[event_count], NULL );
|
||||
test_error( error, "Unable to queue wait for events" );
|
||||
} else if (test_barrier) {
|
||||
if (DEBUG_OUT) log_info("clEnqueueBarrierWithWaitList(queues[%d])\n", queue_to_use);
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 0, NULL, NULL);
|
||||
test_error( error, "Unable to queue barrier" );
|
||||
if (test_enqueue_wait_for_events)
|
||||
{
|
||||
if (DEBUG_OUT)
|
||||
log_info("clEnqueueBarrierWithWaitList(queues[%d], 1, "
|
||||
"&event[%d], NULL)\n",
|
||||
queue_to_use, event_count);
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 1,
|
||||
&event[event_count], NULL);
|
||||
test_error(error, "Unable to queue wait for events");
|
||||
}
|
||||
else if (test_barrier)
|
||||
{
|
||||
if (DEBUG_OUT)
|
||||
log_info("clEnqueueBarrierWithWaitList(queues[%d])\n",
|
||||
queue_to_use);
|
||||
error = clEnqueueBarrierWithWaitList(queues[queue_to_use], 0, NULL,
|
||||
NULL);
|
||||
test_error(error, "Unable to queue barrier");
|
||||
}
|
||||
}
|
||||
|
||||
// Now finish up everything
|
||||
if (two_queues) {
|
||||
if (two_queues)
|
||||
{
|
||||
error = clFlush(queues[1]);
|
||||
test_error( error, "clFlush[1] failed");
|
||||
test_error(error, "clFlush[1] failed");
|
||||
}
|
||||
|
||||
error = clEnqueueReadBuffer(queues[0], data, CL_TRUE, 0, TEST_SIZE*sizeof(cl_int), values, 1, &event[event_count], NULL);
|
||||
error = clEnqueueReadBuffer(queues[0], data, CL_TRUE, 0,
|
||||
TEST_SIZE * sizeof(cl_int), values, 1,
|
||||
&event[event_count], NULL);
|
||||
|
||||
test_error(error, "clEnqueueReadBuffer failed");
|
||||
|
||||
failed = 0;
|
||||
for (i=0; i<(int)TEST_SIZE; i++)
|
||||
if (values[i] != expected_value) {
|
||||
for (i = 0; i < (int)TEST_SIZE; i++)
|
||||
if (values[i] != expected_value)
|
||||
{
|
||||
failed = 1;
|
||||
log_info("\tvalues[%d] = %d, expected %d (If only queue 1 accessed memory: %d only queue 2 accessed memory: %d)\n",
|
||||
i, values[i], expected_value, expected_if_only_queue[0], expected_if_only_queue[1]);
|
||||
log_info("\tvalues[%d] = %d, expected %d (If only queue 1 accessed "
|
||||
"memory: %d only queue 2 accessed memory: %d)\n",
|
||||
i, values[i], expected_value, expected_if_only_queue[0],
|
||||
expected_if_only_queue[1]);
|
||||
break;
|
||||
}
|
||||
|
||||
free(values);
|
||||
if (two_devices)
|
||||
free(two_device_ids);
|
||||
if (two_devices) free(two_device_ids);
|
||||
|
||||
return failed;
|
||||
}
|
||||
|
||||
int test( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements,
|
||||
int two_queues, int two_devices,
|
||||
int test_enqueue_wait_for_events, int test_barrier, int use_waitlists, int use_marker)
|
||||
int test(cl_device_id deviceID, cl_context context, cl_command_queue queue,
|
||||
int num_elements, int two_queues, int two_devices,
|
||||
int test_enqueue_wait_for_events, int test_barrier, int use_waitlists,
|
||||
int use_marker)
|
||||
{
|
||||
if( !checkDeviceForQueueSupport( deviceID, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE ) )
|
||||
if (!checkDeviceForQueueSupport(deviceID,
|
||||
CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE))
|
||||
{
|
||||
log_info( "WARNING: Device does not support out-of-order exec mode; skipping test.\n" );
|
||||
log_info("WARNING: Device does not support out-of-order exec mode; "
|
||||
"skipping test.\n");
|
||||
return 0;
|
||||
}
|
||||
|
||||
log_info("Running test for baseline results to determine if out-of-order execution can be detected...\n");
|
||||
int baseline_results = test_event_enqueue_wait_for_events_run_test(deviceID, context, queue, num_elements, two_queues, two_devices, 0, 0, 0, 0);
|
||||
if (baseline_results == 0) {
|
||||
log_info("Running test for baseline results to determine if out-of-order "
|
||||
"execution can be detected...\n");
|
||||
int baseline_results = test_event_enqueue_wait_for_events_run_test(
|
||||
deviceID, context, queue, num_elements, two_queues, two_devices, 0, 0,
|
||||
0, 0);
|
||||
if (baseline_results == 0)
|
||||
{
|
||||
if (test_enqueue_wait_for_events)
|
||||
log_info("WARNING: could not detect any out-of-order execution without using clEnqueueBarrierWithWaitList, so this test is not a valid test of out-of-order event dependencies.\n");
|
||||
log_info(
|
||||
"WARNING: could not detect any out-of-order execution without "
|
||||
"using clEnqueueBarrierWithWaitList, so this test is not a "
|
||||
"valid test of out-of-order event dependencies.\n");
|
||||
if (test_barrier)
|
||||
log_info("WARNING: could not detect any out-of-order execution without using clEnqueueBarrierWithWaitList, so this test is not a valid test of out-of-order event dependencies.\n");
|
||||
log_info(
|
||||
"WARNING: could not detect any out-of-order execution without "
|
||||
"using clEnqueueBarrierWithWaitList, so this test is not a "
|
||||
"valid test of out-of-order event dependencies.\n");
|
||||
if (use_waitlists)
|
||||
log_info("WARNING: could not detect any out-of-order execution without using waitlists, so this test is not a valid test of out-of-order event dependencies.\n");
|
||||
log_info("WARNING: could not detect any out-of-order execution "
|
||||
"without using waitlists, so this test is not a valid "
|
||||
"test of out-of-order event dependencies.\n");
|
||||
if (use_marker)
|
||||
log_info("WARNING: could not detect any out-of-order execution without using clEnqueueMarker, so this test is not a valid test of out-of-order event dependencies.\n");
|
||||
} else if (baseline_results == 1) {
|
||||
log_info("WARNING: could not detect any out-of-order execution "
|
||||
"without using clEnqueueMarker, so this test is not a "
|
||||
"valid test of out-of-order event dependencies.\n");
|
||||
}
|
||||
else if (baseline_results == 1)
|
||||
{
|
||||
if (test_enqueue_wait_for_events)
|
||||
log_info("Detected incorrect execution (possibly out-of-order) without clEnqueueBarrierWithWaitList. Test can be a valid test of out-of-order event dependencies.\n");
|
||||
log_info("Detected incorrect execution (possibly out-of-order) "
|
||||
"without clEnqueueBarrierWithWaitList. Test can be a "
|
||||
"valid test of out-of-order event dependencies.\n");
|
||||
if (test_barrier)
|
||||
log_info("Detected incorrect execution (possibly out-of-order) without clEnqueueBarrierWithWaitList. Test can be a valid test of out-of-order event dependencies.\n");
|
||||
log_info("Detected incorrect execution (possibly out-of-order) "
|
||||
"without clEnqueueBarrierWithWaitList. Test can be a "
|
||||
"valid test of out-of-order event dependencies.\n");
|
||||
if (use_waitlists)
|
||||
log_info("Detected incorrect execution (possibly out-of-order) without waitlists. Test can be a valid test of out-of-order event dependencies.\n");
|
||||
log_info("Detected incorrect execution (possibly out-of-order) "
|
||||
"without waitlists. Test can be a valid test of "
|
||||
"out-of-order event dependencies.\n");
|
||||
if (use_marker)
|
||||
log_info("Detected incorrect execution (possibly out-of-order) without clEnqueueMarker. Test can be a valid test of out-of-order event dependencies.\n");
|
||||
} else if( baseline_results == -1942 ) {
|
||||
log_info("Detected incorrect execution (possibly out-of-order) "
|
||||
"without clEnqueueMarker. Test can be a valid test of "
|
||||
"out-of-order event dependencies.\n");
|
||||
}
|
||||
else if (baseline_results == -1942)
|
||||
{
|
||||
// Just ignore and return (out-of-order exec mode not supported)
|
||||
return 0;
|
||||
} else {
|
||||
}
|
||||
else
|
||||
{
|
||||
print_error(baseline_results, "Baseline run failed");
|
||||
return baseline_results;
|
||||
}
|
||||
log_info("Running test for actual results...\n");
|
||||
return test_event_enqueue_wait_for_events_run_test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test_event_enqueue_wait_for_events_run_test(
|
||||
deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
}
|
||||
|
||||
|
||||
int test_out_of_order_event_waitlist_single_queue( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
int test_out_of_order_event_waitlist_single_queue(cl_device_id deviceID,
|
||||
cl_context context,
|
||||
cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
int two_queues = 0;
|
||||
int two_devices = 0;
|
||||
@@ -397,10 +543,15 @@ int test_out_of_order_event_waitlist_single_queue( cl_device_id deviceID, cl_con
|
||||
int test_barrier = 0;
|
||||
int use_waitlists = 1;
|
||||
int use_marker = 0;
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices, test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists,
|
||||
use_marker);
|
||||
}
|
||||
|
||||
int test_out_of_order_event_waitlist_multi_queue( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
int test_out_of_order_event_waitlist_multi_queue(cl_device_id deviceID,
|
||||
cl_context context,
|
||||
cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
int two_queues = 1;
|
||||
int two_devices = 0;
|
||||
@@ -408,10 +559,14 @@ int test_out_of_order_event_waitlist_multi_queue( cl_device_id deviceID, cl_cont
|
||||
int test_barrier = 0;
|
||||
int use_waitlists = 1;
|
||||
int use_marker = 0;
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices, test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists,
|
||||
use_marker);
|
||||
}
|
||||
|
||||
int test_out_of_order_event_waitlist_multi_queue_multi_device( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
int test_out_of_order_event_waitlist_multi_queue_multi_device(
|
||||
cl_device_id deviceID, cl_context context, cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
int two_queues = 1;
|
||||
int two_devices = 1;
|
||||
@@ -419,11 +574,15 @@ int test_out_of_order_event_waitlist_multi_queue_multi_device( cl_device_id devi
|
||||
int test_barrier = 0;
|
||||
int use_waitlists = 1;
|
||||
int use_marker = 0;
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices, test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists,
|
||||
use_marker);
|
||||
}
|
||||
|
||||
|
||||
int test_out_of_order_event_enqueue_wait_for_events_single_queue( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
int test_out_of_order_event_enqueue_wait_for_events_single_queue(
|
||||
cl_device_id deviceID, cl_context context, cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
int two_queues = 0;
|
||||
int two_devices = 0;
|
||||
@@ -431,10 +590,14 @@ int test_out_of_order_event_enqueue_wait_for_events_single_queue( cl_device_id d
|
||||
int test_barrier = 0;
|
||||
int use_waitlists = 0;
|
||||
int use_marker = 0;
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices, test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists,
|
||||
use_marker);
|
||||
}
|
||||
|
||||
int test_out_of_order_event_enqueue_wait_for_events_multi_queue( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
int test_out_of_order_event_enqueue_wait_for_events_multi_queue(
|
||||
cl_device_id deviceID, cl_context context, cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
int two_queues = 1;
|
||||
int two_devices = 0;
|
||||
@@ -442,11 +605,15 @@ int test_out_of_order_event_enqueue_wait_for_events_multi_queue( cl_device_id de
|
||||
int test_barrier = 0;
|
||||
int use_waitlists = 0;
|
||||
int use_marker = 0;
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices, test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists,
|
||||
use_marker);
|
||||
}
|
||||
|
||||
|
||||
int test_out_of_order_event_enqueue_wait_for_events_multi_queue_multi_device( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
int test_out_of_order_event_enqueue_wait_for_events_multi_queue_multi_device(
|
||||
cl_device_id deviceID, cl_context context, cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
int two_queues = 1;
|
||||
int two_devices = 1;
|
||||
@@ -454,13 +621,16 @@ int test_out_of_order_event_enqueue_wait_for_events_multi_queue_multi_device( cl
|
||||
int test_barrier = 0;
|
||||
int use_waitlists = 0;
|
||||
int use_marker = 0;
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices, test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists,
|
||||
use_marker);
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
int test_out_of_order_event_enqueue_barrier_single_queue( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
int test_out_of_order_event_enqueue_barrier_single_queue(cl_device_id deviceID,
|
||||
cl_context context,
|
||||
cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
int two_queues = 0;
|
||||
int two_devices = 0;
|
||||
@@ -468,11 +638,16 @@ int test_out_of_order_event_enqueue_barrier_single_queue( cl_device_id deviceID,
|
||||
int test_barrier = 1;
|
||||
int use_waitlists = 0;
|
||||
int use_marker = 0;
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices, test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists,
|
||||
use_marker);
|
||||
}
|
||||
|
||||
|
||||
int test_out_of_order_event_enqueue_marker_single_queue( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
int test_out_of_order_event_enqueue_marker_single_queue(cl_device_id deviceID,
|
||||
cl_context context,
|
||||
cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
int two_queues = 0;
|
||||
int two_devices = 0;
|
||||
@@ -480,10 +655,15 @@ int test_out_of_order_event_enqueue_marker_single_queue( cl_device_id deviceID,
|
||||
int test_barrier = 0;
|
||||
int use_waitlists = 0;
|
||||
int use_marker = 1;
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices, test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists,
|
||||
use_marker);
|
||||
}
|
||||
|
||||
int test_out_of_order_event_enqueue_marker_multi_queue( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
int test_out_of_order_event_enqueue_marker_multi_queue(cl_device_id deviceID,
|
||||
cl_context context,
|
||||
cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
int two_queues = 1;
|
||||
int two_devices = 0;
|
||||
@@ -491,11 +671,15 @@ int test_out_of_order_event_enqueue_marker_multi_queue( cl_device_id deviceID, c
|
||||
int test_barrier = 0;
|
||||
int use_waitlists = 0;
|
||||
int use_marker = 1;
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices, test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists,
|
||||
use_marker);
|
||||
}
|
||||
|
||||
|
||||
int test_out_of_order_event_enqueue_marker_multi_queue_multi_device( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
|
||||
int test_out_of_order_event_enqueue_marker_multi_queue_multi_device(
|
||||
cl_device_id deviceID, cl_context context, cl_command_queue queue,
|
||||
int num_elements)
|
||||
{
|
||||
int two_queues = 1;
|
||||
int two_devices = 1;
|
||||
@@ -503,7 +687,7 @@ int test_out_of_order_event_enqueue_marker_multi_queue_multi_device( cl_device_i
|
||||
int test_barrier = 0;
|
||||
int use_waitlists = 0;
|
||||
int use_marker = 1;
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices, test_enqueue_wait_for_events, test_barrier, use_waitlists, use_marker);
|
||||
return test(deviceID, context, queue, num_elements, two_queues, two_devices,
|
||||
test_enqueue_wait_for_events, test_barrier, use_waitlists,
|
||||
use_marker);
|
||||
}
|
||||
|
||||
|
||||
|
||||
Reference in New Issue
Block a user