j dZddlZddlZddlmZddlmZejddgddgdej fdZ d Z ejddgddgd ej dej fd Z dd Z ddlmZedS)aAlias-free SnakeBeta activation of the H3 audio decoder (BigVGAN) in one Triton kernel: 2x upsampling (replicate padding, 12-tap transposed Kaiser-sinc filter), SnakeBeta, then 2x downsampling (replicate padding, 12-tap low-pass filter). The unfused module runs a dozen kernels per activation (the decoder has about a hundred and thirty of them) and writes two intermediates of twice its length.N) libdevicelengthchannels)do_not_specializedo_not_specialize_on_alignmentBLOCKc Btjd} tjd|ztjd|z} | |k} | |z} tjtj|| ztj} dtjtj|| ztjdzz }|| tj|zz}tj|gtj}tj dD]}tj tj d| z|zdz dd|zdz }|d zdz}|d zd|zz }tj|gtj}tj d D]}tj tj ||z dz d|dz }|tj||z| d tjtj||zd|zztjzz }d |z}tj | |z}||||z|zztj||ztjzz }tj|| tj|zz| z||jj| dS)Nr?& .>)dtype maskotherg@r)tl program_idarangeexploadtofloat32int64zeros static_rangeminimummaximumrsinstorer element_ty)XYALPHABETAUPDOWNrrrrownvalidchannelalphainv_betax_rowoutkmlasttapupjsourcesines CC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\audio_triton.py_snake_beta_aar> s -  C a5 29Q#6#66A JEHnG F2757?++..rz:: ; ;EbfRWTG^4477 CCDDtKLH rx  6) )E (E7"* - - -C _R P P Jrz!a%!)a-33QZ!^ D DB1}"fq4x XugRZ 0 0 0## { {AZ 4!8a< ; ;VaZHHF "'%&.uC@@@CCBJOORTRYZ\_bZbefijejZjRkRkRnRnoqoyRzRzz zBB 2X}URZ(( X_t++rwtax/@/@/C/CBJ/O/OOOHQ!!F* *Q .qw7I0J0JQVWWWWWWc $|}tj|}|j\}}}t ||zt j|df|||||d|d||d |S)z{x [batch, channels, length], alpha and beta [channels] (log space), 12-tap filters: the activation's output, shaped like x.i)r) contiguoustorch empty_likeshaper>tritoncdivreshape) xr1beta up_filter down_filterr4batchrrs r= snake_beta_aarN$s A  1  CgE8VEH$fk&#&>&>?@CPTV_VgVghjVkVkmxnAnABDnEnEGMAIQTVVVV Jr? LOG_SCALEcRtjd}tjd|ztjd|z}||k} ||z} tj|| ztj} tj|| ztj} |r(t j| } t j| } tjd| dz} ||tj |zz|z}tj|| dtj}t j | |z}tj ||||z| zz|j j | dS) zx += sin(a x)^2 / (b + 1e-9) in place, one row of (batch * channels) per program along axis 0: a and b per channel, stored in log scale (exp taken here in FP32) or directly. FP32 arithmetic from the loaded values, one rounding at the store.rr r r rrrN)rrrrrrrrdiv_rnrr$r%r r&)r'r)r*rrrOrr-r.r/r0r1rJr2ptrrIwaves r=_snakerT.sd -  C a5 29Q#6#66A JEHnG GEGO $ $ ' ' 3 3E 74'> " " % %bj 1 1D# e$$}T""ydTk**H cffRX' '! +C %s+++..rz::A = # #DHS1td{X--11!'2DEEERRRRRRr?Tc |j\}}}t||ztj|df||d|d|||ddt jj||S)zaIn place on x (batch, channels, length), contiguous; alpha, beta (1, channels, 1) or (channels,).irAF)rOrenable_fp_fusion) rErTrFrGrHrCautogradgraphincrement_version)rIr1rJ log_scalerMrrs r=snake_r[BsgE8V EH fk&$77 89!U]]2=N=NPTP\P\]_P`P`bhjrHPTLQSSSS N**1--- Hr?)!install_triton_compilation_logger)T)__doc__rCrFtriton.languagelanguagertriton.language.extrarjit constexprr>rNrTr[%shared.kernels.triton_compilation_logr\r?r=res:GG ++++++ x4V^`jUklllXXXXmlX2 x4V^`jUklllS SR\SSSmlS&    TSSSSS!!#####r?