2e56847d72
Introduces cutlass::epilogue::thread::Snake, a two-operand activation functor implementing Snake_a(x) = x + (1/a) * sin^2(a*x) from Ziyin et al. 2020 (arXiv:2006.08195). The per-channel learnable frequency `a` flows through an EVT child (e.g. Sm90RowBroadcast), composing into Sm90EVT<Sm90Compute<Snake, ...>, x_node, alpha_node> for fused GEMM+Snake epilogues used in neural vocoders. Adds unit tests in test/unit/epilogue/thread/activation.cu covering f32 and bf16 paths, validated against float64 reference goldens. Closes #3141
653 lines
34 KiB
Plaintext
653 lines
34 KiB
Plaintext
/***************************************************************************************************
|
|
* 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;
|
|
}
|
|
}
|
|
|
|
/////////////////////////////////////////////////////////////////////////////////////////////////
|