/*
// Copyright (c) 2016 Intel Corporation
//
// 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 "common_kernel_base.h"
#include <iostream>

#if defined __INTEL_COMPILER
#pragma warning disable: 177
#endif

namespace kernel_selector 
{
    namespace {

        class CodeBuilder
        {
            std::ostringstream oss;
            std::string code;
            std::vector<std::string> defined_macroses;

            CodeBuilder& register_macro(const std::string& name)
            {
                assert(std::count(defined_macroses.begin(), defined_macroses.end(), name) == 0);
                defined_macroses.push_back(name);
                return *this;
            }

        public:
            CodeBuilder& set_code(const std::string& c)
            {
                assert(code.empty());
                code = c;
                return *this;
            }

            CodeBuilder& add_line(const std::string& line) {
                oss << line << "\n";
                return *this;
            }

            CodeBuilder& decoration_macro(const std::string& name, const std::string& prefix, const std::string& postfix, const std::string& name_prefix = std::string())
            {
                oss << "#define " << name << "(name) " << prefix << " " + name_prefix + "_##" + "name" << (postfix.empty() ? "" : "##_") << postfix << std::endl;
                return register_macro(name);
            }


            CodeBuilder& value_macro(const std::string& name, const std::string& value)
            {
                oss << "#define " << name << " " << value << std::endl;
                return register_macro(name.substr(0, name.find('(')));
            }

            std::string str()
            {
                std::ostringstream os;
                os << oss.str();
                os << code << std::endl;
                return os.str();
            }
        };
    }

    std::string common_kernel_base::GetEntryPoint(const std::string& templateName, const std::string& layerID, const optional_params& options) const
    {
        std::string kernelID = layerID;

        if (kernelID.empty() || !options.meaningfulKernelsNames)
        {
            kernelID = templateName;
        }

        std::replace(kernelID.begin(), kernelID.end(), '.', '_');
        std::replace(kernelID.begin(), kernelID.end(), '/', '_');

        kernelID += "_" + std::to_string(UniqeID());

        return kernelID;
    }

    std::string common_kernel_base::CreateJit(const std::string& template_name, const JitConstants& constants, const std::string& kernel_id) const
    {
        class CodeBuilder code;
        code.add_line("\n//====================================================")
            .add_line("// Kernel template: " + template_name + " ")
            .add_line("// Kernel name: " + kernel_id)
            .value_macro("KERNEL(name)", "__kernel void " + kernel_id)
            .decoration_macro("FUNC", "", kernel_id)
            .decoration_macro("FUNC_CALL", "", kernel_id);
        
        for (auto& definition : constants.GetDefinitions())
        {
            code.value_macro(definition.first, definition.second);
        }

        std::string jit = code.str();

        return jit;
    }

    Arguments common_kernel_base::GetArgsDesc(uint32_t num_of_input, bool use_weights, bool use_bias, bool use_quantization, bool use_output_calibration) const
    {
        Arguments args;

        for (uint32_t i = 0; i < num_of_input; i++)
        {
            args.push_back({ ArgumentDescriptor::Types::INPUT, i });
        }

        args.push_back({ ArgumentDescriptor::Types::OUTPUT, 0 });

        if (use_weights)
        {
            args.push_back({ ArgumentDescriptor::Types::WEIGHTS, 0 });
        }

        if (use_bias)
        {
            args.push_back({ ArgumentDescriptor::Types::BIAS, 0 });
        }

        if (use_quantization && use_weights)
        {
            args.push_back({ ArgumentDescriptor::Types::WEIGHTS_QUANTIZATION_FACTORS, 0 });
        }

        if (use_output_calibration)
        {
            args.push_back({ ArgumentDescriptor::Types::OUTPUT_CALIBRATION_FACTORS, 0 });
        }

        return args;
    }

    std::shared_ptr<KernelString> common_kernel_base::GetKernelString(const std::string& name, const std::string& jit, const std::string& entry_point, const EngineInfo& engine_info, const std::string& exe_mode) const
    {
        std::shared_ptr<KernelString> kernel_string = std::make_shared<KernelString>();

        auto codes = db.get(name);

        if (codes.size())
        {
            kernel_string->str = codes[0];
            kernel_string->jit = jit;
            kernel_string->options = exe_mode + " -cl-mad-enable";
            if (engine_info.bIMMADSupport)
                kernel_string->options += " -DMMAD_SUPPORTED=1";
            if (engine_info.bIMADSupport)
                kernel_string->options += " -DIMAD_SUPPORTED=1";
            kernel_string->entry_point = entry_point;
            kernel_string->batch_compilation = true;
        }

        return kernel_string;
    }

    static void Check_RunInfoData(const std::string &kernelName, const kernel_selector::CommonDispatchData &runInfo)
    {
        if (runInfo.lws0 * runInfo.lws1 * runInfo.lws2 > 256)
        {
            std::cout << "ERROR: dispatch data for kernel: " << kernelName <<  " LWS cannot be greater than 256!\n" << std::endl;
        }
        if (runInfo.gws0 == 0 || runInfo.gws1 == 0 || runInfo.gws2 == 0 || runInfo.lws0 == 0 || runInfo.lws1 == 0 || runInfo.lws2 == 0)
        {
            std::cout << "ERROR: dispatch data for kernel: " << kernelName << " dispatch data cannot contain zeros!" << std::endl;
        }
        if (runInfo.gws0 % runInfo.lws0 != 0)
        {
            std::cout << "ERROR: dispatch data for kernel: " << kernelName << " is incorrect: GWS0: " << runInfo.gws0 << " LWS0: " << runInfo.lws0 << std::endl;
        }
        if (runInfo.gws0 % runInfo.lws0 != 0)
        {
            std::cout << "ERROR: dispatch data for kernel: " << kernelName << " is incorrect: GWS1: " << runInfo.gws1 << " LWS1: " << runInfo.lws1 << std::endl;
        }
        if (runInfo.gws0 % runInfo.lws0 != 0)
        {
            std::cout << "ERROR: dispatch data for kernel: " << kernelName << " is incorrect: GWS2: " << runInfo.gws2 << " LWS2: " << runInfo.lws2 << std::endl;
        }
    }

   void common_kernel_base::FillCLKernelData(clKernelData& kernel, const CommonDispatchData& runInfo, const EngineInfo& engine_info,
        const std::string& kernelMapName, const std::string& jit, const std::string& entryPoint, const std::string& exeMode, bool weights, bool bias, int number_of_inputs, bool quantization, bool calibration) const
    {
        Check_RunInfoData(kernelMapName, runInfo);
        kernel.workGroups.global = { runInfo.gws0, runInfo.gws1, runInfo.gws2 };
        kernel.workGroups.local = { runInfo.lws0, runInfo.lws1, runInfo.lws2 };
        kernel.kernelString = GetKernelString(kernelMapName, jit, entryPoint, engine_info, exeMode);
        kernel.arguments = GetArgsDesc(number_of_inputs, weights, bias, quantization, calibration);
    }
}