#include "testBase.h"
#include "harness/testHarness.h"
const char *write_kernels[] = {
"__kernel void write_up(__global int *dst, int length)\n"
"{\n"
"\n"
" dst[get_global_id(0)] *= 2;\n"
"\n"
"}\n"
"__kernel void write_down(__global int *dst, int length)\n"
"{\n"
"\n"
" dst[get_global_id(0)]--;\n"
"\n"
"}\n"
};
#define TEST_SIZE 10000
#define TEST_COUNT 100
#define RANDOMIZE 1
#define DEBUG_OUT 0
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.
*/
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};
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];
clCommandQueueWrapper queueWrappers[2];
clContextWrapper context_to_use;
clMemWrapper data;
clProgramWrapper program;
clKernelWrapper kernel1[TEST_COUNT], kernel2[TEST_COUNT];
clEventWrapper event[TEST_COUNT*4+2];
if (test_enqueue_wait_for_events)
log_info("\tTesting with clEnqueueBarrierWithWaitList as barrier function.\n");
if (test_barrier)
log_info("\tTesting with clEnqueueBarrierWithWaitList as barrier function.\n");
if (use_waitlist)
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");
return -1;
}
error = clGetPlatformIDs(1, &platform, NULL);
test_error(error, "clGetPlatformIDs failed.");
cl_device_id *two_device_ids;
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) {
log_info("Failed to obtain two devices. Test can not run.\n");
free(two_device_ids);
return 0;
}
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.");
if (type & CL_DEVICE_TYPE_CPU)
log_info("\tDevice %d is CL_DEVICE_TYPE_CPU.\n", i);
if (type & CL_DEVICE_TYPE_GPU)
log_info("\tDevice %d is CL_DEVICE_TYPE_GPU.\n", i);
if (type & CL_DEVICE_TYPE_ACCELERATOR)
log_info("\tDevice %d is CL_DEVICE_TYPE_ACCELERATOR.\n", i);
if (type & CL_DEVICE_TYPE_DEFAULT)
log_info("\tDevice %d is CL_DEVICE_TYPE_DEFAULT.\n", i);
}
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);
test_error(error, "clCreateContext failed for one device.");
log_info("\tTesting with one device.\n");
}
cl_command_queue_properties props = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
if (two_queues) {
if (two_devices)
{
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" );
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.");
}
else
{
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);
test_error(error, "clCreateCommandQueue for second queue failed.");
}
queues[0] = queueWrappers[0];
queues[1] = queueWrappers[1];
log_info("\tTesting with two queues.\n");
}
else
{
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];
log_info("\tTesting with one queue.\n");
}
data = clCreateBuffer(context_to_use, CL_MEM_READ_WRITE, TEST_SIZE*sizeof(cl_int), NULL, &error);
test_error( error, "clCreateBuffer 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;
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");
for (i=1; i<TEST_COUNT; i++) {
kernel1[i] = clCreateKernel(program, "write_up", &error);
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");
}
for (i=0; i<TEST_COUNT; i++) {
kernel2[i] = clCreateKernel(program, "write_down", &error);
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");
}
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");
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" );
}
for (loop_count=0; loop_count<TEST_COUNT; loop_count++) {
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 (error) {
log_info("\tLoop count %d\n", loop_count);
print_error( error, "clEnqueueNDRangeKernel for kernel 1 failed");
return error;
}
expected_value *= 2;
expected_if_only_queue[queue_to_use] *= 2;
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 (RANDOMIZE)
queue_to_use = rand()%2;
else
queue_to_use = (queue_to_use + 1)%2;
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" );
}
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 (error) {
log_info("\tLoop count %d\n", loop_count);
print_error( error, "clEnqueueNDRangeKernel for kernel 2 failed");
return error;
}
expected_value--;
expected_if_only_queue[queue_to_use]--;
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 (RANDOMIZE)
queue_to_use = rand()%2;
else
queue_to_use = (queue_to_use + 1)%2;
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 (two_queues) {
error = clFlush(queues[1]);
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);
test_error(error, "clEnqueueReadBuffer failed");
failed = 0;
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]);
break;
}
free(values);
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)
{
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" );
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) {
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");
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");
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");
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) {
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");
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");
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");
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 ) {
return 0;
} 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);
}
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;
int test_enqueue_wait_for_events = 0;
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);
}
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;
int test_enqueue_wait_for_events = 0;
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);
}
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;
int test_enqueue_wait_for_events = 0;
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);
}
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;
int test_enqueue_wait_for_events = 1;
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);
}
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;
int test_enqueue_wait_for_events = 1;
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);
}
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;
int test_enqueue_wait_for_events = 1;
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);
}
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;
int test_enqueue_wait_for_events = 0;
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);
}
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;
int test_enqueue_wait_for_events = 0;
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);
}
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;
int test_enqueue_wait_for_events = 0;
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);
}
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;
int test_enqueue_wait_for_events = 0;
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);
}