§
    ‘M§j�  ã                   ó  — d Z ddlZddlmZ ddlmZ ej        dej        dej        fd„¦   «         Z	ej        dej        dej        d	ej        d
ej        dej        dej        dej        fd„¦   «         Z
d„ ZddlmZ  e¦   «          dS )zITile-local ConvRot for small INT8 decode inputs; no activation workspace.é    N)Ú	libdeviceÚBMÚSTRIDEc                 ón  — |                       |dd|z  z  d|f¦  «        }t          j        dt          j        dd¦  «        z
  d d d d …d f         |dd|z  z  d|f¦  «        }t          j        |d¬¦  «        d d …d d …d d d …f         dt          j        ||d¬¦  «        z  z
                        |df¦  «        S )Né   é   é   r   é   ©Úaxisg       @)ÚreshapeÚtlÚbroadcast_toÚarangeÚsumÚgather)Úar   r   ÚvÚreverses        úJC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\convrot_int8_triton.pyÚ_regular_stepr      sÍ   € à	�	Š	�2�s˜q 6™zÑ*¨A¨vÐ6Ñ7Ô7€AÝŒo˜q¥2¤9¨Q°¡?¤?Ñ2°D¸$ÀÀÀÀ4Ð4GÔHÈ2ÈsÐWXÐ[aÑWaÑObÐdeÐgmÐJnÑoÔo€GÝŒF�1˜1ÐÑÔ˜a˜a˜a    D¨!¨!¨!˜mÔ,¨sµR´Y¸qÀ'ÐPQÐ5RÑ5RÔ5RÑ/RÑR×[Ò[Ð]_ÐadÐ\eÑfÔfÐfó    ÚSX0ÚSX1ÚSW0ÚSW1ÚBNÚHAS_BIASc                 óä	  — t          j        d¦  «        }t          j        ||¦  «        }||z  |z  t          j        d|¦  «        z   }||z  |z  t          j        d|¦  «        z   }t          j        dd¦  «        }t          j        dd¦  «        }t          j        ||fdt           j        ¦  «        }t          d|d¦  «        D �]}t          j        | |d d …d f         |z  z   ||d d d …f         z   |	z  z   |d d …d f         |k     d¦  «                             t           j        ¦  «        }t          ||d¦  «        }t          ||d¦  «        }t          ||d¦  «        }t          ||d¦  «        }|dz                       | j
        j        ¦  «                             t           j        ¦  «                             |ddf¦  «        }t          j        t          j        |¦  «        d¬	¦  «        d
z  }t          j        |dk    |d¦  «        }t!          j        ||d d …d d …d f         z  ¦  «        }t          j        t          j        |d¦  «        d¦  «                             t           j        ¦  «                             |df¦  «        }t          j        d¦  «        D �]6}t          j        |t          j        |dz  |z   d d d …f         |df¦  «        d¬	¦  «        }t          j        |t          j        |df|t           j        ¦  «        d¬	¦  «                             |f¦  «        }t          j        ||d d d …f         |
z  z   ||dz  z   |d d …d f         z   |z  z   |d d d …f         |k     d¦  «                             t           j        ¦  «        }|t          j        ||¦  «                             t           j        ¦  «        |d d …d f         z  z  }�Œ8�Œ|t          j        ||z   ||k     d¦  «        d d d …f                              t           j        ¦  «        z  }|r�|                     |j
        j        ¦  «                             t           j        ¦  «        t          j        ||z   ||k     d¦  «        d d d …f                              t           j        ¦  «        z   }t          j        ||d d …d f         |z  z   |d d d …f         z   ||d d …d f         |k     |d d d …f         |k     z  ¦  «         d S )Nr   r   é@   é   r   é   g      °?r
   r   g     À_@g      ð?é   i€ÿÿÿ)r   Ú
program_idÚcdivr   ÚfullÚfloat32ÚrangeÚloadÚtor   ÚdtypeÚ
element_tyr   ÚmaxÚabsÚwherer   ÚrintÚmaximumÚminimumÚint8Ústatic_ranger   r   Úint32ÚdotÚstore)ÚxÚwÚscaleÚoutÚbiasÚMÚNÚKr   r   r   r   r   r   r   ÚpidÚtiles_mÚrowsÚcolsÚkkÚqkÚaccÚbaser   ÚscalesÚquantÚgÚqaÚsÚbÚresults                                  r   Ú_convrot_int8_kernelrO      s©  € õ Œ-˜Ñ
Ô
€CÝŒg�a˜‰nŒn€GØ�'‰M˜RÑ¥"¤)¨A¨rÑ"2Ô"2Ñ2€DØ�7‰N˜bÑ ¥2¤9¨Q°Ñ#3Ô#3Ñ3€DÝ	Œ�1�cÑ	Ô	€BÝ	Œ�1�bÑ	Ô	€BÝ
Œ'�2�r�(˜A�rœzÑ
*Ô
*€CÝ�a˜˜CÑ Ô ð =ñ =ˆÝŒG�A˜˜Q˜Q˜Q ˜Wœ¨Ñ+Ñ+¨t°b¸¸q¸q¸q¸´kÑ/AÀSÑ.HÑHÈ$ÈqÈqÈqÐRVÈwÌ-ÐZ[ÒJ[Ð]^Ñ_Ô_×bÒbÕceÔcmÑnÔnˆÝ˜!˜R Ñ#Ô#ˆÝ˜!˜R Ñ#Ô#ˆÝ˜!˜R Ñ$Ô$ˆÝ˜!˜R Ñ$Ô$ˆà�‰Z�OŠO˜AœGÔ.Ñ/Ô/×2Ò2µ2´:Ñ>Ô>×FÒFÈÈAÈrÀ{ÑSÔSˆÝ”�œ˜q™	œ	¨Ð*Ñ*Ô*¨UÑ2ˆÝ”˜& 1š* f¨cÑ2Ô2ˆÝ”˜q 6¨!¨!¨!¨Q¨Q¨Q°¨*Ô#5Ñ5Ñ6Ô6ˆÝ”
�2œ: e¨SÑ1Ô1°4Ñ8Ô8×;Ò;½B¼GÑDÔD×LÒLÈbÐRUÈYÑWÔWˆÝ” Ñ#Ô#ð 	=ñ 	=ˆAÝ”˜5¥"¤/°1°r±6¸B±;ÀÀaÀaÀaÀÔ2HÈ2ÈrÈ(Ñ"SÔ"SÐZ[Ð\Ñ\Ô\ˆBÝ”	˜&¥"¤'¨2¨q¨'°1µb´hÑ"?Ô"?ÀaÐHÑHÔH×PÒPÐRTÐQVÑWÔWˆAÝ”˜˜D  q q q œM¨CÑ/Ñ/°4¸!¸b¹&±=À2ÀaÀaÀaÈÀgÄ;Ñ3NÐRUÑ2UÑUÐW[Ð\`ÐbcÐbcÐbcÐ\cÔWdÐghÒWhÐjkÑlÔl×oÒoÕprÔpwÑxÔxˆAØ•2”6˜"˜a‘=”=×#Ò#¥B¤JÑ/Ô/°!°A°A°A°t°G´*Ñ<Ñ<ˆC‰Cñ		=ð
 •2”7˜5 4™<¨°ª°1Ñ5Ô5°d¸A¸A¸A°gÔ>×AÒAÅ"Ä*ÑMÔMÑM€FØð |Ø—’˜3œ9Ô/Ñ0Ô0×3Ò3µB´JÑ?Ô?Å"Ä'È$ÐQUÉ+ÐW[Ð^_ÒW_ÐabÑBcÔBcÐdhÐjkÐjkÐjkÐdkÔBl×BoÒBoÕprÔpzÑB{ÔB{Ñ{ˆÝ„HˆS�4˜˜˜˜4˜”= 1Ñ$Ñ$ t¨D°!°!°!¨G¤}Ñ4°f¸tÀAÀAÀAÀtÀG¼}ÈqÒ?PÐUYÐZ^Ð`aÐ`aÐ`aÐZaÔUbÐefÒUfÑ>gÑhÔhÐhÐhÐhr   c           
      ó0  — | j         \  }}|j         d         }|\  }	}
}}}t          t          j        ||	¦  «        t          j        ||
¦  «        z  f         | |||||||g|                      ¦   «         ¢|                     ¦   «         ¢|	‘|
‘|d u‘R ||dœŽ d S )Nr   )Ú	num_warpsÚ
num_stages)ÚshaperO   Útritonr%   Ústride)r8   Úweightr:   r;   r<   ÚcfgÚmÚkÚnÚbmÚbnÚ_ÚwarpsÚstagess                 r   Úconvrot_int8_mmr`   0   s  € ØŒ7�D€A€qØŒ�QŒ€AØ"Ñ€BˆˆAˆu�fÝ�&œ+ a¨Ñ,Ô,­v¬{¸1¸bÑ/AÔ/AÑAÐCÔDÀQÈÐPUÐWZÐ\`ÐbcÐefÐhið  HÐlm×ltÒltÑlvÔlvð  HÐy÷  zGò  zGñ  zIô  zIð  Hð  KMð  Hð  OQð  Hð  SWð  _cð  Scð  Hð  Hð  otð  AGð  Hð  Hð  Hð  Hð  Hð  Hr   )Ú!install_triton_compilation_logger)Ú__doc__rT   Útriton.languageÚlanguager   Útriton.language.extra.cudar   ÚjitÚ	constexprr   rO   r`   Ú%shared.kernels.triton_compilation_logra   © r   r   ú<module>rj      s#  ðØ OÐ OØ €€€Ø Ð Ð Ð Ð Ð Ø 0Ð 0Ð 0Ð 0Ð 0Ð 0ð „ðg˜œð g¨r¬|ð gð gð gñ „ðgð „ðiØ œlðiØ13´ðiØDFÄLðiØWYÔWcðiàœ\ðià/1¬|ðiàGIÄ|ðið ið iñ „ðiðBHð Hð Hð TÐ SÐ SÐ SÐ SÐ SØ !Ð !Ñ #Ô #Ð #Ð #Ð #r   