//
// 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
//
//    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/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];    // 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

    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.");

    // 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);
        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");
    }

    // 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) {
        // Get a second queue
        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
        {
            // 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);
            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
        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);
        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");
    }


    // 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");


    // 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");
    expected_value = 0;

    // Build the kernels
    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");
    }

    // 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");
    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++) {
        // 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 (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 we are using a marker, it needs to go in the same queue
        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

        }

        // Pick the next queue to run
        if (RANDOMIZE)
            queue_to_use = rand()%2;
        else
            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" );
        }

        // 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 (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 we are using a marker, it needs to go in the same queue
        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
        }

        // Pick the next queue to run
        if (RANDOMIZE)
            queue_to_use = rand()%2;
        else
            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" );
        }
    }

    // Now finish up everything
    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 ) {
        // Just ignore and return (out-of-order exec mode not supported)
        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);
}