+
    "Ý³j¼  ã                   óx  € ^ RI t ^ RIt^ RIHt ^ RItRP                  4       tRtRt. ROt	. ROt
. ROt. R OtRR.tR R	 ltR
 R ltR R ltR R ltR R ltR!R R lltR!R R lltR R ltR!R R llt]R8X  d=   R]P0                  R,          9   d   ]! ]! 4       '       d   ^ 4      h^4      h]! 4        R# R# )"é    N©ÚPathz“
// auto generated by generate_kernels.py
// clang-format off

#include "kernel.h"
#include "marlin_template.h"

namespace MARLIN_NAMESPACE_NAME {
z}
a!  template __global__ void Marlin<{{scalar_t}}, {{w_type_id}}, {{s_type_id}}, {{threads}}, {{thread_m_blocks}}, {{thread_n_blocks}}, {{thread_k_blocks}}, {{'true' if m_block_size_8 else 'false'}}, {{stages}}, {{group_blocks}}, {{'true' if is_zp_float else 'false'}}>( MARLIN_KERNEL_PARAMS );Úfp16Úbf16c                ó0   € V ^8„  d   QhR\         R\        /# )é   Úscalar_typeÚreturn)ÚstrÚbool)Úformats   "Úf/home/luca/screen-checking/.venv/lib/python3.14/site-packages/gptqmodel_ext/marlin/generate_kernels.pyÚ__annotate__r   5   s   € ÷ ñ ¤ð ¬ñ ó    c                 ó   € V R9   # )ú	vllm::kU4>   r   úvllm::kU4B8úvllm::kFE2M1f© )r	   s   &r   Ú_is_4bit_weightr   5   s   € Øð ñ ð r   c                ó   € V ^8„  d   QhRR/# )r   r
   Nr   )r   s   "r   r   r   =   s   € ÷ %ñ %˜Dñ %r   c                  óŒ   € \        \        4      P                  p V P                  R 4       F  pVP	                  RR7       K  	  R# )úkernel_*.cuT©Ú
missing_okN)r   Ú__file__ÚparentÚglobÚunlink)ÚrootÚpaths     r   Úremove_old_kernelsr"   =   s2   € Ü”‹>× Ñ €DØ—	‘	˜-Ö(ˆØ�‰˜tˆÖ$ó )r   c                óH   € V ^8„  d   QhR\         R\        R\        R\         /# )r   r    r	   Údtyper
   )r   r   )r   s   "r   r   r   C   s(   € ÷ 7ñ 7œdð 7´ð 7¼Sð 7ÄTñ 7r   c                 óž   € R V9   d(   VP                  R ^4      ^,          P                  4       MVP                  4       pV RV RV R2,          # )z::Úkernel_Ú_z.cu)ÚsplitÚlower)r    r	   r$   Úscalar_suffixs   &&& r   Ú_kernel_output_pathr+   C   sO   € Ø=AÀ[Ô=P�K×%Ñ% d¨AÓ.¨qÕ1×7Ñ7Ô9ÐVa×VgÑVgÓVi€MØ�G˜E˜7 ! M ?°#Ð6Õ6Ð6r   c                ó^   € V ^8„  d   QhR\         R\         R\        \         ,          R\         /# )r   r	   r$   Ú	templatesr
   ©r   Úlist)r   s   "r   r   r   H   s-   € ÷ ñ ¬#ð ´cð ÄdÌ3Åið ÔTWñ r   c                 óÐ   € \         R RV RV  2R .pVP                  RP                  V4      4       VP                  R 4       VP                  \        4       RP                  V4      # )Ú z// Instantiations for dtype=z	, weight=Ú
)Ú	FILE_HEADÚappendÚjoinÚ	FILE_TAIL)r	   r$   r-   Úliness   &&& r   Ú_render_kernel_file_textr8   H   s[   € Ü˜Ð:¸5¸'ÀÈ;È-ÐXÐZ\Ð]€EØ	‡L�L�—‘˜9Ó%Ô&Ø	‡L�L�ÔØ	‡L�L”ÔØ�9‰9�UÓÐr   c          
      ój   € V ^8„  d   QhR\         R\        R\        R\        \        ,          R\         /# )r   r    r	   r$   r-   r
   )r   r   r/   )r   s   "r   r   r   P   s5   € ÷ ñ œTð ´ð ¼Cð ÌDÔQTÍIð ÔZ^ñ r   c                 óX   € \        WV4      pVP                  \        WV4      R R7       V# )úutf-8©Úencoding)r+   Ú
write_textr8   )r    r	   r$   r-   Úoutput_paths   &&&& r   Ú_write_kernel_filer@   P   s0   € Ü% d¸Ó?€KØ×ÑÔ3°KÈ	ÓRÐ]dÐÔeØÐr   c                ó`   € V ^8„  d   QhR\         R,          R\        \         \        3,          /# ©r   r    Nr
   )r   Údictr   )r   s   "r   r   r   V   s&   € ÷ ñ ¤¨¥ð ¼¼TÄ3¸Y½ñ r   c                 ó$  € T ;'       g    \        \        4      P                  p / p\        P                  ! \
        \        4       F4  w  r#\        W#4      pV'       g   K  \        WV4      p\        W#V4      W&   K6  	  V'       g   \        R 4      hV# )z?No marlin kernels were generated; check template configuration.)r   r   r   Ú	itertoolsÚproductÚSCALAR_TYPESÚDTYPESÚrender_templates_for_combor+   r8   ÚRuntimeError)r    Úexpectedr	   r$   r-   r?   s   &     r   Úbuild_expected_kernelsrL   V   sz   € Ø×(Ð(”4œ“>×(Ñ(€DØ "€HÜ'×/Ò/´¼fÖEÑˆÜ.¨{ÓBˆ	ßÙÜ)¨$¸UÓCˆÜ 8¸ÈYÓ WˆÓñ F÷ ÜÐ\Ó]Ð]Ø€Or   c                ó>   € V ^8„  d   QhR\         R,          R\        /# rB   )r   r   )r   s   "r   r   r   e   s   € ÷ ñ ¬¨t­ð ¼tñ r   c                 óœ  € T ;'       g    \        \        4      P                  p \        V 4      pV Uu0 uF  q"P                  kK  	  ppV P                  R 4       Uu0 uF  q"P                  kK  	  ppWC8w  d   R# VP                  4        F!  w  rV VP                  RR7      pYv8w  g   K    R# 	  R# u upi u upi   \         d      R# i ; i)r   Fr;   r<   T)	r   r   r   rL   Únamer   ÚitemsÚ	read_textÚOSError)r    rK   r!   Úexpected_namesÚexisting_namesr?   Úexpected_textÚcurrent_texts   &       r   Úgenerated_kernels_are_currentrW   e   s»   € Ø×(Ð(”4œ“>×(Ñ(€DÜ% dÓ+€HÙ,4Ó5©H D—i”i©H€NÐ5Ø,0¯I©I°mÔ,DÓEÑ,D D—i”iÑ,D€NÐEØÔ'Ùà&.§n¡nÖ&6Ñ"ˆð	Ø&×0Ñ0¸'Ð0ÓBˆLð Ö(Úñ '7ñ ùò 6ùÚEøô ô 	Ûð	ús   ³B1ÁB6ÂB;Â;CÃ
Cc                óR   € V ^8„  d   QhR\         R\         R\        \         ,          /# )r   r	   r$   r
   r.   )r   s   "r   r   r   w   s*   € ÷ [ñ [¬Cð [¼ð [ÄÄSÅ	ñ [r   c                 óP  € . pR .pVR8X  d   VP                  ^ ^4       \        P                  ! \        \        \
        4       EFÞ  w  rEpV^ 8X  d
   V R9  d   K  V^,          R8X  d/   V^8:  d   V^ ,          ^€8w  d   K<  V^8”  d   V^ ,          ^@8w  d   KS  V R8X  d
   VR9  d   Kc  V R8X  d
   VR9  d   Ks  V R8w  d
   V^8X  d   Kƒ  V^ ,          ^,          pV^,          ^,          pV^,          p	VR8X  d   RMRp
R.pVR8X  d    V R8X  d   V^8X  d   VP                  R	4       V R8X  d   V^8X  d   RpMIV R8X  d   V^8X  d   R
pVR8X  d   EK  M.V R8X  d   V^8X  d   R
pVR8X  d   EK&  MVR8X  d   RpM	VR8X  d   RpV Fž  pV F•  pV^8X  d,   \        V 4      '       d   \        V^4      ^,          V8”  d   K5  \        P                  ! \        4      P                  V
V R,           XR,           V	\        V^4      VVVR8H  VVVR7      pVP                  V4       K—  	  K   	  EKá  	  V# )Úpipe_stagesr   é   úvllm::kFE4M3fnr   ÚhalfÚnv_bfloat16Fr   Tzvllm::kFE8M0fnuzvllm::kFloat16r   zvllm::kBFloat16z.id()ç      à?)Úscalar_tÚ	w_type_idÚ	s_type_idÚthreadsÚthread_m_blocksÚthread_n_blocksÚthread_k_blocksÚm_block_size_8ÚstagesÚgroup_blocksÚis_zp_float)r   úvllm::kU8B128)éÿÿÿÿr   é   )é   r   )ÚinsertrE   rF   ÚGROUP_BLOCKSÚTHREAD_M_BLOCKSÚTHREAD_CONFIGSr4   r   ÚmaxÚjinja2ÚTemplateÚTEMPLATEÚrender)r	   r$   ÚresultsÚstage_valuesri   Úm_blocksÚthread_configsÚk_blocksÚn_blocksrc   Úc_dtypeÚis_zp_float_listÚs_typerj   Ústage_valueÚtemplate_strs   &&              r   rI   rI   w   s]  € Ø€GØ!�?€LØ�„à×Ñ˜A˜qÔ!Ü2;×2CÒ2CÜœ/¬>÷3;Ñ.ˆ ð ˜1Ô ð 5
ô "
ñ Ø˜!Õ Ô#ð ˜1Œ} °Õ!2°cÔ!9ÙØ˜!Œ| ¨qÕ 1°RÔ 7Ùð Ð*Ô*¨|À:Ô/MÙð ˜/Ô)¨lÀ&Ô.HÙà˜/Ô)¨l¸aÔ.?Ùà! !Õ$¨Õ*ˆØ! !Õ$¨Õ*ˆØ  Õ#ˆà! VœO‘&°ˆà!˜7ÐØ�FŒ?˜{¨kÔ9Ø Ô!ð ×#Ñ# DÔ)à˜/Ô)¨l¸aÔ.?Ø%‰FØ˜OÔ+°ÀÔ0AØ&ˆFØ˜Œâð ð Ð,Ô,°ÀÔ1BØ&ˆFØ˜Œâð ð �fŒ_Ø%‰FØ�fŒ_Ø&ˆFã+ˆKÛ+�à 1Ô$Ü'¨×4Ò4Ü˜H aÓ(¨1Õ,¨xÔ7ñ
 Ü%Ÿš¬xÓ8×?Ñ?Ø$Ø)¨GÕ3Ø$ wÕ.Ø#Ü$'¨°!Ó$4Ø$,Ø$,Ø#+¨s¡?Ø&Ø!-Ø +ð  @ó  �ð —‘˜|Ö,ó1  ,ô ,ñu3;ðj €Nr   c                ó6   € V ^8„  d   QhR\         R,          RR/# rB   r   )r   s   "r   r   r   Õ   s   € ÷ ;ñ ;œt d�{ð ;°dñ ;r   c                 ó  € T ;'       g    \        \        4      P                  p \        V 4      pV P	                  R 4       F  pW!9  g   K  VP                  RR7       K  	  VP                  4        F  w  r4VP                  VRR7       K  	  R# )r   Tr   r;   r<   N)r   r   r   rL   r   r   rP   r>   )r    rK   r!   r?   Úrendereds   &    r   Úgenerate_new_kernelsr†   Õ   ss   € Ø×(Ð(”4œ“>×(Ñ(€DÜ% dÓ+€HØ—	‘	˜-Ö(ˆØÖØ�K‰K 4ˆKÖ(ñ )ð "*§¡Ö!1ÑˆØ×Ñ˜x°'ÐÖ:ó "2r   Ú__main__z--check:rn   NN)r   r   rk   r\   r   ))é€   rˆ   r[   )é@   r[   r[   )r‰   rˆ   rˆ   )rˆ   r‰   rˆ   )r_   rn   r   é   é   )r   rn   rl   r   r‹   rm   )N)rE   ÚsysÚpathlibr   rt   Ústripr3   r6   rv   rG   rr   rq   rp   rH   r   r"   r+   r8   r@   rL   rW   rI   r†   Ú__name__ÚargvÚ
SystemExitr   r   r   Ú<module>r’      sÊ   ðó Û 
Ý ã ð÷ 
�EƒGð 
ð €	ð(€ò €ò"€ò $€ò
 #€Ø
�&Ð	€õõ%õ7õ
õ÷÷õ$[÷|;ð ˆzÔØ�C—H‘H˜R•LÔ ÙÑ;×=Ò=˜ÓEÐEÀ1ÓEÐEÙÖñ r   