/***************************************************************************************************
 * Copyright (c) 2017 - 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 * SPDX-License-Identifier: BSD-3-Clause
 *
 * Redistribution and use in source and binary forms, with or without
 * modification, are permitted provided that the following conditions are met:
 *
 * 1. Redistributions of source code must retain the above copyright notice, this
 * list of conditions and the following disclaimer.
 *
 * 2. Redistributions in binary form must reproduce the above copyright notice,
 * this list of conditions and the following disclaimer in the documentation
 * and/or other materials provided with the distribution.
 *
 * 3. Neither the name of the copyright holder nor the names of its
 * contributors may be used to endorse or promote products derived from
 * this software without specific prior written permission.
 *
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
 * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
 * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 * DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
 * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
 * SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
 * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
 * OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 *
 **************************************************************************************************/
/*! \file
    \brief Unit tests for thread-level GEMM
*/

#include "../../common/cutlass_unit_test.h"

#include "cutlass/layout/layout.h"
#include "cutlass/epilogue/thread/activation.h"

#include "cutlass/util/host_tensor.h"

/////////////////////////////////////////////////////////////////////////////////////////////////

template <typename T, int N, typename Func>
__global__ void test_Epilogue_thread_activation(T *out, T *in) {

  cutlass::Array<T, N> *vec_out = reinterpret_cast<cutlass::Array<T, N> *>(out);
  cutlass::Array<T, N> *vec_in = reinterpret_cast<cutlass::Array<T, N> *>(in);

  Func func;
  vec_out[threadIdx.x] = func(vec_in[threadIdx.x]);
}

template <typename T, int N, typename Func>
__global__ void test_Epilogue_thread_activation_binary(T *out, T *x, T *alpha){
    cutlass::Array<T, N> *vec_out = reinterpret_cast<cutlass::Array<T, N> *>(out);
    cutlass::Array<T, N> *vec_x = reinterpret_cast<cutlass::Array<T, N> *>(x);
    cutlass::Array<T, N> *vec_alpha = reinterpret_cast<cutlass::Array<T, N> *>(alpha);

    Func func;
    vec_out[threadIdx.x] = func(vec_x[threadIdx.x], vec_alpha[threadIdx.x]);
}

/////////////////////////////////////////////////////////////////////////////////////////////////

//
// Reference
//

static double GELU_golden_input[] = {
    1.587425827980,  1.157652974129,  0.750432848930, -0.965980410576,
    -0.388184845448,  0.014422321692,  0.353164494038,  1.354383468628,
     0.167588576674,  0.272798538208, -0.377032428980,  1.923444747925,
     0.308164477348, -0.341318070889,  0.278338819742, -0.292668998241,
    -1.051743745804, -0.814175724983,  0.112737402320,  1.262938618660,
    -1.582363605499,  0.722016870975,  1.053453564644, -0.659764587879,
     0.734917521477,  0.091274201870,  0.604461073875, -0.219043627381,
    -0.136795744300,  0.960650205612, -1.805408835411,  0.091029644012,
    -1.023343324661,  0.147713735700, -0.499895423651,  1.351878166199,
    -1.631091356277, -0.336171895266, -1.612408638000,  0.090832948685,
    -0.658132910728, -0.326727777719, -1.986387014389,  0.787685871124,
    -1.015677452087, -0.225094825029,  0.876752018929,  0.744826257229,
     0.870290279388, -0.757595360279,  1.510331749916,  0.750012576580,
     0.906444966793, -0.915759027004,  1.260277032852, -0.158465340734,
    -0.109191477299, -0.817102134228,  0.391305118799, -0.524910449982,
     0.351349592209,  0.801979541779,  0.446691334248, -0.741077482700,
     1.205966711044, -0.910210072994,  0.945986449718,  0.784096539021,
     1.670521497726,  0.344931513071, -0.301411420107,  0.309870749712,
    -0.879704594612, -1.951189517975, -0.805817663670, -0.661812782288,
    -0.505914270878, -1.836273789406, -0.381845980883, -0.554707705975,
    -0.375447630882, -0.516645610332,  0.509586095810,  1.087131023407,
     2.664817094803, -1.558295488358, -0.076461032033, -0.504621028900,
     1.327111959457, -1.819981694221,  1.350415468216, -2.074112653732,
     1.501431345940, -1.339013576508,  0.162817999721, -1.473457217216,
     0.357770472765,  0.188413277268,  1.601302266121, -0.653882205486,
     0.856162548065,  0.763102591038, -0.526283502579,  0.581961452961,
     0.089969776571,  1.968745589256,  0.545802056789, -1.168786048889,
     1.206663012505, -0.109096683562, -1.223938226700,  0.744599223137,
    -1.779406785965,  0.766436159611, -0.579044401646, -1.002057313919,
    -0.715845823288, -0.562508940697,  0.886768460274,  2.327786445618,
    -0.148763969541, -0.918884515762, -0.367678701878, -1.105021238327,
    -0.461237311363,  0.158228352666, -0.254040330648,  1.427477598190,
     0.277530491352,  0.046293262392, -0.535557329655, -1.486695051193,
    -0.953706681728, -1.040495038033, -0.314667612314,  0.348172843456,
     0.522773325443,  0.025960063562, -0.482472360134,  1.993084549904,
    -0.253064930439, -0.012146313675, -2.166327714920,  0.398040622473,
    -0.022238900885, -0.443580865860, -0.898376941681, -0.571689844131,
     1.666979670525, -0.831176340580, -0.671057403088,  0.481970995665,
    -1.096243023872, -1.493894338608,  0.596651911736, -0.229505166411,
     1.165976166725,  0.905094027519,  0.049716457725, -1.362933635712,
    -0.366948783398,  1.461613893509, -0.718411505222,  0.895385026932,
    -0.763122260571,  1.329716682434,  1.366570711136, -0.086544901133,
     0.059739742428,  0.940766513348, -0.272854357958, -1.738811373711,
    -0.361239165068,  0.696977972984,  1.288442254066,  1.264815807343,
    -0.573566436768, -1.141678214073,  0.081865988672, -0.886228799820,
    -0.236933603883,  1.050115466118, -0.538952171803,  0.651773929596,
    -0.220034509897, -1.198960781097,  1.247478365898, -0.053529661149,
     0.639809548855,  1.672434806824,  0.511088073254, -1.179364681244,
    -0.730427742004,  0.157630980015,  0.389369845390, -0.925578773022,
    -0.093250080943, -0.391062080860,  0.852983593941,  1.868778109550,
    -1.198786258698,  0.604997038841, -1.482687234879, -2.469333171844,
     0.718807697296, -0.559609353542,  2.187228441238, -2.927527904510,
     0.148535788059, -0.097280368209,  0.674131810665, -1.137645959854,
     0.792729616165, -1.166317462921, -0.498791724443,  1.675866723061,
    -0.137909621000, -0.653263568878, -2.281216144562,  0.296096831560,
     2.002410173416,  1.083609819412,  0.933580815792, -1.504760265350,
     2.185185909271,  0.286121010780, -1.035485863686, -0.216372340918,
    -0.274334043264, -0.849510788918, -1.397169828415, -0.407644748688,
     0.159476816654, -0.170650705695,  0.335193097591, -0.156852483749,
     0.036168430001,  0.858105242252, -1.086121797562,  0.404813349247,
    -0.481496721506, -0.389882832766,  0.020690204576, -0.772020936012,
    -0.758921504021,  0.323482036591,  0.115715265274, -0.811228036880,
    -0.882436633110,  0.176811277866,  1.678015947342,  0.379081040621,
    -0.842976212502,  0.346952259541, -0.545828759670,  1.632800459862
};

static double GELU_golden_output[] = {
    1.498199582100,  1.014679551125,  0.580462038517, -0.161344811320,
    -0.135453075171,  0.007294139825,  0.225325092673,  1.235459089279,
     0.094946734607,  0.165724009275, -0.133120641112,  1.871103763580,
     0.191376730800, -0.125069886446,  0.169681981206, -0.112644664943,
    -0.154036879539, -0.169163048267,  0.061428427696,  1.132469892502,
    -0.089851818979,  0.552240371704,  0.899579226971, -0.168043658137,
     0.565008401871,  0.048956073821,  0.439583092928, -0.090532489121,
    -0.060955654830,  0.798911273479, -0.064101703465,  0.048816055059,
    -0.156645998359,  0.082529976964, -0.154254898429,  1.232632875443,
    -0.083896033466, -0.123835846782, -0.086161509156,  0.048703473061,
    -0.167972877622, -0.121522113681, -0.046670529991,  0.617986679077,
    -0.157319813967, -0.092503339052,  0.709896743298,  0.574865520000,
     0.703132867813, -0.169963955879,  1.411436080933,  0.580042064190,
     0.741154611111, -0.164741978049,  1.129479527473, -0.069256491959,
    -0.049848672003, -0.169087052345,  0.255214750767, -0.157380074263,
     0.223928079009,  0.632535398006,  0.300378054380, -0.169946283102,
     1.068588852882, -0.165071934462,  0.783203184605,  0.614346146584,
     1.591325283051,  0.219006344676, -0.115003645420,  0.192637458444,
    -0.166712537408, -0.049788996577, -0.169361919165, -0.168130636215,
    -0.155041679740, -0.060888241976, -0.134137839079, -0.160614117980,
    -0.132782235742, -0.156389534473,  0.354075312614,  0.936574816704,
     2.654553413391, -0.092845752835, -0.035900454968, -0.154874503613,
     1.204704761505, -0.062572605908,  1.230982899666, -0.039479542524,
     1.401402950287, -0.120890334249,  0.091938301921, -0.103604510427,
     0.228880971670,  0.108285568655,  1.513783097267, -0.167782157660,
     0.688394129276,  0.593158841133, -0.157540664077,  0.418839782476,
     0.048209801316,  1.920528769493,  0.386099845171, -0.141709372401,
     1.069367766380, -0.049809500575, -0.135230198503,  0.574639260769,
    -0.066881760955,  0.596510827541, -0.162873372436, -0.158483341336,
    -0.169686436653, -0.161375194788,  0.720409095287,  2.304597616196,
    -0.065585561097, -0.164551988244, -0.131098195910, -0.148708447814,
    -0.148663327098,  0.089060656726, -0.101548098028,  1.317959904671,
     0.169103100896,  0.024001283571, -0.158595800400, -0.101909510791,
    -0.162240833044, -0.155090972781, -0.118474565446,  0.221488356590,
     0.365645468235,  0.013248858973, -0.151851043105,  1.946992278099,
    -0.101253561676, -0.006014300976, -0.032804865390,  0.260597169399,
    -0.010922161862, -0.145792976022, -0.165743649006, -0.162226170301,
     1.587365984917, -0.168676435947, -0.168497130275,  0.330191940069,
    -0.149622067809, -0.100989677012,  0.432351946831, -0.093922272325,
     1.023946166039,  0.739726305008,  0.025843897834, -0.117827951908,
    -0.130937814713,  1.356489539146, -0.169726014137,  0.729478538036,
    -0.169943705201,  1.207641005516,  1.249209761620, -0.040288090706,
     0.031292784959,  0.777626037598, -0.107090584934, -0.071350336075,
    -0.129670530558,  0.527676224709,  1.161149263382,  1.134579420090,
    -0.162394225597, -0.144757837057,  0.043603736907, -0.166386902332,
    -0.096278958023,  0.895924389362, -0.158969298005,  0.484089732170,
    -0.090857118368, -0.138206124306,  1.115107178688, -0.025622237474,
     0.472724437714,  1.593463659286,  0.355387806892, -0.140493586659,
    -0.169871479273,  0.088687323034,  0.253673940897, -0.164135158062,
    -0.043161027133, -0.136040985584,  0.685087263584,  1.811169505119,
    -0.138226687908,  0.440080583096, -0.102422207594, -0.016713079065,
     0.549075841904, -0.161096408963,  2.155813455582, -0.005001218989,
     0.083037458360, -0.044870752841,  0.505522191525, -0.145202502608,
     0.623111069202, -0.141991063952, -0.154108211398,  1.597298502922,
    -0.061391282827, -0.167753636837, -0.025704355910,  0.182520583272,
     1.957115054131,  0.932696640491,  0.769961357117, -0.099604383111,
     2.153636932373,  0.175279796124, -0.155551761389, -0.089653611183,
    -0.107515335083, -0.168032020330, -0.113423995674, -0.139319628477,
     0.089841812849, -0.073763631284,  0.211594089866, -0.068651281297,
     0.018605981022,  0.690416753292, -0.150658726692,  0.266040354967,
    -0.151710823178, -0.135800719261,  0.010515870526, -0.169883996248,
    -0.169960290194,  0.202769815922,  0.063187584281, -0.169236257672,
    -0.166577890515,  0.100812792778,  1.599699616432,  0.245525524020,
    -0.168275654316,  0.220552831888, -0.159705042839,  1.549110531807
};

static double Snake_golden_x[] = {
     0.029053429184,  -0.073700162105,  -0.957751321690,  -0.229464130144,
    -1.054150496919,  -0.438869051576,   0.058027082664,  -0.003119327765,
     0.432405505892,   0.530008876465,  -0.193385668869,  -0.186124514126,
     0.987970150471,  -0.990709745322,   0.417275110485,   0.806415034595,
     0.555878890939,   1.091370569098,  -0.652624600678,  -0.305256813078,
    -3.677834378094,   0.020385454985,   1.029275814765,  -1.658167845616,
     1.657781136611,  -0.627017252985,  -0.956601541851,   1.061830856399,
    -1.996895384427,  -0.796726103109,  -0.620245955882,  -0.216337758007,
    -0.970205843319,  -0.762363183314,  -0.263468892540,   0.706260146861,
    -0.547733438907,   1.028736714398,   0.532700417479,   0.219131887583,
     0.410371090789,  -0.860321949693,  -0.134011925815,   0.378068516506,
    -0.397935138621,  -0.281645814281,  -0.674965764600,   0.075066377376,
     0.591349731308,  -0.986103386873,   0.540573540728,  -0.195186025501,
     0.382170818241,  -0.336124810010,   0.514623317515,  -0.633037271176,
    -2.225114959496,  -0.424278574433,   0.864951591244,  -0.748100501207,
    -0.814455285078,   1.399324223175,   2.550507919197,   0.112022056199,
     0.501069610182,   0.615049456991,  -2.455269089844,  -0.133235353840,
    -1.048553628708,   0.614094770306,   1.135570854607,   2.135839794654,
    -1.132154443249,   0.416582342787,   0.199975770530,   0.870957070562,
    -1.067125609377,  -0.317593855414,  -1.365204049444,   0.581168074356,
    -1.034777548687,  -1.060219263238,  -0.952406555466,  -0.003488014184,
     1.072586326741,  -0.397457304592,  -0.479034742410,   1.616406066530,
    -0.340577642369,   1.352103052797,  -0.691795383481,  -0.623777129584,
    -0.774545149758,   1.277156048075,   0.098359200811,  -0.686292493575,
    -0.469269332700,  -1.084840021288,  -0.732852705424,   0.869019317992,
    -0.073697571167,  -1.893281354625,  -1.110974735000,   0.558214090747,
    -0.451057560756,   0.447461089338,   0.197494571084,   1.828269271566,
     0.827678321070,   0.798278975336,  -0.586337889400,   0.344753541998,
     0.815051610629,  -0.614818900479,   1.582063437842,   0.651619160959,
     1.581119442075,   1.303423829832,   1.186057802117,  -1.537848796338,
     0.439386823326,  -1.989887737521,  -0.530844704054,  -0.092983961880,
     2.731809826905,  -0.818315266683,   0.307108151906,   0.419630822952,
    -0.228842566535,  -0.215408664510,  -0.288090633921,  -1.365735055710,
    -1.513522512402,   1.790313840521,   0.365196032814,  -0.428147708497,
     0.944347054741,   1.889699917802,  -1.242820812714,  -0.295221915419,
     0.078524208465,   1.054135848649,   1.027491739965,  -0.638342920129,
     1.076303143523,  -0.431866264152,   0.582247177956,  -1.448992543035,
    -0.756821186630,  -0.148059611738,   0.413561474504,  -0.685984494800,
    -0.145935889932,   2.010263640011,   0.703426641971,  -2.445829502306,
    -1.085952602443,  -1.030567244448,  -0.778586822390,   1.409630762219,
     0.277775313817,  -0.817427176209,   1.535450399874,   0.689783169875,
    -0.558834477613,  -1.653817347161,   1.712497692950,   0.069523477922,
     0.629550163029,  -1.335994370169,   1.938423995860,   0.110381170776,
    -0.572298694654,  -0.510668568690,  -0.182994362066,  -1.180016882135,
     0.265853520200,   0.193906608187,   1.428012477429,   0.191521664853,
     1.540343614656,   0.351226457409,  -0.508023007883,   0.285515852982,
     2.007526227190,   0.513980525975,   1.671017405384,   1.609615930737,
     0.791780795594,  -1.086131492859,  -0.252483101516,  -0.456344955137,
    -0.457268634699,   0.032448480614,  -0.002629782119,   0.121148264944,
    -0.909985276560,  -0.615054078531,  -0.523624031453,   0.614206124249,
    -0.370579889227,   1.692268634801,  -0.032660060202,   1.486452809603,
    -0.111649745792,   0.493170630381,  -1.490800986310,  -0.468854631740,
     0.573790416548,  -2.243287177026,   0.778386044561,  -0.376802969344,
     1.936419879235,  -1.017792188928,   1.318757680373,   0.712829317167,
    -1.937630450484,   0.758523630020,  -0.624877514003,   0.228800672781,
     0.876842599715,  -0.973031770011,  -0.502147893064,  -0.839962712295,
     0.233819743409,   1.727011814368,   0.580076431287,  -0.219354728997,
    -0.644087881572,   0.259351024516,  -0.659000482282,  -0.717816416909,
     0.611838700073,   2.284100566159,   0.976363940000,  -0.168807743207,
     0.224294957866,   0.919491656813,   0.089184696718,  -0.742632491463,
     1.165372466386,  -1.195659887110,   0.368437873632,   1.075869087616,
    -0.070620875083,  -2.492959152949,   0.851985725585,   0.398781627880,
     0.710014437585,  -1.460378031581,  -1.532444654966,   2.996609973191,
     0.570902491622,   0.123891040827,  -0.820849107832,   0.255297073043,
};

static double Snake_golden_alpha[] = {
     1.314910917070,   0.147520434923,   0.622555704901,   0.524100402483,
     1.499295306912,   1.385729026104,   1.795141178639,   0.265183781996,
     0.901651457402,   0.156614716932,   0.515412152127,   1.060175047396,
     0.150418342399,   0.477791536305,   1.334780431781,   1.135388813146,
     0.518837181877,   1.219604799364,   1.637917867688,   0.112347643388,
     1.631056578482,   1.426464850478,   0.746475981384,   0.395411049642,
     1.918704837193,   0.739529635714,   0.276217102422,   0.283761115984,
     1.710239296060,   1.247079459597,   1.633543719221,   1.486490394718,
     1.118833373764,   1.948919951561,   0.819215316696,   1.148877199419,
     1.675868862081,   1.275187529492,   1.737243110590,   1.196969075988,
     1.438686488808,   0.187066328946,   0.533006723738,   0.649837130844,
     0.251604756155,   0.542302684086,   0.291902715878,   0.628149845909,
     1.307800444102,   0.793181140043,   0.803343837522,   0.498063358466,
     0.607257861893,   1.879643716654,   1.331267231969,   1.257348910767,
     0.425163431576,   1.485340916106,   0.410464738148,   0.820965339340,
     1.980094366210,   1.315999543723,   1.158204513172,   1.400767076881,
     1.701418648361,   1.574399831938,   0.535191336732,   0.160990463418,
     0.699360791312,   0.608707664354,   0.500867402814,   1.891528457237,
     1.765098490298,   0.697887973517,   1.345333464060,   0.851700612015,
     1.837640420507,   0.971818519916,   0.603272316346,   0.568592264619,
     1.166599454910,   0.599209056194,   1.210713381425,   1.805863478845,
     0.858860959767,   0.516709442399,   1.995321452341,   1.068099957985,
     0.272727883130,   0.189521113307,   0.308333347666,   1.292147479236,
     1.604950792290,   0.902103936919,   0.220702641689,   0.825076644362,
     1.992630622456,   1.105317255688,   1.945048917466,   1.735481434246,
     0.121813941691,   1.469371456784,   1.395249701150,   1.120243627777,
     0.606967860910,   1.317827417302,   0.311949129832,   0.926053976271,
     0.962075042025,   1.912250262290,   1.764120586719,   0.600439196427,
     1.051113614796,   0.439438573007,   1.833992894755,   1.753985282690,
     0.667045103752,   1.314004040245,   1.257043401733,   0.390394610244,
     1.548770520143,   1.124820157227,   1.579390309398,   1.107671977171,
     0.101086602643,   0.715896508309,   0.137005810533,   1.865287370903,
     1.769571567864,   1.680164505786,   0.684276838265,   0.210057816339,
     1.768218238488,   1.899203946066,   0.262741558929,   1.023381880302,
     0.231503785090,   1.545144113989,   1.555085415683,   0.343943782550,
     1.003036518388,   1.144626827640,   0.603607594986,   1.757622778062,
     0.903962086382,   0.502416590340,   1.124662568681,   1.486869031271,
     0.482187020440,   0.692260953472,   1.990783777656,   1.334768309515,
     0.932390159438,   1.083394097968,   0.329907972150,   0.526924940360,
     0.742362568080,   1.217786565069,   0.537217991933,   0.518413030458,
     0.234886863417,   1.299095618813,   0.534989389241,   1.820298024712,
     1.733307260482,   0.234628964788,   0.552208805301,   1.371057778763,
     0.507049934004,   0.351392512578,   1.877477057103,   1.184981877318,
     0.998074949992,   1.590776906152,   1.634244295757,   0.461778837288,
     0.284168547035,   0.918997246572,   0.904799383738,   0.987346869270,
     1.485244113974,   1.379392639857,   1.969913901595,   0.286993955189,
     0.864980435994,   0.744674950250,   1.737177819070,   0.572447034449,
     0.461396926038,   0.952365740883,   0.901575115685,   0.629235774867,
     0.574632250976,   1.854204638624,   0.941948415602,   1.736563299047,
     1.145618093655,   0.196117826097,   1.998636689984,   1.688452411652,
     1.941092888841,   1.860097267715,   1.712521895387,   0.415991110147,
     1.022718138356,   0.506119868478,   0.861976555844,   0.211407259947,
     0.820048926056,   1.972086803181,   0.603885810527,   1.589734143702,
     0.964515897944,   0.903714223381,   1.918903517633,   1.991303110036,
     1.155959814471,   1.464975723063,   0.394113968021,   0.663744868440,
     1.940547793441,   1.200442552551,   1.130170882611,   1.521153564720,
     0.208614018524,   1.209937429472,   1.055415727547,   1.720167794892,
     0.399122183085,   1.925479916221,   0.252211783957,   0.453067425863,
     1.230566702255,   1.382903851848,   0.546887400502,   0.327784566500,
     1.791545896846,   0.567809160798,   1.229586391714,   1.276824869631,
     0.896527339138,   1.208977349653,   1.093287159511,   1.875941889699,
     0.488092478905,   1.460764421500,   0.553503309970,   0.851993108903,
     1.376211423624,   0.669994451618,   0.700736672917,   1.528542535587,
     0.237831917537,   0.970742492975,   1.997063437623,   1.992583250925,
};

static double Snake_golden_output[] = {
     0.030162807937,  -0.072898904881,  -0.451235357962,  -0.202001042643,
    -0.387233066601,  -0.203293172314,   0.064049747723,  -0.003116747473,
     0.592620610942,   0.573902473192,  -0.174174024192,  -0.149871786124,
     1.133713590934,  -0.555749922155,   0.626624953928,   1.360172568431,
     0.711804070287,   1.865072258447,  -0.183375257350,  -0.294792169803,
    -3.629557790529,   0.020978079314,   1.676228747578,  -0.718106263419,
     1.658581601367,  -0.356520741146,  -0.709666913330,   1.372201772507,
    -1.954214035750,  -0.233634934140,  -0.179485640924,  -0.149132420195,
    -0.270902259126,  -0.252957737752,  -0.207479952329,   1.164095231296,
    -0.171221404601,   1.761510017151,   0.900069003977,   0.275302814603,
     0.625778121370,  -0.723055354236,  -0.124455822242,   0.469099634932,
    -0.358225882828,  -0.238961475856,  -0.543692958091,   0.078603354731,
     0.964472741676,  -0.359814582475,   0.760935446722,  -0.176270715394,
     0.469282660995,  -0.150556971280,   0.815370048789,  -0.226984684910,
    -0.677733423654,  -0.190473406374,   1.159349608758,  -0.343584035169,
    -0.310315044483,   2.104861670423,   2.580534178202,   0.129456397669,
     0.834279493899,   1.046249631450,  -0.707274523119,  -0.130377944208,
    -0.407949718281,   0.833151490387,   1.714738405775,   2.459414719751,
    -0.663025883113,   0.534320399265,   0.252490523020,   1.406886887967,
    -0.601681821943,  -0.222643876025,  -0.473145016659,   0.766324510773,
    -0.286018010333,  -0.472545537912,  -0.262369684875,  -0.003466043901,
     1.810937227027,  -0.316972658667,  -0.144678092813,   2.530137759850,
    -0.309033953907,   1.691064488623,  -0.546457367230,  -0.220875718256,
    -0.216018047603,   2.202461263004,   0.100494060288,  -0.337481553094,
    -0.144343811696,  -0.299469854864,  -0.229514049722,   1.442971110350,
    -0.073035976981,  -1.808978356092,  -0.394564481565,   0.864091753499,
    -0.330622849956,   0.682118910180,   0.209646481704,   2.892053970254,
     1.358707200042,   1.320198072576,  -0.167539638392,   0.415105207584,
     1.358353624039,  -0.452712667833,   1.612896990941,   1.123591800211,
     2.715150738504,   2.049233543250,   1.976511261696,  -0.720306639514,
     0.694995256047,  -1.441499342653,  -0.180764329912,  -0.083440827570,
     3.467215173315,  -0.391315325873,   0.320022288481,   0.686256722794,
    -0.141128137179,  -0.140792612702,  -0.232030129207,  -0.984560009377,
    -1.399641042245,   1.824741858878,   0.400129978571,  -0.252253096938,
     1.147531900195,   1.921001492668,  -0.680375923993,  -0.265348040001,
     0.084696204485,   1.816932450233,   1.587114829655,  -0.176514350448,
     1.832056765146,  -0.339622624154,   0.912045912589,  -0.980690637506,
    -0.492679505498,  -0.132937177582,   0.683745005253,  -0.214917977569,
    -0.126200761750,   2.632874230636,   0.863758613512,  -0.695016654463,
    -0.384453837757,  -0.288605904012,  -0.471480605518,   2.268948289718,
     0.295873275711,  -0.230375350268,   2.537377334201,   1.186352685088,
    -0.167005613002,  -1.043649922380,   2.903221727544,   0.076130458029,
     0.823777276176,  -0.753543706804,   2.059843094880,   0.124736844178,
    -0.279443963185,  -0.179416250318,  -0.129880671289,  -0.598191370944,
     0.285899830656,   0.228096484580,   2.449564888553,   0.227308464352,
     1.922903611519,   0.508484639786,  -0.148227747872,   0.308859084627,
     3.132178423292,   0.701284943286,   1.703208304556,   2.717753607115,
     1.068396770315,  -0.310346108239,  -0.195995603805,  -0.328868361179,
    -0.339856050311,   0.034398425285,  -0.002623267849,   0.146261819239,
    -0.258880864936,  -0.541223489008,  -0.148655781444,   1.053168999018,
    -0.146930666868,   1.692289244233,  -0.030835252178,   2.294266447178,
    -0.098956195752,   0.613732471310,  -0.422860765158,  -0.422534094080,
     0.824433000317,  -1.777238421851,   1.118104980062,  -0.176829371659,
     2.884464242244,  -0.317590030795,   1.490282214541,   1.203599102469,
    -1.405324950154,   1.306823510371,  -0.474073201204,   0.263281246252,
     1.383400285807,  -0.267965517976,  -0.246482570724,  -0.237437585109,
     0.245215978030,   2.350318084309,   0.893000796635,  -0.140440535337,
    -0.482127721991,   0.378450318066,  -0.550474425828,  -0.492483912395,
     0.991771804327,   2.284311962680,   1.450002778639,  -0.159476705205,
     0.309676599108,   1.357490057435,   0.098925576361,  -0.225763312326,
     1.999603147308,  -0.381426939450,   0.508992666191,   1.509134957604,
    -0.068187571496,  -2.335591589999,   1.224848727676,   0.529137461176,
     1.209259963395,  -0.433076347361,  -0.429737367104,   3.639504604020,
     0.647943774070,   0.138719258661,  -0.322459202822,   0.374343913903,
};



/////////////////////////////////////////////////////////////////////////////////////////////////

TEST(Epilogue_thread_snake, device_f32) {

    int const kN = 256;
    int const kV = 4;

    using Element = float;
    using Func = cutlass::epilogue::thread::Snake<cutlass::Array<Element, kV>>;

    double tolerance = 1e-5;

    //
    // Construct workspace
    //
    cutlass::HostTensor<Element, cutlass::layout::RowMajor> tensor_Destination({1, kN});
    cutlass::HostTensor<Element, cutlass::layout::RowMajor> tensor_X({1, kN});
    cutlass::HostTensor<Element, cutlass::layout::RowMajor> tensor_Alpha({1, kN});

    for (int i = 0; i < kN; ++i) {
        tensor_X.host_data(i) = Element(Snake_golden_x[i]);
        tensor_Alpha.host_data(i) = Element(Snake_golden_alpha[i]);
    }

    tensor_Destination.sync_device();
    tensor_X.sync_device();
    tensor_Alpha.sync_device();

    //
    // Launch the kernel
    //
    dim3 grid(1,1,1);
    dim3 block(kN / kV, 1, 1);

    test_Epilogue_thread_activation_binary<Element, kV, Func><<< grid, block >>>(
        tensor_Destination.device_data(),
        tensor_X.device_data(),
        tensor_Alpha.device_data());

    tensor_Destination.sync_host();

    //
    // Verify
    //

    for (int i = 0; i < kN; ++i) {
        Element x_in = Element(Snake_golden_x[i]);
        Element alpha_in = Element(Snake_golden_alpha[i]);
        Element got = tensor_Destination.host_data(i);
        Element expected = Element(Snake_golden_output[i]);

        double rel_error = (double(got) - double(expected)) / double(expected);

        EXPECT_LT(std::abs(rel_error), tolerance)
            << "Input[" << i << "]: x=" << x_in << ", alpha=" << alpha_in
            << ", Got: " << got << ", expected: " << expected;
    }
}

/////////////////////////////////////////////////////////////////////////////////////////////////

TEST(Epilogue_thread_snake, device_bf16) {

    int const kN = 256;
    int const kV = 8;

    using Element = cutlass::bfloat16_t;
    using Func = cutlass::epilogue::thread::Snake<cutlass::Array<Element, kV>>;

    double tolerance = 0.02;

    //
    // Construct workspace
    //
    cutlass::HostTensor<Element, cutlass::layout::RowMajor> tensor_Destination({1, kN});
    cutlass::HostTensor<Element, cutlass::layout::RowMajor> tensor_X({1, kN});
    cutlass::HostTensor<Element, cutlass::layout::RowMajor> tensor_Alpha({1, kN});

    for (int i = 0; i < kN; ++i) {
        tensor_X.host_data(i) = Element(Snake_golden_x[i]);
        tensor_Alpha.host_data(i) = Element(Snake_golden_alpha[i]);
    }

    tensor_Destination.sync_device();
    tensor_X.sync_device();
    tensor_Alpha.sync_device();

    //
    // Launch the kernel
    //
    dim3 grid(1,1,1);
    dim3 block(kN / kV, 1, 1);

    test_Epilogue_thread_activation_binary<Element, kV, Func><<< grid, block >>>(
        tensor_Destination.device_data(),
        tensor_X.device_data(),
        tensor_Alpha.device_data());

    tensor_Destination.sync_host();

    //
    // Verify
    //

    for (int i = 0; i < kN; ++i) {
        Element x_in = Element(Snake_golden_x[i]);
        Element alpha_in = Element(Snake_golden_alpha[i]);
        Element got = tensor_Destination.host_data(i);
        Element expected = Element(Snake_golden_output[i]);

        double rel_error = (double(got) - double(expected)) / double(expected);

        EXPECT_LT(std::abs(rel_error), tolerance)
            << "Input[" << i << "]: x=" << x_in << ", alpha=" << alpha_in
            << ", Got: " << got << ", expected: " << expected;
    }
}

/////////////////////////////////////////////////////////////////////////////////////////////////

TEST(Epilogue_thread_gelu_taylor, device_f32) {

    int const kN = 256;
    int const kV = 4;

    using Element = float;
    using Func = cutlass::epilogue::thread::GELU_taylor<cutlass::Array<Element, kV>>;

    double tolerance = 0.005;
    
    //
    // Construct workspace
    //
    cutlass::HostTensor<Element, cutlass::layout::RowMajor> tensor_Destination({1, kN});
    cutlass::HostTensor<Element, cutlass::layout::RowMajor> tensor_Source({1, kN});

    for (int i = 0; i < kN; ++i) {
        tensor_Source.host_data(i) = Element(GELU_golden_input[i]);
    }

    tensor_Destination.sync_device();
    tensor_Source.sync_device();

    //
    // Launch the kernel
    //
    dim3 grid(1,1,1);
    dim3 block(kN / kV, 1, 1);

    test_Epilogue_thread_activation<Element, kV, Func><<< grid, block >>>(
        tensor_Destination.device_data(),
        tensor_Source.device_data());

    tensor_Destination.sync_host();

    //
    // Verify
    //

    for (int i = 0; i < kN; ++i) {
        Element input = Element(GELU_golden_input[i]);
        Element got = tensor_Destination.host_data(i);
        Element expected = Element(GELU_golden_output[i]);

        double rel_error = (double(got) - double(expected)) / double(expected);

        double tolerance_override = tolerance;

        switch (i) {
            case 142: tolerance_override = 0.008; break;
            case 203: tolerance_override = 0.03; break;
            case 207: tolerance_override = 0.09; break;
            case 218: tolerance_override = 0.013; break;
        }

        EXPECT_LT(std::abs(rel_error), tolerance_override)
            << "Input[" << i << "]: " << input << ", Got: " << got << ", expected: " << expected;
    }
}

/////////////////////////////////////////////////////////////////////////////////////////////////

TEST(Epilogue_thread_gelu_taylor, device_f16) {

    int const kN = 256;
    int const kV = 8;

    using Element = cutlass::half_t;
    using Func = cutlass::epilogue::thread::GELU_taylor<cutlass::Array<Element, kV>>;

    double tolerance = 0.005;

    //
    // Construct workspace
    //
    cutlass::HostTensor<Element, cutlass::layout::RowMajor> tensor_Destination({1, kN});
    cutlass::HostTensor<Element, cutlass::layout::RowMajor> tensor_Source({1, kN});

    for (int i = 0; i < kN; ++i) {
        tensor_Source.host_data(i) = Element(GELU_golden_input[i]);
    }

    tensor_Destination.sync_device();
    tensor_Source.sync_device();

    //
    // Launch the kernel
    //
    dim3 grid(1,1,1);
    dim3 block(kN / kV, 1, 1);

    test_Epilogue_thread_activation<Element, kV, Func><<< grid, block >>>(
        tensor_Destination.device_data(),
        tensor_Source.device_data());

    tensor_Destination.sync_host();

    //
    // Verify
    //

    for (int i = 0; i < kN; ++i) {
        Element input = Element(GELU_golden_input[i]);
        Element got = tensor_Destination.host_data(i);
        Element expected = Element(GELU_golden_output[i]);

        double rel_error = (double(got) - double(expected)) / double(expected);
        
        double tolerance_override = tolerance;

        switch (i) {
            case 36: tolerance_override = 0.006; break;
            case 77: tolerance_override = 0.009; break;
            case 95: tolerance_override = 0.008; break;
            case 112: tolerance_override = 0.007; break;
            case 171: tolerance_override = 0.006; break;
            case 203: tolerance_override = 0.03; break;
            case 207: tolerance_override = 0.15; break;
        }

        EXPECT_LT(std::abs(rel_error), tolerance_override) 
            << "Input[" << i << "]: " << input << ", Got: " << got << ", expected: " << expected;
    }
}

/////////////////////////////////////////////////////////////////////////////////////////////////