//
// 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/os_helpers.h"

const char *preprocessor_test_kernel[] = {
"__kernel void sample_test(__global int *dst)\n"
"{\n"
"    dst[0] = TEST_MACRO;\n"
"\n"
"}\n" };

const char *preprocessor_existence_test_kernel[] = {
    "__kernel void sample_test(__global int *dst)\n"
    "{\n"
    "#ifdef TEST_MACRO\n"
    "    dst[0] = 42;\n"
    "#else\n"
    "    dst[0] = 24;\n"
    "#endif\n"
    "\n"
    "}\n" };

const char *include_test_kernel[] = {
"#include \"./testIncludeFile.h\"\n"
"__kernel void sample_test(__global int *dst)\n"
"{\n"
"    dst[0] = HEADER_FOUND;\n"
"\n"
"}\n" };

const char *options_test_kernel[] = {
    "__kernel void sample_test(__global float *src, __global int *dst)\n"
    "{\n"
    "    size_t tid = get_global_id(0);\n"
    "    dst[tid] = (int)src[tid];\n"
    "}\n"
};

const char *optimization_options[] = {
    "-cl-single-precision-constant",
    "-cl-denorms-are-zero",
    "-cl-opt-disable",
    "-cl-mad-enable",
    "-cl-no-signed-zeros",
    "-cl-unsafe-math-optimizations",
    "-cl-finite-math-only",
    "-cl-fast-relaxed-math",
    "-w",
    "-Werror",
    };

cl_int get_result_from_program( cl_context context, cl_command_queue queue, cl_program program, cl_int *outValue )
{
    cl_int error;
    clKernelWrapper kernel = clCreateKernel( program, "sample_test", &error );
    test_error( error, "Unable to create kernel from program" );

    clMemWrapper outStream;
    outStream = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL,
                               &error);
    test_error( error, "Unable to create test buffer" );

    error = clSetKernelArg( kernel, 0, sizeof( outStream ), &outStream );
    test_error( error, "Unable to set kernel argument" );

    size_t threads[1] = { 1 };

    error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
    test_error( error, "Unable to execute test kernel" );

    error = clEnqueueReadBuffer( queue, outStream, true, 0, sizeof( cl_int ), outValue, 0, NULL, NULL );
    test_error( error, "Unable to read output array!" );

    return CL_SUCCESS;
}

int test_options_build_optimizations(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
    int error;
    cl_build_status status;

    for(size_t i = 0; i < sizeof(optimization_options) / (sizeof(char*)); i++) {

        clProgramWrapper program;
        error = create_single_kernel_helper_create_program(context, &program, 1, options_test_kernel, optimization_options[i]);
        if( program == NULL || error != CL_SUCCESS )
        {
            log_error( "ERROR: Unable to create reference program!\n" );
            return -1;
        }

        /* Build with the macro defined */
        log_info("Testing optimization option '%s'\n", optimization_options[i]);
        error = clBuildProgram( program, 1, &deviceID, optimization_options[i], NULL, NULL );
        test_error( error, "Test program did not properly build" );

        error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
        test_error( error, "Unable to get program build status" );

        if( (int)status != CL_BUILD_SUCCESS )
        {
            log_info("Building with optimization option '%s' failed to compile!\n", optimization_options[i]);
            print_error( error, "Failed to build with optimization defined")
            return -1;
        }
    }
    return 0;
}

int test_options_build_macro(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
    int error;
    clProgramWrapper program;
    cl_build_status status;


    program = clCreateProgramWithSource( context, 1, preprocessor_test_kernel, NULL, &error );
    if( program == NULL || error != CL_SUCCESS )
    {
        log_error( "ERROR: Unable to create reference program!\n" );
        return -1;
    }

    /* Build with the macro defined */
    error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO=1 ", NULL, NULL );
    test_error( error, "Test program did not properly build" );

    error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
    test_error( error, "Unable to get program build status" );

    if( (int)status != CL_BUILD_SUCCESS )
    {
        print_error( error, "Failed to build with macro defined" );
        return -1;
    }


    // Go ahead and run the program to verify results
    cl_int firstResult, secondResult;

    error = get_result_from_program( context, queue, program, &firstResult );
    test_error( error, "Unable to get result from first program" );

    if( firstResult != 1 )
    {
        log_error( "ERROR: Result from first program did not validate! (Expected 1, got %d)\n", firstResult );
        return -1;
    }

    // Rebuild with a different value for the define macro, to make sure caching behaves properly
    error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO=5 ", NULL, NULL );
    test_error( error, "Test program did not properly rebuild" );

    error = get_result_from_program( context, queue, program, &secondResult );
    test_error( error, "Unable to get result from second program" );

    if( secondResult != 5 )
    {
        if( secondResult == firstResult )
            log_error( "ERROR: Program result did not change with device macro change (program was not recompiled)!\n" );
        else
            log_error( "ERROR: Result from second program did not validate! (Expected 5, got %d)\n", secondResult );
        return -1;
    }

    return 0;
}

int test_options_build_macro_existence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
    int error;
    clProgramWrapper program;


    // In this case, the program should still run without the macro, but it should give a different result
    program = clCreateProgramWithSource( context, 1, preprocessor_existence_test_kernel, NULL, &error );
    if( program == NULL || error != CL_SUCCESS )
    {
        log_error( "ERROR: Unable to create reference program!\n" );
        return -1;
    }

    /* Build without the macro defined */
    error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
    test_error( error, "Test program did not properly build" );

    // Go ahead and run the program to verify results
    cl_int firstResult, secondResult;

    error = get_result_from_program( context, queue, program, &firstResult );
    test_error( error, "Unable to get result from first program" );

    if( firstResult != 24 )
    {
        log_error( "ERROR: Result from first program did not validate! (Expected 24, got %d)\n", firstResult );
        return -1;
    }

    // Now compile again with the macro defined and verify a change in results
    error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO", NULL, NULL );
    test_error( error, "Test program did not properly build" );

    error = get_result_from_program( context, queue, program, &secondResult );
    test_error( error, "Unable to get result from second program" );

    if( secondResult != 42 )
    {
        if( secondResult == firstResult )
            log_error( "ERROR: Program result did not change with device macro addition (program was not recompiled)!\n" );
        else
            log_error( "ERROR: Result from second program did not validate! (Expected 42, got %d)\n", secondResult );
        return -1;
    }

    return 0;
}

int test_options_include_directory(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
    int error;

    std::string sep  = dir_sep();
    std::string path = exe_dir();    // Directory where test executable is located.
    std::string include_dir;

    clProgramWrapper program;
    cl_build_status status;

    /* Try compiling the program first without the directory included Should fail. */
    program = clCreateProgramWithSource( context, 1, include_test_kernel, NULL, &error );
    if( program == NULL || error != CL_SUCCESS )
    {
        log_error( "ERROR: Unable to create reference program!\n" );
        return -1;
    }

    /* Build with the include directory defined */
    include_dir = "-I " + path + sep + "includeTestDirectory";

//    log_info("%s\n", include_dir);
    error = clBuildProgram( program, 1, &deviceID, include_dir.c_str(), NULL, NULL );
    test_error( error, "Test program did not properly build" );

    error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
    test_error( error, "Unable to get program build status" );

    if( (int)status != CL_BUILD_SUCCESS )
    {
        print_error( error, "Failed to build with include directory" );
        return -1;
    }

    // Go ahead and run the program to verify results
    cl_int firstResult, secondResult;

    error = get_result_from_program( context, queue, program, &firstResult );
    test_error( error, "Unable to get result from first program" );

    if( firstResult != 12 )
    {
        log_error( "ERROR: Result from first program did not validate! (Expected 12, got %d)\n", firstResult );
        return -1;
    }

    // Rebuild with a different include directory
    include_dir = "-I " + path + sep + "secondIncludeTestDirectory";
    error = clBuildProgram( program, 1, &deviceID, include_dir.c_str(), NULL, NULL );
    test_error( error, "Test program did not properly rebuild" );

    error = get_result_from_program( context, queue, program, &secondResult );
    test_error( error, "Unable to get result from second program" );

    if( secondResult != 42 )
    {
        if( secondResult == firstResult )
            log_error( "ERROR: Program result did not change with include path change (program was not recompiled)!\n" );
        else
            log_error( "ERROR: Result from second program did not validate! (Expected 42, got %d)\n", secondResult );
        return -1;
    }

    return 0;
}

const char *denorm_test_kernel[] = {
    "__kernel void sample_test( float src1, float src2, __global float *dst)\n"
    "{\n"
    "    dst[ 0 ] = src1 + src2;\n"
    "\n"
    "}\n" };

cl_int get_float_result_from_program( cl_context context, cl_command_queue queue, cl_program program, cl_float inA, cl_float inB, cl_float *outValue )
{
    cl_int error;

    clKernelWrapper kernel = clCreateKernel( program, "sample_test", &error );
    test_error( error, "Unable to create kernel from program" );

    clMemWrapper outStream = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                            sizeof(cl_float), NULL, &error);
    test_error( error, "Unable to create test buffer" );

    error = clSetKernelArg( kernel, 0, sizeof( cl_float ), &inA );
    test_error( error, "Unable to set kernel argument" );

    error = clSetKernelArg( kernel, 1, sizeof( cl_float ), &inB );
    test_error( error, "Unable to set kernel argument" );

    error = clSetKernelArg( kernel, 2, sizeof( outStream ), &outStream );
    test_error( error, "Unable to set kernel argument" );

    size_t threads[1] = { 1 };

    error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
    test_error( error, "Unable to execute test kernel" );

    error = clEnqueueReadBuffer( queue, outStream, true, 0, sizeof( cl_float ), outValue, 0, NULL, NULL );
    test_error( error, "Unable to read output array!" );

    return CL_SUCCESS;
}

int test_options_denorm_cache(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
    int error;

    clProgramWrapper program;
    cl_build_status status;


    // If denorms aren't even supported, testing this flag is pointless
    cl_device_fp_config floatCaps = 0;
    error = clGetDeviceInfo( deviceID, CL_DEVICE_SINGLE_FP_CONFIG, sizeof(floatCaps), &floatCaps,  NULL);
    test_error( error, "Unable to get device FP config" );
    if( ( floatCaps & CL_FP_DENORM ) == 0 )
    {
        log_info( "Device does not support denormalized single-precision floats; skipping test.\n" );
        return 0;
    }

    program = clCreateProgramWithSource( context, 1, denorm_test_kernel, NULL, &error );
    test_error( error, "Unable to create test program" );

    // Build first WITH the denorm flush flag
    error = clBuildProgram( program, 1, &deviceID, "-cl-denorms-are-zero", NULL, NULL );
    test_error( error, "Test program did not properly build" );

    error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
    test_error( error, "Unable to get program build status" );

    if( (int)status != CL_BUILD_SUCCESS )
    {
        print_error( error, "Failed to build with include directory" );
        return -1;
    }

    // Note: the following in floating point is a subnormal number, thus adding two of them together
    // should give us a subnormalized result. If denormals are flushed to zero, however, it'll give us zero instead
    uint32_t intSubnormal = 0x00000001;
    cl_float *input = (cl_float *)&intSubnormal;
    cl_float firstResult, secondResult;

    error = get_float_result_from_program( context, queue, program, *input, *input, &firstResult );
    test_error( error, "Unable to get result from first program" );

    // Note: since -cl-denorms-are-zero is a HINT, not a requirement, the result we got could
    // either be subnormal (hint ignored) or zero (hint respected). Since either is technically
    // valid, there isn't anything we can to do validate results for now

    // Rebuild without flushing flag set
    error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
    test_error( error, "Test program did not properly rebuild" );

    error = get_float_result_from_program( context, queue, program, *input, *input, &secondResult );
    test_error( error, "Unable to get result from second program" );

    // Now, there are three possiblities here:
    // 1. The denorms-are-zero hint is not respected, in which case the first and second result will be identical
    // 2. The hint is respected, and the program was properly rebuilt, in which case the first result will be zero and the second non-zero
    // 3. The hint is respected, but the program was not properly rebuilt, in which case both results will be zero
    // 3 is the only error condition we need to look for
    uint32_t *fPtr = (uint32_t *)&firstResult;
    uint32_t *sPtr = (uint32_t *)&secondResult;

    if( ( *fPtr == 0 ) && ( *sPtr == 0 ) )
    {
        log_error( "ERROR: Program result didn't change when -cl-denorms-are-zero flag was removed.\n"
                  "First result (should be zero): 0x%08x, Second result (should be non-zero): 0x%08x\n",
                  *fPtr, *sPtr );
        return -1;
    }

    return 0;
}