//
// 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/conversions.h"
#include "harness/typeWrappers.h"
#include "harness/testHarness.h"

#include "structs.h"

#include "defines.h"

#include "type_replacer.h"


/*
 test_step_type,
 test_step_var,
 test_step_typedef_type,
 test_step_typedef_var,
 */


int test_step_internal(cl_device_id deviceID, cl_context context,
                       cl_command_queue queue, const char* pattern,
                       const char* testName)
{
    int err;
    int typeIdx, vecSizeIdx;

    char tempBuffer[2048];

    clState* pClState = newClState(deviceID, context, queue);
    bufferStruct* pBuffers =
        newBufferStruct(BUFFER_SIZE, BUFFER_SIZE, pClState);

    if (pBuffers == NULL)
    {
        destroyClState(pClState);
        vlog_error("%s : Could not create buffer\n", testName);
        return -1;
    }

    // detect whether profile of the device is embedded
    char profile[1024] = "";
    err = clGetDeviceInfo(deviceID, CL_DEVICE_PROFILE, sizeof(profile), profile,
                          NULL);
    if (err)
    {
        print_error(err, "clGetDeviceInfo for CL_DEVICE_PROFILE failed\n");
        return -1;
    }
    gIsEmbedded = NULL != strstr(profile, "EMBEDDED_PROFILE");

    for (typeIdx = 0; types[typeIdx] != kNumExplicitTypes; ++typeIdx)
    {
        if (types[typeIdx] == kDouble)
        {
            // If we're testing doubles, we need to check for support first
            if (!is_extension_available(deviceID, "cl_khr_fp64"))
            {
                log_info("Not testing doubles (unsupported on this device)\n");
                continue;
            }
        }

        if (types[typeIdx] == kLong || types[typeIdx] == kULong)
        {
            // If we're testing long/ulong, we need to check for embedded
            // support
            if (gIsEmbedded
                && !is_extension_available(deviceID, "cles_khr_int64"))
            {
                log_info("Not testing longs (unsupported on this embedded "
                         "device)\n");
                continue;
            }
        }

        char srcBuffer[2048];

        doSingleReplace(tempBuffer, 2048, pattern, ".EXTENSIONS.",
                        types[typeIdx] == kDouble
                            ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable"
                            : "");

        for (vecSizeIdx = 0; vecSizeIdx < NUM_VECTOR_SIZES; ++vecSizeIdx)
        {
            doReplace(srcBuffer, 2048, tempBuffer, ".TYPE.",
                      g_arrTypeNames[typeIdx], ".NUM.",
                      g_arrVecSizeNames[vecSizeIdx]);

            if (srcBuffer[0] == '\0')
            {
                vlog_error("%s: failed to fill source buf for type %s%s\n",
                           testName, g_arrTypeNames[typeIdx],
                           g_arrVecSizeNames[vecSizeIdx]);
                destroyBufferStruct(pBuffers, pClState);
                destroyClState(pClState);
                return -1;
            }

            err = clStateMakeProgram(pClState, srcBuffer, testName);
            if (err)
            {
                vlog_error("%s: Error compiling \"\n%s\n\"", testName,
                           srcBuffer);
                destroyBufferStruct(pBuffers, pClState);
                destroyClState(pClState);
                return -1;
            }

            err = pushArgs(pBuffers, pClState);
            if (err != 0)
            {
                vlog_error("%s: failed to push args %s%s\n", testName,
                           g_arrTypeNames[typeIdx],
                           g_arrVecSizeNames[vecSizeIdx]);
                destroyBufferStruct(pBuffers, pClState);
                destroyClState(pClState);
                return -1;
            }

            // now we run the kernel
            err = runKernel(pClState, 1024);
            if (err != 0)
            {
                vlog_error("%s: runKernel fail (%ld threads) %s%s\n", testName,
                           pClState->m_numThreads, g_arrTypeNames[typeIdx],
                           g_arrVecSizeNames[vecSizeIdx]);
                destroyBufferStruct(pBuffers, pClState);
                destroyClState(pClState);
                return -1;
            }

            err = retrieveResults(pBuffers, pClState);
            if (err != 0)
            {
                vlog_error("%s: failed to retrieve results %s%s\n", testName,
                           g_arrTypeNames[typeIdx],
                           g_arrVecSizeNames[vecSizeIdx]);
                destroyBufferStruct(pBuffers, pClState);
                destroyClState(pClState);
                return -1;
            }

            err = checkCorrectnessStep(pBuffers, pClState,
                                       g_arrTypeSizes[typeIdx],
                                       g_arrVecSizes[vecSizeIdx]);

            if (err != 0)
            {
                vlog_error("%s: incorrect results %s%s\n", testName,
                           g_arrTypeNames[typeIdx],
                           g_arrVecSizeNames[vecSizeIdx]);
                vlog_error("%s: Source was \"\n%s\n\"", testName, srcBuffer);
                destroyBufferStruct(pBuffers, pClState);
                destroyClState(pClState);
                return -1;
            }
        }
    }

    destroyBufferStruct(pBuffers, pClState);

    destroyClState(pClState);


    // vlog_error("%s : implementation incomplete : FAIL\n", testName);
    return 0; // -1; // fails on account of not being written.
}

static const char* patterns[] = {
    ".EXTENSIONS.\n"
    "__kernel void test_step_type(__global .TYPE..NUM. *source, __global int "
    "*dest)\n"
    "{\n"
    "    int  tid = get_global_id(0);\n"
    "    dest[tid] = vec_step(.TYPE..NUM.);\n"
    "\n"
    "}\n",

    ".EXTENSIONS.\n"
    "__kernel void test_step_var(__global .TYPE..NUM. *source, __global int "
    "*dest)\n"
    "{\n"
    "    int  tid = get_global_id(0);\n"
    "    dest[tid] = vec_step(source[tid]);\n"
    "\n"
    "}\n",

    ".EXTENSIONS.\n"
    " typedef .TYPE..NUM. TypeToTest;\n"
    "__kernel void test_step_typedef_type(__global TypeToTest *source, "
    "__global int *dest)\n"
    "{\n"
    "    int  tid = get_global_id(0);\n"
    "    dest[tid] = vec_step(TypeToTest);\n"
    "\n"
    "}\n",

    ".EXTENSIONS.\n"
    " typedef .TYPE..NUM. TypeToTest;\n"
    "__kernel void test_step_typedef_var(__global TypeToTest *source, __global "
    "int *dest)\n"
    "{\n"
    "    int  tid = get_global_id(0);\n"
    "    dest[tid] = vec_step(source[tid]);\n"
    "\n"
    "}\n",
};

/*
 test_step_type,
 test_step_var,
 test_step_typedef_type,
 test_step_typedef_var,
 */

int test_step_type(cl_device_id deviceID, cl_context context,
                   cl_command_queue queue, int num_elements)
{
    return test_step_internal(deviceID, context, queue, patterns[0],
                              "test_step_type");
}

int test_step_var(cl_device_id deviceID, cl_context context,
                  cl_command_queue queue, int num_elements)
{
    return test_step_internal(deviceID, context, queue, patterns[1],
                              "test_step_var");
}

int test_step_typedef_type(cl_device_id deviceID, cl_context context,
                           cl_command_queue queue, int num_elements)
{
    return test_step_internal(deviceID, context, queue, patterns[2],
                              "test_step_typedef_type");
}

int test_step_typedef_var(cl_device_id deviceID, cl_context context,
                          cl_command_queue queue, int num_elements)
{
    return test_step_internal(deviceID, context, queue, patterns[3],
                              "test_step_typedef_var");
}