/* file: low_order_moments_kernels_all.h */
/*******************************************************************************
* Copyright 2014-2021 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.
*******************************************************************************/

/*
//++
//  Implementation of low order moments kernels.
//--
*/

#ifndef __low_order_moments_kernels_all__
#define __low_order_moments_kernels_all__

static const char * low_order_moments_kernels_all_cl =
    "\n"
    "#define CONCAT(n, suff) n ## suff\n"
    "#define FULLNAME(n, p) CONCAT(n, p)\n"
    "\n"
    "#define singlePassBlockProcessor FULLNAME(singlePassBlockProcessor, FNAMESUFF)\n"
    "#define singlePass FULLNAME(singlePass, FNAMESUFF)\n"
    "#define blockProcessor FULLNAME(blockProcessor, FNAMESUFF)\n"
    "#define processBlocks FULLNAME(processBlocks, FNAMESUFF)\n"
    "#define mergeBlocks FULLNAME(mergeBlocks, FNAMESUFF)\n"
    "#define finalize FULLNAME(finalize, FNAMESUFF)\n"
    "\n"
    "/* single pass kernels common */\n"
    "void singlePassBlockProcessor(__global const algorithmFPType* vectors,\n"
    "                              const uint nVectors,\n"
    "                              const uint vectorSize\n"
    "                #if (defined _ONLINE_)\n"
    "                              ,const algorithmFPType nObservations\n"
    "                #endif\n"
    "                      #if (defined _RMIN_)\n"
    "                              ,__global algorithmFPType* gMin\n"
    "                      #endif\n"
    "                      #if (defined _RMAX_)\n"
    "                              ,__global algorithmFPType* gMax\n"
    "                      #endif\n"
    "                      #if (defined _RSUM_) || (defined _ONLINE_) && ((defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined "
    "_RSTDEV_) || (defined _RVART_))\n"
    "                              ,__global algorithmFPType* gSum\n"
    "                      #endif\n"
    "                      #if (defined _RSUM2_) || (defined _ONLINE_) && (defined _RSORM_)\n"
    "                              ,__global algorithmFPType* gSum2\n"
    "                      #endif\n"
    "                      #if (defined _RSUM2C_) || (defined _ONLINE_) && ((defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_))\n"
    "                              ,__global algorithmFPType* gSum2Cent\n"
    "                      #endif\n"
    "                #if !(defined _ONLINE_)\n"
    "                      #if (defined _RMEAN_)\n"
    "                              ,__global algorithmFPType* gMean\n"
    "                      #endif\n"
    "                      #if (defined _RSORM_)\n"
    "                              ,__global algorithmFPType* gSecondOrderRawMoment\n"
    "                      #endif\n"
    "                      #if (defined _RVARC_)\n"
    "                              ,__global algorithmFPType* gVariance\n"
    "                      #endif\n"
    "                      #if (defined _RSTDEV_)\n"
    "                              ,__global algorithmFPType* gStDev\n"
    "                      #endif\n"
    "                      #if (defined _RVART_)\n"
    "                              ,__global algorithmFPType* gVariation\n"
    "                      #endif\n"
    "                #endif\n"
    "                              ,const uint rowPartIndex,\n"
    "                              const uint rowParts,\n"
    "                              const uint colPartIndex,\n"
    "                              const uint colParts,\n"
    "                              const uint tid,\n"
    "                              const uint tnum)\n"
    "{\n"
    "    const uint colOffset = colPartIndex * tnum;\n"
    "    const uint x = tid + colOffset;\n"
    "\n"
    "    if (x < nVectors)\n"
    "    {\n"
    "        uint rowPartSize = (vectorSize + rowParts - 1) / rowParts;\n"
    "        const uint rowOffset = rowPartSize * rowPartIndex;\n"
    "\n"
    "        if (rowPartSize + rowOffset > vectorSize)\n"
    "        {\n"
    "            rowPartSize = vectorSize - rowOffset;\n"
    "        }\n"
    "\n"
    "#if (defined _ONLINE_)\n"
    "    // for online mode initial values of min/max are defined later depending on nObservations \n"
    "    #if (defined _RMIN_)\n"
    "            algorithmFPType min      = (algorithmFPType)0;\n"
    "    #endif\n"
    "    #if (defined _RMAX_)\n"
    "            algorithmFPType max      = (algorithmFPType)0;\n"
    "    #endif\n"
    "#else\n"
    "    #if (defined _RMIN_)\n"
    "            algorithmFPType min      = vectors[rowOffset * nVectors + x];\n"
    "    #endif\n"
    "    #if (defined _RMAX_)\n"
    "            algorithmFPType max      = vectors[rowOffset * nVectors + x];\n"
    "    #endif\n"
    "#endif\n"
    "\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        algorithmFPType sum      = (algorithmFPType)0;  \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "        algorithmFPType sum2     = (algorithmFPType)0;  \n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        algorithmFPType sum2Cent = (algorithmFPType)0;  \n"
    "        algorithmFPType mean     = (algorithmFPType)0;  \n"
    "#endif\n"
    "\n"
    "#if (defined _ONLINE_)\n"
    "        if((algorithmFPType)0 == nObservations)\n"
    "        {\n"
    "    #if (defined _RMIN_)\n"
    "            min      = vectors[rowOffset * nVectors + x];\n"
    "    #endif\n"
    "    #if (defined _RMAX_)\n"
    "            max      = vectors[rowOffset * nVectors + x];\n"
    "    #endif\n"
    "    #if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            sum      = (algorithmFPType)0;  \n"
    "    #endif\n"
    "    #if (defined _RSUM2_) || (defined _RSORM_)\n"
    "            sum2     = (algorithmFPType)0;  \n"
    "    #endif\n"
    "    #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            sum2Cent = (algorithmFPType)0;  \n"
    "            mean     = (algorithmFPType)0;  \n"
    "    #endif\n"
    "        }\n"
    "        else\n"
    "        {\n"
    "    #if (defined _RMIN_)\n"
    "            min      = gMin [x * rowParts + rowPartIndex];\n"
    "    #endif\n"
    "    #if (defined _RMAX_)\n"
    "            max      = gMax [x * rowParts + rowPartIndex];\n"
    "    #endif\n"
    "    #if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            sum      = gSum [x * rowParts + rowPartIndex];  \n"
    "    #endif\n"
    "    #if (defined _RSUM2_) || (defined _RSORM_)\n"
    "            sum2     = gSum2[x * rowParts + rowPartIndex];  \n"
    "    #endif\n"
    "    #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            sum2Cent = gSum2Cent[x * rowParts + rowPartIndex];  \n"
    "            mean     = sum/nObservations;  \n"
    "    #endif\n"
    "        }\n"
    "#endif \n"
    "\n"
    "        for (int row = 0; row < rowPartSize; row++)\n"
    "        {\n"
    "            const uint y = (row + rowOffset) * nVectors;\n"
    "            const algorithmFPType el = vectors[y + x];\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "    #if (defined _ONLINE_)\n"
    "            algorithmFPType invN   = ((algorithmFPType)1) / (nObservations + (algorithmFPType)(row + 1));\n"
    "    #else            \n"
    "            algorithmFPType invN   = ((algorithmFPType)1) / (algorithmFPType)(row + 1);\n"
    "    #endif\n"
    "            algorithmFPType delta  = el - mean;\n"
    "#endif\n"
    "\n"
    "#if (defined _RMIN_)\n"
    "            min       = fmin(el, min);\n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "            max       = fmax(el, max);\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            sum      += el; \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "            sum2     += el * el; \n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            mean     += delta  * invN;\n"
    "            sum2Cent += delta  * (el - mean);\n"
    "#endif\n"
    "        }\n"
    "\n"
    "#if (defined _RMIN_)\n"
    "        gMin [x * rowParts + rowPartIndex] = min;\n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "        gMax [x * rowParts + rowPartIndex] = max;\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _ONLINE_) && ((defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined "
    "_RVART_))\n"
    "        gSum [x * rowParts + rowPartIndex] = sum;  \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _ONLINE_) && (defined _RSORM_)\n"
    "        gSum2[x * rowParts + rowPartIndex] = sum2;\n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _ONLINE_) && ((defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_))\n"
    "        gSum2Cent[x * rowParts + rowPartIndex] = sum2Cent;\n"
    "#endif\n"
    "\n"
    "#if !(defined _ONLINE_)\n"
    "    // common vars calculation\n"
    "    #if (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        algorithmFPType variance = sum2Cent / (rowPartSize - (algorithmFPType)1);\n"
    "    #endif\n"
    "    #if (defined _RSTDEV_) || (defined _RVART_)\n"
    "        algorithmFPType stDev    = (algorithmFPType)sqrt(variance);\n"
    "    #endif\n"
    "\n"
    "    // output assignment\n"
    "    #if (defined _RMEAN_)\n"
    "        gMean[x * rowParts + rowPartIndex] = mean;\n"
    "    #endif\n"
    "    #if (defined _RSORM_)\n"
    "        gSecondOrderRawMoment[x * rowParts + rowPartIndex] = sum2/rowPartSize;\n"
    "    #endif\n"
    "    #if (defined _RVARC_)\n"
    "        gVariance[x * rowParts + rowPartIndex]  = variance;\n"
    "    #endif\n"
    "    #if (defined _RSTDEV_)\n"
    "        gStDev[x * rowParts + rowPartIndex]     = stDev; \n"
    "    #endif\n"
    "    #if (defined _RVART_)\n"
    "        gVariation[x * rowParts + rowPartIndex] = stDev/mean;\n"
    "    #endif\n"
    "#endif\n"
    "    }\n"
    "}\n"
    "\n"
    "__kernel void singlePass(__global const algorithmFPType* vectors,\n"
    "                            const uint nVectors,\n"
    "                            const uint vectorSize\n"
    "                #if (defined _ONLINE_)\n"
    "                              ,const algorithmFPType nObservations\n"
    "                #endif\n"
    "                      #if (defined _RMIN_)\n"
    "                              ,__global algorithmFPType* gMin\n"
    "                      #endif\n"
    "                      #if (defined _RMAX_)\n"
    "                              ,__global algorithmFPType* gMax\n"
    "                      #endif\n"
    "                      #if (defined _RSUM_) || (defined _ONLINE_) && ((defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined "
    "_RSTDEV_) || (defined _RVART_))\n"
    "                              ,__global algorithmFPType* gSum\n"
    "                      #endif\n"
    "                      #if (defined _RSUM2_) || (defined _ONLINE_) && (defined _RSORM_)\n"
    "                              ,__global algorithmFPType* gSum2\n"
    "                      #endif\n"
    "                      #if (defined _RSUM2C_) || (defined _ONLINE_) && ((defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_))\n"
    "                              ,__global algorithmFPType* gSum2Cent\n"
    "                      #endif\n"
    "                #if !(defined _ONLINE_)\n"
    "                      #if (defined _RMEAN_)\n"
    "                              ,__global algorithmFPType* gMean\n"
    "                      #endif\n"
    "                      #if (defined _RSORM_)\n"
    "                              ,__global algorithmFPType* gSecondOrderRawMoment\n"
    "                      #endif\n"
    "                      #if (defined _RVARC_)\n"
    "                              ,__global algorithmFPType* gVariance\n"
    "                      #endif\n"
    "                      #if (defined _RSTDEV_)\n"
    "                              ,__global algorithmFPType* gStDev\n"
    "                      #endif\n"
    "                      #if (defined _RVART_)\n"
    "                              ,__global algorithmFPType* gVariation\n"
    "                      #endif\n"
    "                #endif\n"
    "                            )\n"
    "{\n"
    "    const uint tid  = get_local_id(0);\n"
    "    const uint tnum = get_local_size(0);\n"
    "    const uint gid  = get_group_id(0);\n"
    "    const uint gnum = get_num_groups(0);\n"
    "\n"
    "    const uint colParts = (nVectors + tnum - 1) / tnum;\n"
    "    const uint rowParts = gnum / colParts;\n"
    "\n"
    "    const uint rowPartIndex = gid / colParts;\n"
    "    const uint colPartIndex = gid - rowPartIndex * colParts;\n"
    "\n"
    "    singlePassBlockProcessor(vectors, nVectors, vectorSize \n"
    "                #if (defined _ONLINE_)\n"
    "                             ,nObservations\n"
    "                #endif\n"
    "                      #if (defined _RMIN_)\n"
    "                             ,gMin\n"
    "                      #endif\n"
    "                      #if (defined _RMAX_)\n"
    "                             ,gMax\n"
    "                      #endif\n"
    "                      #if (defined _RSUM_) || (defined _ONLINE_) && ((defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined "
    "_RSTDEV_) || (defined _RVART_))\n"
    "                             ,gSum\n"
    "                      #endif\n"
    "                      #if (defined _RSUM2_) || (defined _ONLINE_) && (defined _RSORM_)\n"
    "                             ,gSum2\n"
    "                      #endif\n"
    "                      #if (defined _RSUM2C_) || (defined _ONLINE_) && ((defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_))\n"
    "                             ,gSum2Cent\n"
    "                      #endif\n"
    "                #if !(defined _ONLINE_)\n"
    "                      #if (defined _RMEAN_)\n"
    "                             ,gMean\n"
    "                      #endif\n"
    "                      #if (defined _RSORM_)\n"
    "                             ,gSecondOrderRawMoment\n"
    "                      #endif\n"
    "                      #if (defined _RVARC_)\n"
    "                             ,gVariance\n"
    "                      #endif\n"
    "                      #if (defined _RSTDEV_)\n"
    "                             ,gStDev\n"
    "                      #endif\n"
    "                      #if (defined _RVART_)\n"
    "                             ,gVariation\n"
    "                      #endif\n"
    "                #endif\n"
    "                             ,rowPartIndex, rowParts,\n"
    "                             colPartIndex, colParts,\n"
    "                             tid, tnum);\n"
    "}\n"
    "\n"
    "/* common kernels for blocks processing */\n"
    "\n"
    "void blockProcessor(__global const algorithmFPType* vectors,\n"
    "                    const uint nVectors,\n"
    "                    const uint vectorSize\n"
    "            #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                    ,__global uint* bNVec\n"
    "            #endif\n"
    "            #if (defined _RMIN_)\n"
    "                    ,__global algorithmFPType* bMin\n"
    "            #endif\n"
    "            #if (defined _RMAX_)\n"
    "                    ,__global algorithmFPType* bMax\n"
    "            #endif\n"
    "            #if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                    ,__global algorithmFPType* bSum\n"
    "            #endif\n"
    "            #if (defined _RSUM2_) || (defined _RSORM_)\n"
    "                    ,__global algorithmFPType* bSum2\n"
    "            #endif\n"
    "            #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                    ,__global algorithmFPType* bSum2Cent\n"
    "            #endif\n"
    "                    ,const uint rowPartIndex,\n"
    "                    const uint rowParts,\n"
    "                    const uint colPartIndex,\n"
    "                    const uint colParts,\n"
    "                    const uint tid,\n"
    "                    const uint tnum)\n"
    "{\n"
    "    const uint colOffset = colPartIndex * tnum;\n"
    "    const uint x = tid + colOffset;\n"
    "\n"
    "    if (x < nVectors)\n"
    "    {\n"
    "        uint rowPartSize     = (vectorSize + rowParts - 1) / rowParts;\n"
    "        const uint rowOffset = rowPartSize * rowPartIndex;\n"
    "\n"
    "        if (rowPartSize + rowOffset > vectorSize)\n"
    "        {\n"
    "            rowPartSize = vectorSize - rowOffset;\n"
    "        }\n"
    "\n"
    "#if (defined _RMIN_)\n"
    "        algorithmFPType min      = vectors[rowOffset * nVectors + x];\n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "        algorithmFPType max      = vectors[rowOffset * nVectors + x];\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        algorithmFPType sum      = (algorithmFPType)0;  \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "        algorithmFPType sum2     = (algorithmFPType)0;  \n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        algorithmFPType sum2Cent = (algorithmFPType)0;  \n"
    "        algorithmFPType mean     = (algorithmFPType)0;  \n"
    "#endif\n"
    "\n"
    "        for (int row = 0; row < rowPartSize; row++)\n"
    "        {\n"
    "            const uint              y = (row + rowOffset) * nVectors;\n"
    "            const algorithmFPType el = vectors[y + x];\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            algorithmFPType invN   = ((algorithmFPType)1) / (algorithmFPType)(row + 1);\n"
    "            algorithmFPType delta  = el - mean;\n"
    "#endif\n"
    "\n"
    "#if (defined _RMIN_)\n"
    "            min       = fmin(el, min);\n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "            max       = fmax(el, max);\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            sum      += el; \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "            sum2     += el * el; \n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            mean     += delta  * invN;\n"
    "            sum2Cent += delta  * (el - mean);\n"
    "#endif\n"
    "        }\n"
    "\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        bNVec[x * rowParts + rowPartIndex] = (uint)rowPartSize;\n"
    "#endif\n"
    "#if (defined _RMIN_)\n"
    "        bMin [x * rowParts + rowPartIndex] = min;\n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "        bMax [x * rowParts + rowPartIndex] = max;\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        bSum [x * rowParts + rowPartIndex] = sum;  \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "        bSum2[x * rowParts + rowPartIndex] = sum2;\n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        bSum2Cent[x * rowParts + rowPartIndex] = sum2Cent;\n"
    "#endif\n"
    "    }\n"
    "}\n"
    "\n"
    "__kernel void processBlocks(__global const algorithmFPType* vectors,\n"
    "                            const    uint             nVectors,\n"
    "                            const    uint             vectorSize\n"
    "                    #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                            ,__global uint* bNVec\n"
    "                    #endif\n"
    "                    #if (defined _RMIN_)\n"
    "                            ,__global algorithmFPType* bMin\n"
    "                    #endif\n"
    "                    #if (defined _RMAX_)\n"
    "                            ,__global algorithmFPType* bMax\n"
    "                    #endif\n"
    "                    #if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined "
    "_RVART_)\n"
    "                            ,__global algorithmFPType* bSum\n"
    "                    #endif\n"
    "                    #if (defined _RSUM2_) || (defined _RSORM_)\n"
    "                            ,__global algorithmFPType* bSum2\n"
    "                    #endif\n"
    "                    #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                            ,__global algorithmFPType* bSum2Cent\n"
    "                    #endif\n"
    "                            )\n"
    "{\n"
    "    const uint tid  = get_local_id(0);\n"
    "    const uint tnum = get_local_size(0);\n"
    "    const uint gid  = get_group_id(0);\n"
    "    const uint gnum = get_num_groups(0);\n"
    "\n"
    "    const uint colParts = (nVectors + tnum - 1) / tnum;\n"
    "    const uint rowParts = gnum / colParts;\n"
    "\n"
    "    const uint rowPartIndex = gid / colParts;\n"
    "    const uint colPartIndex = gid - rowPartIndex * colParts;\n"
    "\n"
    "    blockProcessor(vectors, nVectors, vectorSize\n"
    "           #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                   ,bNVec\n"
    "           #endif\n"
    "           #if (defined _RMIN_)\n"
    "                   ,bMin\n"
    "           #endif\n"
    "           #if (defined _RMAX_)\n"
    "                   ,bMax\n"
    "           #endif\n"
    "           #if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                   ,bSum\n"
    "           #endif\n"
    "           #if (defined _RSUM2_) || (defined _RSORM_)\n"
    "                   ,bSum2\n"
    "           #endif\n"
    "           #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                   ,bSum2Cent\n"
    "           #endif\n"
    "                   ,rowPartIndex, rowParts,\n"
    "                   colPartIndex, colParts,\n"
    "                   tid, tnum);\n"
    "}\n"
    "\n"
    "/* merge blocks kernel */\n"
    "__kernel void mergeBlocks(const    uint             vectorSize\n"
    "                #if (defined _ONLINE_)\n"
    "                              ,const algorithmFPType nObservations\n"
    "                #endif\n"
    "                    #if (defined _RMIN_)\n"
    "                             ,__global algorithmFPType* gMin\n"
    "                    #endif\n"
    "                    #if (defined _RMAX_)\n"
    "                             ,__global algorithmFPType* gMax\n"
    "                    #endif\n"
    "                    #if (defined _RSUM_) || (defined _ONLINE_) && ((defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined "
    "_RSTDEV_) || (defined _RVART_))\n"
    "                             ,__global algorithmFPType* gSum\n"
    "                    #endif\n"
    "                    #if (defined _RSUM2_) || (defined _ONLINE_) && (defined _RSORM_)\n"
    "                             ,__global algorithmFPType* gSum2\n"
    "                    #endif\n"
    "                    #if (defined _RSUM2C_) || (defined _ONLINE_) && ((defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_))\n"
    "                             ,__global algorithmFPType* gSum2Cent\n"
    "                    #endif\n"
    "                #if !(defined _ONLINE_)\n"
    "                      #if (defined _RMEAN_)\n"
    "                              ,__global algorithmFPType* gMean\n"
    "                      #endif\n"
    "                      #if (defined _RSORM_)\n"
    "                              ,__global algorithmFPType* gSecondOrderRawMoment\n"
    "                      #endif\n"
    "                      #if (defined _RVARC_)\n"
    "                              ,__global algorithmFPType* gVariance\n"
    "                      #endif\n"
    "                      #if (defined _RSTDEV_)\n"
    "                              ,__global algorithmFPType* gStDev\n"
    "                      #endif\n"
    "                      #if (defined _RVART_)\n"
    "                              ,__global algorithmFPType* gVariation\n"
    "                      #endif\n"
    "                #endif\n"
    "                    #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                             ,__global uint* bNVec\n"
    "                    #endif\n"
    "                    #if (defined _RMIN_)\n"
    "                             ,__global algorithmFPType* bMin\n"
    "                    #endif\n"
    "                    #if (defined _RMAX_)\n"
    "                             ,__global algorithmFPType* bMax\n"
    "                    #endif\n"
    "                    #if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined "
    "_RVART_)\n"
    "                             ,__global algorithmFPType* bSum\n"
    "                    #endif\n"
    "                    #if (defined _RSUM2_) || (defined _RSORM_)\n"
    "                             ,__global algorithmFPType* bSum2\n"
    "                    #endif\n"
    "                    #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                             ,__global algorithmFPType* bSum2Cent\n"
    "                    #endif\n"
    "                            )\n"
    "{\n"
    "#if (defined _RMIN_)\n"
    "    __local algorithmFPType lMin[LOCAL_BUFFER_SIZE];\n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "    __local algorithmFPType lMax[LOCAL_BUFFER_SIZE];\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "    __local algorithmFPType lSum[LOCAL_BUFFER_SIZE];\n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "    __local algorithmFPType lSum2[LOCAL_BUFFER_SIZE];\n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "    __local uint            lNVec[LOCAL_BUFFER_SIZE];\n"
    "    __local algorithmFPType lSum2Cent[LOCAL_BUFFER_SIZE];\n"
    "    __local algorithmFPType lMean[LOCAL_BUFFER_SIZE];\n"
    "#endif\n"
    "\n"
    "    const uint localSize = get_local_size(0);\n"
    "    const uint globalDim = vectorSize;\n"
    "    const uint localDim  = 1;\n"
    "    const uint itemId    = get_local_id(0);\n"
    "    const uint groupId   = get_group_id(0);\n"
    "\n"
    "#if (defined _RMIN_)\n"
    "    algorithmFPType mrgMin      = bMin[groupId*globalDim + itemId*localDim];\n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "    algorithmFPType mrgMax      = bMax[groupId*globalDim + itemId*localDim];\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "    algorithmFPType mrgSum      = (algorithmFPType)0;\n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "    algorithmFPType mrgSum2     = (algorithmFPType)0;\n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "    algorithmFPType mrgVectors  = (algorithmFPType)0;\n"
    "    algorithmFPType mrgSum2Cent = (algorithmFPType)0;\n"
    "    algorithmFPType mrgMean     = (algorithmFPType)0;\n"
    "#endif\n"
    "\n"
    "#if (defined _ONLINE_)\n"
    "    if(0 == itemId && (algorithmFPType)0 != nObservations)\n"
    "    {\n"
    "        // item 0 in each group performs merge of previous results\n"
    "    #if (defined _RMIN_)\n"
    "        mrgMin      = gMin[groupId];\n"
    "    #endif\n"
    "    #if (defined _RMAX_)\n"
    "        mrgMax      = gMax[groupId];\n"
    "    #endif\n"
    "    #if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        mrgSum      = gSum[groupId];\n"
    "    #endif\n"
    "    #if (defined _RSUM2_) || (defined _RSORM_)\n"
    "        mrgSum2     = gSum2[groupId];\n"
    "    #endif\n"
    "    #if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        mrgVectors  = nObservations;  \n"
    "        mrgSum2Cent = gSum2Cent[groupId];\n"
    "        mrgMean     = mrgSum/mrgVectors;\n"
    "    #endif\n"
    "    }\n"
    "#endif\n"
    "\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "    lNVec[itemId]     =  mrgVectors; \n"
    "#endif\n"
    "\n"
    "    for(uint i = itemId; i < vectorSize; i+= localSize)\n"
    "    {\n"
    "        uint offset = groupId*globalDim + i*localDim;\n"
    "\n"
    "#if (defined _RMIN_)\n"
    "        algorithmFPType min      = bMin[offset];  \n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "        algorithmFPType max      = bMax[offset];\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        algorithmFPType sum      = bSum[offset];  \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "        algorithmFPType sum2     = bSum2[offset];  \n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        uint            nVec     = bNVec[offset];\n"
    "        algorithmFPType sum2Cent = bSum2Cent[offset];  \n"
    "        algorithmFPType mean     = sum/(algorithmFPType)nVec;\n"
    "        \n"
    "        algorithmFPType sumN1N2        = mrgVectors  + (algorithmFPType)nVec;\n"
    "        algorithmFPType mulN1N2        = mrgVectors  * (algorithmFPType)nVec;\n"
    "        algorithmFPType deltaScale     = mulN1N2 / sumN1N2;\n"
    "        algorithmFPType meanScale      = (algorithmFPType)1 / sumN1N2;\n"
    "        algorithmFPType delta          = mean - mrgMean;\n"
    "#endif\n"
    "\n"
    "#if (defined _RMIN_)\n"
    "        mrgMin      = fmin(min, mrgMin);\n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "        mrgMax      = fmax(max, mrgMax);\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        mrgSum     += sum;  \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "        mrgSum2    += sum2;\n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        mrgSum2Cent = mrgSum2Cent + sum2Cent + delta*delta*deltaScale;\n"
    "        mrgMean     = (mrgMean * mrgVectors + mean * (algorithmFPType)nVec)* meanScale;\n"
    "        mrgVectors  = sumN1N2;\n"
    "#endif\n"
    "\n"
    "#if (defined _RMIN_)\n"
    "        lMin[itemId]      = mrgMin;     \n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "        lMax[itemId]      = mrgMax;     \n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        lSum[itemId]      = mrgSum;     \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "        lSum2[itemId]     = mrgSum2;    \n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        lNVec[itemId]    += nVec; \n"
    "        lSum2Cent[itemId] = mrgSum2Cent;\n"
    "        lMean[itemId]     = mrgMean;\n"
    "#endif\n"
    "    }\n"
    "\n"
    "    barrier(CLK_LOCAL_MEM_FENCE);\n"
    "\n"
    "    for (uint stride = localSize / 2; stride > 0; stride /= 2)\n"
    "    {\n"
    "        if (stride > itemId)\n"
    "        {\n"
    "            uint offset = itemId + stride;\n"
    "\n"
    "#if (defined _RMIN_)\n"
    "            algorithmFPType min      = lMin[offset];  \n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "            algorithmFPType max      = lMax[offset];\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            algorithmFPType sum      = lSum[offset];  \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "            algorithmFPType sum2     = lSum2[offset];  \n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            uint            nVec     = lNVec[offset];\n"
    "            algorithmFPType sum2Cent = lSum2Cent[offset];  \n"
    "            algorithmFPType mean     = lMean[offset];\n"
    "\n"
    "            algorithmFPType sumN1N2        = mrgVectors  + (algorithmFPType) nVec;\n"
    "            algorithmFPType mulN1N2        = mrgVectors  * (algorithmFPType) nVec;\n"
    "            algorithmFPType deltaScale     = mulN1N2 / sumN1N2;\n"
    "            algorithmFPType meanScale      = (algorithmFPType)1 / sumN1N2;\n"
    "            algorithmFPType delta          = mean - mrgMean;\n"
    "#endif\n"
    "\n"
    "#if (defined _RMIN_)\n"
    "            mrgMin      = fmin(min, mrgMin);\n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "            mrgMax      = fmax(max, mrgMax);\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            mrgSum     += sum;  \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "            mrgSum2    += sum2;\n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "            mrgSum2Cent = mrgSum2Cent + sum2Cent + delta*delta*deltaScale;\n"
    "            mrgMean     = (mrgMean * mrgVectors + mean * (algorithmFPType)nVec)* meanScale;\n"
    "            mrgVectors  = sumN1N2;\n"
    "#endif\n"
    "\n"
    "            // item 0 collects all results in private vars\n"
    "            // but all others need to store it\n"
    "            if(0 < itemId)\n"
    "            {\n"
    "#if (defined _RMIN_)\n"
    "                lMin[itemId]      = mrgMin;     \n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "                lMax[itemId]      = mrgMax;     \n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                lSum[itemId]      = mrgSum;     \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _RSORM_)\n"
    "                lSum2[itemId]     = mrgSum2;    \n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                lNVec[itemId]    += nVec; \n"
    "                lSum2Cent[itemId] = mrgSum2Cent;\n"
    "                lMean[itemId]     = mrgMean;\n"
    "#endif\n"
    "            }\n"
    "        }\n"
    "        barrier(CLK_LOCAL_MEM_FENCE);\n"
    "    }\n"
    "     \n"
    "    if (0 == itemId)\n"
    "    {\n"
    "#if (defined _RMIN_)\n"
    "        gMin[groupId]       = mrgMin;\n"
    "#endif\n"
    "#if (defined _RMAX_)\n"
    "        gMax[groupId]       = mrgMax;\n"
    "#endif\n"
    "#if (defined _RSUM_) || (defined _ONLINE_) && ((defined _RMEAN_) || (defined _RSUM2C_) || (defined _RVARC_) || (defined _RSTDEV_) || (defined "
    "_RVART_))\n"
    "        gSum[groupId]       = mrgSum;  \n"
    "#endif\n"
    "#if (defined _RSUM2_) || (defined _ONLINE_) && (defined _RSORM_)\n"
    "        gSum2[groupId]      = mrgSum2;\n"
    "#endif\n"
    "#if (defined _RSUM2C_) || (defined _ONLINE_) && ((defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_))\n"
    "        gSum2Cent[groupId]  = mrgSum2Cent;\n"
    "#endif\n"
    "\n"
    "#if !(defined _ONLINE_)\n"
    "    // common vars calculation\n"
    "    #if (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "        algorithmFPType mrgVariance         = mrgSum2Cent / (mrgVectors - (algorithmFPType)1);\n"
    "    #endif\n"
    "    #if (defined _RSTDEV_) || (defined _RVART_)\n"
    "        algorithmFPType mrgStDev            = (algorithmFPType)sqrt(mrgVariance);\n"
    "    #endif\n"
    "\n"
    "    // output assignment\n"
    "    #if (defined _RMEAN_)\n"
    "        gMean[groupId] = mrgMean;\n"
    "    #endif\n"
    "    #if (defined _RSORM_)\n"
    "        gSecondOrderRawMoment[groupId] = mrgSum2/mrgVectors;\n"
    "    #endif\n"
    "    #if (defined _RVARC_)\n"
    "        gVariance[groupId]  = mrgVariance;\n"
    "    #endif\n"
    "    #if (defined _RSTDEV_)\n"
    "        gStDev[groupId]     = mrgStDev; \n"
    "    #endif\n"
    "    #if (defined _RVART_)\n"
    "        gVariation[groupId] = mrgStDev/mrgMean;\n"
    "    #endif\n"
    "#endif\n"
    "    }\n"
    "}\n"
    "\n"
    "/* finalize kernel */\n"
    "\n"
    "__kernel void finalize(const algorithmFPType     nObservations\n"
    "                      #if (defined _RMIN_)\n"
    "                          ,__global algorithmFPType* gMin\n"
    "                      #endif\n"
    "                      #if (defined _RMAX_)\n"
    "                          ,__global algorithmFPType* gMax\n"
    "                      #endif\n"
    "                      #if (defined _RMEAN_)\n"
    "                          ,__global algorithmFPType* gSum\n"
    "                      #endif\n"
    "                      #if (defined _RSORM_)\n"
    "                          ,__global algorithmFPType* gSum2\n"
    "                      #endif\n"
    "                      #if (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "                          ,__global algorithmFPType* gSum2Cent\n"
    "                      #endif\n"
    "                      #if (defined _RMEAN_)\n"
    "                          ,__global algorithmFPType* gMean\n"
    "                      #endif\n"
    "                      #if (defined _RSORM_)\n"
    "                          ,__global algorithmFPType* gSecondOrderRawMoment\n"
    "                      #endif\n"
    "                      #if (defined _RVARC_)\n"
    "                          ,__global algorithmFPType* gVariance\n"
    "                      #endif\n"
    "                      #if (defined _RSTDEV_)\n"
    "                          ,__global algorithmFPType* gStDev\n"
    "                      #endif\n"
    "                      #if (defined _RVART_)\n"
    "                          ,__global algorithmFPType* gVariation\n"
    "                      #endif\n"
    "                         )\n"
    "{\n"
    "    const uint tid  = get_global_id(0);\n"
    "\n"
    "#if (defined _RMEAN_)\n"
    "    algorithmFPType sum      = gSum [tid];  \n"
    "#endif\n"
    "#if (defined _RSORM_)\n"
    "    algorithmFPType sum2     = gSum2[tid];  \n"
    "#endif\n"
    "#if (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "    algorithmFPType sum2Cent = gSum2Cent[tid];  \n"
    "#endif\n"
    "#if (defined _RMEAN_) || (defined _RVART_)\n"
    "    algorithmFPType mean     = sum / nObservations;\n"
    "#endif\n"
    "#if (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "    algorithmFPType variance = (algorithmFPType)0;  \n"
    "#endif\n"
    "#if (defined _RSTDEV_) || (defined _RVART_)\n"
    "    algorithmFPType stDev    = (algorithmFPType)0;  \n"
    "#endif\n"
    "\n"
    "// common vars calculation\n"
    "#if (defined _RVARC_) || (defined _RSTDEV_) || (defined _RVART_)\n"
    "    variance      = sum2Cent / (nObservations - (algorithmFPType)1);\n"
    "#endif\n"
    "#if (defined _RSTDEV_) || (defined _RVART_)\n"
    "    stDev         = (algorithmFPType)sqrt(variance);\n"
    "#endif\n"
    "\n"
    "// output assignment\n"
    "#if (defined _RMEAN_)\n"
    "    gMean[tid] = mean;\n"
    "#endif\n"
    "#if (defined _RSORM_)\n"
    "    gSecondOrderRawMoment[tid] = sum2 / nObservations;\n"
    "#endif\n"
    "#if (defined _RVARC_)\n"
    "    gVariance[tid]  = variance;\n"
    "#endif\n"
    "#if (defined _RSTDEV_)\n"
    "    gStDev[tid]     = stDev; \n"
    "#endif\n"
    "#if (defined _RVART_)\n"
    "    gVariation[tid] = stDev / mean;\n"
    "#endif\n"
    "}\n"
    "\n";
#endif
