+
    "Ý³j&  ã                   óà   € ^ RI Ht ^ RIHtHt ^ RIt^RIHtHt ^RI	H
t
Ht ^RIHtHt ^RIHt ^RIHt ^R	IHt ^R
IHt ]! 4       tRR R llt ! R R]4      tR.tR R lt]R 4       tR# )é    )Ú	lru_cache)ÚOptionalÚTupleN)ÚAdapterÚLora)ÚDEVICEÚPLATFORM)ÚFORMATÚMETHOD)ÚBACKEND)Úsetup_logger)Úhas_gil_disabled)ÚTorchLinearÚTritonV2Linearc                ó”   € V ^8„  d   QhR\         \        P                  ,          R\         \        P                  ,          R\        RR/# )é   Úg_idxÚscalesÚ
layer_nameÚreturnN)r   ÚtorchÚTensorÚstr)Úformats   "Úf/home/luca/screen-checking/.venv/lib/python3.14/site-packages/gptqmodel/nn_modules/qlinear/tritonv2.pyÚ__annotate__r      sB   € ÷ 
ñ 
Ü”E—L‘LÕ!ð
ä”U—\‘\Õ"ð
ô ð
ð 
ñ	
ó    c                ó2  € V e   Ve   V P                  4       ^ 8X  d   R# VP                  ^ ,          p\        V P                  4       4      p\        V P	                  4       4      pWC) 8  g   WS8¼  d&   \        V RV RV RV RV)  RV^,
           R24      hR# )a�  Reject checkpoint ``g_idx`` that would index the group buffers out of bounds.

The Triton dequant kernel resolves each row's group as
``groups = where(g_idx < 0, g_idx + num_groups, g_idx)`` and then loads
``scales``/``qzeros`` at that group with no upper-bound guard, where
``num_groups = scales.shape[0]``. A value outside ``[-num_groups, num_groups)``
therefore indexes past the buffer, causing an out-of-bounds device read
(denial of service, possible adjacent-memory disclosure) when an untrusted
checkpoint is loaded. This mirrors the safe faulting of the Torch backend's
``scales[g_idx]`` gather, but validates once at load rather than per forward.
Nz': checkpoint g_idx is out of range for z scale group(s) (got min=z, max=z	, valid [z, z;]); the group index would read scales/qzeros out of bounds.)ÚnumelÚshapeÚintÚminÚmaxÚ
ValueError)r   r   r   Ú
num_groupsÚg_minÚg_maxs   &&&   r   Ú_validate_g_idx_boundsr(      s§   € ð  ‚}˜š¨%¯+©+«-¸1Ô*<Ùà—‘˜a•€JÜ�—	‘	“Ó€EÜ�—	‘	“Ó€EØˆ{Ô˜eÔ1ÜØˆlÐAÀ*Àð N'Ø', g¨V°E°7ð ;Ø�ˆ}˜B˜z¨A�~Ð.ð /+ð,ó
ð 	
ñ 2r   c                   ó  a a€ ] tR t^5t o]P
                  .t]P                  .t	]
P                  ^(]
P                  ^(/t. ROt. ROtRR.tRR.tRtRtRt^ .t^ .t]P.                  .t]P4                  ]P6                  .t]P<                  ]P>                  ]P@                  .t!]".t#]PH                  ]PJ                  .t&Rt'Rt( R]P<                  RR3V3R lV 3R lllt)]*V3R lR l4       t+]*V3R	 lR
 l4       t,V 3R lt-V 3R lt.Rt/Vt0V ;t1# )r   TFÚtritonv2Nc                ó^   <€ V ^8„  d   QhRS[ RS[ RS[RS[RS[RS[P                  RS[RS[/# )	r   ÚbitsÚ
group_sizeÚdesc_actÚsymÚbiasÚ
pack_dtypeÚadapterÚregister_buffers)r!   Úboolr   Údtyper   )r   Ú__classdict__s   "€r   r   ÚTritonV2Linear.__annotate__X   s_   ø€ ÷ ñ áðñ ðñ ð	ñ
 ðñ ðñ Ÿ™ðñ ðñ #ñr   c                ó�   <€ \         SV `  ! RR VRVRVRVRVRVRVRVRVP                  R\        P                  4      R	V	R
V
/VB  R# )r,   r-   r/   r.   Úin_featuresÚout_featuresr0   r1   Úbackendr2   r3   N© )ÚsuperÚ__init__Úpopr   ÚGPTQ_TRITON)Úselfr,   r-   r.   r/   r9   r:   r0   r1   r2   r3   ÚkwargsÚ	__class__s   &&&&&&&&&&&,€r   r>   ÚTritonV2Linear.__init__X   s—   ø€ ô 	‰Òñ 	Øð	à!ð	ð ð	ð ð		ð
 $ð	ð &ð	ð ð	ð "ð	ð —J‘J˜y¬'×*=Ñ*=Ó>ð	ð ð	ð .ð	ð ô	r   c                óF   <€ V ^8„  d   QhRS[ S[S[S[,          3,          /# ©r   r   ©r   r4   r   Ú	Exception)r   r6   s   "€r   r   r7   {   s#   ø€ ÷ ñ ™e¡D©(±9Õ*=Ð$=Õ>ñ r   c                ó  € ^ RI p^ RIHp ^ RIHp ^ RI Hp ^RIHp ^RIH	p VP                  V4      pWsP                  R4      8  d   \        RV 24      h\        4       '       d!   WsP                  R4      8  d   \        R	4      hR
# )r   N)Úversion)Ú__version__©ÚQuantLinearFunction)ÚTritonModuleMixinz2.0.0z*triton version must be >= 2.0.0: actual = z3.4.0zYGIL is disabled and not compatible with current Triton. Please upgrade to Triton >= 3.4.0)TN)ÚtritonÚtriton.languageÚlanguageÚ	packagingrJ   rK   Útriton_utils.dequantrM   Útriton_utils.mixinrN   ÚparseÚImportErrorr   rH   )ÚclsrO   ÚtlrJ   Útriton_versionrM   rN   Útriton_vs   &       r   Úvalidate_onceÚTritonV2Linear.validate_oncez   sp   € ãÝ$Ý%Ý8å>Ý:à—=‘= Ó0ˆà—m‘m GÓ,Ô,ÜÐ JÈ>ÐJZÐ[Ó\Ð\ô ×Ò (¯]©]¸7Ó-CÔ"CÜÐwÓxÐxàÐr   c                óF   <€ V ^8„  d   QhRS[ S[S[S[,          3,          /# rF   rG   )r   r6   s   "€r   r   r7   �   s#   ø€ ÷ %ñ %¡¡t©X±iÕ-@Ð'@Õ!Añ %r   c                ó¬   € VP                  R 4      pV\        P                  8X  d   \        4       '       g   R\	        R4      3# V P
                  ! R/ VB # )ÚdeviceFz½Trying to use the triton backend and xpu device, but it could not be imported. Please install triton by [intel-xpu-backend-for-triton](https://github.com/intel/intel-xpu-backend-for-triton)r<   )Úgetr   ÚXPUÚtriton_xpu_availabler$   Ú	_validate)rW   Úargsr_   s   &, r   ÚvalidateÚTritonV2Linear.validate�   sX   € à—‘˜(Ó#ˆð ”V—Z‘ZÔÔ(<×(>Ò(>Øœ*ð PóQð Qð Qð �}Š}Ñ$˜tÑ$Ð$r   c                óŽ   <€ \         SV `  4        \        V P                  V P                  \        V 4      P                  R 7       R# ))r   N)r=   Ú	post_initr(   r   r   ÚtypeÚ__name__)rA   rC   s   &€r   rh   ÚTritonV2Linear.post_initš   s3   ø€ ô 	‰ÑÔô 	Ø�J‰J˜Ÿ™´°T³
×0CÑ0C÷	
r   c           
     ó¸  <€ ^RI Hp V P                  '       d   \        SV `  V4      # V P                  V4      pVP                  RR V P                  3,           pVP                  VP                  RVP                  R,          4      V P                  V P                  V P                  V P                  V P                  V P                  V P                   4      P                  V4      pV P"                  e   VP%                  V P"                  4       V P&                  '       d   V P&                  P                  WR7      pVP)                  VP*                  R7      # )r   rL   N)ÚxÚout)r5   éÿÿÿÿ)rS   rM   Útrainingr=   ÚforwardÚ_apply_rotation_to_inputr    r:   ÚapplyÚreshapeÚqweightr   Úqzerosr   r,   Úpack_dtype_bitsÚmaxqr0   Úadd_r2   Útor5   )rA   rm   rM   Ú	out_shapern   rC   s   &&   €r   rq   ÚTritonV2Linear.forward°   s  ø€ Ý>à�=�=ˆ=Ü‘7‘? 1Ó%Ð%à×)Ñ)¨!Ó,ˆð —G‘G˜C˜R�L D×$5Ñ$5Ð#7Õ7ˆ	à!×'Ñ'Ø�I‰I�b˜!Ÿ'™' "�+Ó&Ø�L‰LØ�K‰KØ�K‰KØ�J‰JØ�I‰IØ× Ñ Ø�I‰Ió	
÷ ‰'�)Ó
ð 	ð �9‰9Ò Ø�H‰H�T—Y‘YÔà�<�<ˆ<Ø—,‘,×$Ñ$ qÐ$Ó2ˆCà�v‰v˜AŸG™GˆvÓ$Ð$r   r<   )r   é   é   )ro   é   é    é@   é€   é   i   é   )2rj   Ú
__module__Ú__qualname__Ú__firstlineno__r   r@   ÚSUPPORTS_BACKENDSr   ÚGPTQÚSUPPORTS_METHODSr
   ÚGPTQ_V2ÚSUPPORTS_FORMATSÚSUPPORTS_BITSÚSUPPORTS_GROUP_SIZEÚSUPPORTS_DESC_ACTÚSUPPORTS_SYMÚSUPPORTS_SHARDSÚSUPPORTS_TRAININGÚSUPPORTS_AUTO_PADDINGÚ!SUPPORTS_IN_FEATURES_DIVISIBLE_BYÚ"SUPPORTS_OUT_FEATURES_DIVISIBLE_BYr   ÚCUDAÚSUPPORTS_DEVICESr	   ÚLINUXÚWIN32ÚSUPPORTS_PLATFORMr   Úint32Úint16Úint8ÚSUPPORTS_PACK_DTYPESr   ÚSUPPORTS_ADAPTERSÚfloat16Úbfloat16ÚSUPPORTS_DTYPESÚREQUIRES_FORMAT_V2Ú
QUANT_TYPEr>   Úclassmethodr[   re   rh   rq   Ú__static_attributes__Ú__classdictcell__Ú__classcell__)rC   r6   s   @@r   r   r   5   s0  ù‡ € Ø ×,Ñ,Ð-ÐØŸ™�}ÐØŸ™ R¨¯©¸Ð<ÐÚ€MÚ?ÐØ˜u˜ÐØ˜%�=€LØ€OØÐØ ÐØ)+¨Ð%Ø*,¨Ð&ð Ÿ™�}ÐØ!Ÿ™¨¯©Ð8ÐØ!ŸK™K¨¯©°e·j±jÐAÐØ˜Ðà—}‘} e§n¡nÐ5€OàÐð €Jðð  Ø&+§k¡kØ#Ø%)÷õ ðD ÷ó ðð( ÷%ó ð%õ
÷,%õ %r   c                óX   € V ^8„  d   QhR\         P                  R\         P                  /# )r   rm   Úy)r   r   )r   s   "r   r   r   Ö   s"   € ÷ ñ ”u—|‘|ð ¬¯©ñ r   c                 óÔ   aaa€ ^ RI o^ RIHo SP                  V3R lV3R ll4       p\        P
                  ! V 4      pVP                  4       oVV3R lpW$,          ! WVSRR7       V# )r   Nc                ó2   <€ V ^8„  d   QhRSP                   /# )r   Ú
BLOCK_SIZE)Ú	constexpr)r   rX   s   "€r   r   Ú%triton_test_add.<locals>.__annotate__Ü   s   ø€ ÷ ñ ð  "Ÿ|™|ñ	r   c                 óê   <€ SP                  ^ R7      pWT,          pVSP                  ^ V4      ,           pWs8  pSP                  W,           VR7      p	SP                  W,           VR7      p
Wš,           pR# )r   )Úaxis)ÚmaskN)Ú
program_idÚarangeÚload)Úx_ptrÚy_ptrÚ
output_ptrÚ
n_elementsr­   ÚpidÚblock_startÚoffsetsr²   rm   rª   ÚoutputrX   s   &&&&&       €r   Ú
add_kernelÚ#triton_test_add.<locals>.add_kernelÛ   si   ø€ ð �m‰m ˆmÓ#ˆØÕ&ˆØ §	¡	¨!¨ZÓ 8Õ8ˆØÑ#ˆØ�G‰G�E•O¨$ˆGÓ/ˆØ�G‰G�E•O¨$ˆGÓ/ˆØ•Šr   c                 ó8   <€ SP                  SV R ,          4      3# ©r­   )Úcdiv)Úmetar¹   rO   s   &€€r   ÚgridÚtriton_test_add.<locals>.gridí   s   ø€ Ø—‘˜J¨¨\Õ(:Ó;Ð=Ð=r   r„   rÁ   )rO   rP   rQ   Újitr   Ú
empty_liker   )rm   rª   r¾   r½   rÄ   r¹   rX   rO   s   &&   @@@r   Útriton_test_addrÈ   Ö   s`   ú€ ÛÝ ð ‡Z�Z÷ò ó ðô ×Ò˜aÓ €FØ—‘“€Jö>ð Ö�Q˜6 :¸$Õ?Ø€Mr   c                  ó¦   € R p \         P                  ! V RR7      p\         P                  ! V RR7      p \        W4       R#   \         d     R# i ; i)r„   zxpu:0)r_   TF)r   ÚrandrÈ   rH   )Úsizerm   rª   s      r   rb   rb   ô   sI   € à€DÜ�
Š
�4 Ô(€AÜ�
Š
�4 Ô(€AðÜ˜ÔÙøÜô Úðús   ´A ÁAÁA)r   )Ú	functoolsr   Útypingr   r   r   Úadapter.adapterr   r   Úmodels._constr   r	   Úquantizationr
   r   Úutils.backendr   Úutils.loggerr   Úutils.pythonr   r   Úlogr(   r   Ú__all__rÈ   rb   r<   r   r   Ú<module>rÖ      se   ðõ
  ß "ã ç ,ß -ß *Ý $Ý (Ý ,Ý ñ ƒn€÷
ô>Z%�[ô Z%ðz Ð
€õð< ñ	ó ò	r   