§
     ’Þj¦;  ã                   ó˜  — d Z ddlZddlZddlZddlZddlmZ ddlZddlZddl	m
Z dZ eej                             d ej        ¦   «         dz  dz  dz  ¦  «        ¦  «                             ¦   «         Zej                             d	d
¦  «        dk    Zd\  ZZZd\  ZZg d¢g d¢dœg d¢g d¢dœdœZ ej        ¦   «         Zdai Zej        d„ ¦   «         Zej        d„ ¦   «         Zej        d„ ¦   «         Z ej        d„ ¦   «         Z! ej        dg¬¦  «        dej"        dej"        fd„¦   «         Z# ej        dg¬¦  «        dej"        dej"        d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'„ Z&d(„ Z'd)„ Z(d*„ Z)d+„ Z*d,„ Z+d-„ Z,d0d.„Z-dd/l.m/Z/m0Z0  e/¦   «           e0e,¦  «        Z,dS )1aÇ  Weight-only FP8 matmul (W8A16) for block-scaled FP8 weights (shared.qtypes.block_scaled: E4M3 (N, K) with FP32 scales per 128 x 128
block or per output channel and 128-element block of K) on every CUDA GPU, with 16-bit activations: the weights are decoded exactly in
registers, each 128-element block of K accumulated in FP32 on the tensor cores, then promoted with its scales into FP32 accumulators
(the block semantics of DeepGEMM's FP8 GEMM, with exact activations instead of E4M3 ones). The weight matrix is never materialized and
the activations are never quantized: more precise than the reference dequantization + 16-bit matmul (whose dequantized weights are
rounded to the compute dtype) and than the FP8 tensor-core kernels (whose activations are quantized to E4M3).

Tensor cores: sm80+ multiply the activations in their own dtype (BF16 or FP16 MMAs); Turing (sm75) has FP16 MMAs only, so BF16
activations are first written as FP16 with one power-of-two scale per row and 128-element block (exact for BF16 values within 2^28
of their block's maximum: the scaling is exact and FP16 holds BF16's 8-bit mantissa), the scale folded into the promotion. GPUs before
sm89 cannot convert float8 in Triton: the E4M3 bytes are placed as FP16 bits (the value times 2^-8, exact, subnormals included) with
2^8 folded into the promotion, or decoded with integer operations for BF16 MMAs.

Tile configurations are static per GPU generation and tuned once per layer shape and token-count slot on a real call (outside CUDA
graph recording and compilation), cached in ~/.triton/autotune/; the kernel is not specialized on the number of tokens, so a new
resolution compiles nothing. WAN2GP_FP8_TRITON_AUTOTUNE=0 keeps the static configurations.é    N)ÚPathé   ÚWAN2GP_FP8_BLOCK_TRITON_CACHEz.tritonÚautotunezwan2gp_fp8_block_cache.jsonÚWAN2GP_FP8_TRITON_AUTOTUNEÚ1Ú0)r   é   r   )r   é   ))é    é@   é€   é   r   ©r   r   r   r   r   )é   r   r   r   r   )r   r   r   r   r   )r   )r   r   r   r   r   )r   r   r   r   r   ©r   r   r   r   r   )r   r   r   é   r   )r   r   r   r   r   )ÚsmallÚlarge))r   r   r   r   r   r   )r   r   r   r   r   ))r   r   r   r   r   )r   r   r   r   r   )r   r   r   r   r   )r   r   r   r   r   )r   r   r   r   r   )r   é   r   r   r   )Úsm75Úsm80c                 ó„   — |                       t          j        ¦  «        dz                        t          j        d¬¦  «        S )Nr   T©Úbitcast)ÚtoÚtlÚuint32Úfloat32)Úbitss    úGC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\fp8_block_triton.pyÚ_bf16_bits_to_f32r"   .   s0   € à�GŠG•B”IÑÔ "Ñ$×(Ò(­¬¸TÐ(ÑBÔBÐBó    c                 óÔ   — |                       t          j        d¬¦  «        }|dz   |dz	  dz  z   dz	  }t          j        | | k    d|¦  «                              t          j        ¦  «        S )NTr   iÿ  r   r
   iÀ  )r   r   r   ÚwhereÚuint16)Úvaluer    Úroundeds      r!   Ú_f32_to_bf16_bitsr)   3   s_   € ð �8Š8•B”I tˆ8Ñ,Ô,€DØ�f‰} ¨¡¨qÑ 0Ñ1°bÑ8€GÝŒ8�E˜U’N F¨GÑ4Ô4×7Ò7½¼	ÑBÔBÐBr#   c                 ó´  — |                       t          j        ¦  «        }|dz	  dz  }|dz  }t          j        |dk    |                      t          j        ¦  «        dz  |dz   dz  |dz  z                        t          j        d	¬
¦  «        ¦  «        }t          j        |dz  dk    t          d¦  «        |¦  «        }t          j        |dz  dk    | |¦  «        S )z>float8_e4m3fn bytes as fp32 (GPUs without float8 conversions).r   é   é   r   g      `?éx   é   é   Tr   é   Únanr   )r   r   Úint32r%   r   Úfloat)r    ÚbÚeÚmr'   s        r!   Ú_e4m3_bits_to_f32r7   ;   sÃ   € ð 	�Š•”ÑÔ€AØ	
ˆa‰�3‰€AØ	ˆC‰€AÝŒH�Q˜!’V˜QŸTšT¥"¤*Ñ-Ô-°Ñ;ÀÀCÁÈB¹ÐSTÐXZÑSZÑ>[×=_Ò=_Õ`bÔ`jÐtxÐ=_Ñ=yÔ=yÑzÔz€EÝŒH�a˜$‘h 4Ò'­¨u©¬°uÑ=Ô=€EÝŒ8�Q˜‘X !’O e V¨UÑ3Ô3Ð3r#   c                 ó    — |                       t          j        ¦  «        }|dz  dz  |dz  dz  z                        t          j        d¬¦  «        S )zÐfloat8_e4m3fn bytes as fp16 values times 2^-8 (exact: the exponent and mantissa bits move into the low FP16 exponent and high
    mantissa bits, E4M3 subnormals becoming FP16 subnormals of the same mantissa).r   r   r0   r,   Tr   )r   r   r&   Úfloat16)r    r4   s     r!   Ú_e4m3_bits_to_f16_scaledr:   F   sG   € ð 	�Š•”	ÑÔ€AØ�$‰h˜1‰_ ! d¡(¨q¡Ñ1×5Ò5µb´jÈ$Ð5ÑOÔOÐOr#   ÚM)Údo_not_specializeÚBMÚEVEN_Kc
           	      ó  — t          j        d¦  «        t          j        d¦  «        }}
|
|z  t          j        d|¦  «        z   }|dz  t          j        dd¦  «        z   }|dd…df         |k     }|	s||ddd…f         |k     z  }t          t          j        | |dd…df         |z  z   |ddd…f         |z  z   |d¬¦  «        ¦  «        }t          j        t          j        |¦  «        d¬¦  «        }|                     t           j        d¬¦  «        d	z	  d
z  }t          j	        |dk    t          j
        |d¦  «        dz
  d¦  «                             t           j        ¦  «        d	z  }|                     t           j        d¬¦  «        }t          j        |||z  z   |z   |||k     ¬¦  «         t          j        ||dd…df         |z  z   |ddd…f         z   |d|z  dd…df         z                       t           j        ¦  «        |¬¦  «         dS )zƒFP16 rows of BF16 x (bits) with one power-of-two scale per row and 128-element block, the block's maximum scaled into [2^14, 2^15).r   r
   r   N©ÚmaskÚother)ÚaxisTr   r.   éÿ   r+   é   r0   ©rA   g      ð?)r   Ú
program_idÚaranger"   ÚloadÚmaxÚabsr   r   r%   Úmaximumr   Ústorer9   )ÚxÚqÚsar;   ÚKÚsxmÚsxkÚkblocksr=   r>   Úpid_mÚkbÚrmÚrkrA   ÚvÚamaxÚexponentÚ
scale_bitsÚscales                       r!   Ú_prepare_a16_kernelr^   N   sÿ  € õ ”˜aÑ Ô ¥"¤-°Ñ"2Ô"2ˆ2€EØ	�‰•b”i  2Ñ&Ô&Ñ	&€BØ	ˆc‰•B”I˜a Ñ%Ô%Ñ	%€BØˆaˆaˆa�ˆgŒ;˜Š?€DØð (Ø�r˜$   ˜'”{ Q’Ñ'ˆÝ�"œ' ! b¨¨¨¨D¨¤k°CÑ&7Ñ"7¸"¸TÀ1À1À1¸W¼+ÈÑ:KÑ"KÐRVÐ^_Ð`Ñ`Ô`ÑaÔa€AÝŒ6•"”&˜‘)”) !Ð$Ñ$Ô$€DØ—’�œ	¨4�Ñ0Ô0°BÑ6¸$Ñ>€HÝ”˜$ š(¥B¤J¨x¸Ñ$<Ô$<¸rÑ$AÀ3ÑGÔG×JÒJÍ2Ì9ÑUÔUÐY[Ñ[€JØ�MŠM�"œ*¨dˆMÑ3Ô3€EÝ„HˆR�"�w‘,Ñ Ñ# U°°a²Ð8Ñ8Ô8Ð8Ý„HˆQ��A�A�A�t�G”˜q‘Ñ  2 d¨A¨A¨A g¤;Ñ.°°c¸E±kÀ1À1À1ÀdÀ7Ô5KÑ1K×0OÒ0OÕPRÔPZÑ0[Ô0[ÐbfÐgÑgÔgÐgÐgÐgr#   ÚWMODEÚHAS_SAÚHAS_BIASÚDECODEÚBITS_IOÚBNÚBKÚGROUP_Mc           
      ó	  — t          j        d¦  «        }t          j        ||¦  «        t          j        ||¦  «        }}||z  }||z  |z  }t          j        ||z
  |¦  «        }|||z  |z  z   } ||z  |z  }!| |z  t          j        d|¦  «        z   }"|!|z  t          j        d|¦  «        z   }#t          j        d|¦  «        }$|"|k     |#|k     }&}%| |"dd…df         |
z  z   |$ddd…f         |z  z   }'||#ddd…f         |z  z   |$dd…df         |z  z   }(|dk    r||#dz  |	z  z   })n||#|	z  z   })t          j        ||ft           j        ¬¦  «        }*d|z  }+t          d|	¦  «        D �]},t          j        ||ft           j        ¬¦  «        }-t          j        |+¦  «        D �]w}.|rBt          j	        |'|%dd…df         d¬¦  «        }/t          j	        |(|&ddd…f         ¬¦  «        }0nm|$||,|+z  |.z   |z  z
  k     }1t          j	        |'|%dd…df         |1ddd…f         z  d¬¦  «        }/t          j	        |(|&ddd…f         |1dd…df         z  ¬¦  «        }0|dk    rt          |0¦  «        }0ne|dk    r(t          |0¦  «                             |/j        ¦  «        }0n7|0                     t           j        ¦  «                             |/j        ¦  «        }0|s t          j        |1dd…df         |0d	¦  «        }0t          j        |/|0|-¦  «        }-|'||z  z  }'|(||z  z  }(�Œyt          j	        |)|,z   |&d	¬¦  «        ddd…f         }2|dk    r|2d
z  }2|r-|2t          j	        ||"|	z  z   |,z   |%d	¬¦  «        dd…df         z  }2|*|-|2z  z  }*�Œ|rZt          j	        ||#z   |&d¬¦  «        ddd…f         }3|rt!          |3¦  «        }3n|3                     t           j        ¦  «        }3|*|3z  }*||"dd…df         |z  z   |#ddd…f         |z  z   }4|r=t          j        |4t%          |*¦  «        |%dd…df         |&ddd…f         z  ¬¦  «         dS t          j        |4|*                     |j        j        ¦  «        |%dd…df         |&ddd…f         z  ¬¦  «         dS )aØ  out = x @ (w * scales).T + bias, x (M, K) FP16 or BF16 (BITS_IO: the BF16 bias and output are uint16 bits, for GPUs without BF16
    conversions, whose x is FP16), w (N, K) E4M3 bytes (uint8 when DECODE != 0),
    sa (M, kblocks) FP32 row-block scales or none, sb (ceil(N / 128), kblocks) block scales or (N, kblocks) group scales: each 128-element
    block of K is accumulated from zero in FP32 and promoted with sa[row] * sb[block] (times 2^8 for the FP16-bits decode).r   Nr   )Údtyper@   rF   r   r
   g        g      p@)r   rG   ÚcdivÚminimumrH   Úzerosr   ÚrangeÚstatic_rangerI   r:   r7   r   rh   r%   Údotr"   rM   r)   Ú
element_ty)5rN   ÚwrP   ÚsbÚbiasÚoutr;   ÚNrQ   rT   rR   rS   ÚswnÚswkÚsomÚsonr_   r`   ra   rb   rc   r>   r=   rd   re   rf   ÚpidÚtiles_mÚtiles_nÚgroupÚfirst_mÚsize_mrU   Úpid_nrW   ÚrnrX   Úm_maskÚn_maskÚa_ptrÚb_ptrÚsb_ptrÚaccfÚKPBrV   ÚaccÚjÚar4   Úk_maskÚsÚ
bias_valueÚout_ptrs5                                                        r!   Ú_block_kernelr�   `   s  € õ Œ-˜Ñ
Ô
€CÝ”w˜q "‘~”~¥r¤w¨q°"¡~¤~ˆW€GØ�gÑ€EØ�e‰|˜wÑ&€GÝŒZ˜ 'Ñ)¨7Ñ3Ô3€FØ�s˜U‘{ fÑ,Ñ,€EØ�5‰[˜VÑ#€EØ	�‰•b”i  2Ñ&Ô&Ñ	&€BØ	�‰•b”i  2Ñ&Ô&Ñ	&€BÝ	Œ�1�bÑ	Ô	€BØ˜!’V˜R !šVˆF€FØ��1�1�1�d�7”˜cÑ!Ñ! B t¨Q¨Q¨Q w¤K°#Ñ$5Ñ5€EØ��4˜˜˜�7”˜cÑ!Ñ! B q q q¨$ w¤K°#Ñ$5Ñ5€EØ�‚z€zØ�r˜S‘y GÑ+Ñ+ˆˆà�b˜7‘lÑ"ˆÝŒ8�R˜�H¥B¤JÐ/Ñ/Ô/€DØ˜r™	€CÝ�A�wÑÔð ñ ˆÝŒh˜˜B�x¥r¤zÐ2Ñ2Ô2ˆÝ” Ñ%Ô%ð 	ñ 	ˆAØð KÝ”G˜E¨¨q¨q¨q°$¨w¬¸qÐAÑAÔA�Ý”G˜E¨¨t°Q°Q°Q¨w¬Ð8Ñ8Ô8��à˜a 2¨¡8¨a¡<°2Ñ"5Ñ5Ò5�Ý”G˜E¨¨q¨q¨q°$¨w¬¸&ÀÀqÀqÀqÀ¼/Ñ(IÐQRÐSÑSÔS�Ý”G˜E¨¨t°Q°Q°Q¨w¬¸&ÀÀÀÀDÀ¼/Ñ(IÐJÑJÔJ�Ø˜Š{ˆ{Ý,¨QÑ/Ô/��Ø˜1’�Ý% aÑ(Ô(×+Ò+¨A¬GÑ4Ô4��à—D’D�œÑ$Ô$×'Ò'¨¬Ñ0Ô0�Øð 6Ý”H˜V A A A t Gœ_¨a°Ñ5Ô5�Ý”&˜˜A˜sÑ#Ô#ˆCØ�R˜#‘XÑˆEØ�R˜#‘XÑˆE‰EÝŒG�F˜R‘K f°CÐ8Ñ8Ô8¸¸q¸q¸q¸ÔAˆØ�QŠ;ˆ;Ø�E‘	ˆAØð 	UØ•B”G˜B  g¡Ñ-°Ñ2¸ÀsÐKÑKÔKÈAÈAÈAÈtÈGÔTÑTˆAØ��a‘‰ˆ‰Øð Ý”W˜T B™Y¨V¸1Ð=Ñ=Ô=¸dÀAÀAÀA¸gÔFˆ
Øð 	3Ý*¨:Ñ6Ô6ˆJˆJà#Ÿš¥r¤zÑ2Ô2ˆJØ�
ÑˆØ�B�q�q�q˜$�w”K #Ñ%Ñ%¨¨4°°°¨7¬°cÑ(9Ñ9€GØð aÝ
Œ�Õ+¨DÑ1Ô1¸¸q¸q¸qÀ$¸w¼È&ÐQUÐWXÐWXÐWXÐQXÌ/Ñ8YÐZÑZÔZÐZÐZÐZå
Œ�˜$Ÿ'š' #¤)Ô"6Ñ7Ô7¸fÀQÀQÀQÈÀW¼oÐPVÐW[Ð]^Ð]^Ð]^ÐW^ÔP_Ñ>_Ð`Ñ`Ô`Ð`Ð`Ð`r#   c                 ój  — | dz   }|                       ¦   «         dk    r&t          | j        ¦  «        | dz   |fk    rt          S |                       ¦   «         dk    r!t          | j        ¦  «        ||fk    rt          S t          dt          | j        ¦  «        › d| dz   › d|› d|› d|› d�¦  «        ‚)zTWS_BLOCK for (ceil(N / 128), ceil(K / 128)) scales, WS_GROUP for (N, ceil(K / 128)).r   r   zFP8 weight scales of shape z match neither (z, z) blocks nor (z) channel blocks)ÚdimÚtupleÚshapeÚWS_BLOCKÚWS_GROUPÚ
ValueError)r]   ÚnÚkrT   s       r!   Ú
scale_moder™   ¤   sý   € à��c‘	ˆl€GØ‡y‚y�{„{�aÒÐ�E %¤+Ñ.Ô.°a°R¸3±Y°<ÀÐ2IÒIÐIÝˆØ‡y‚y�{„{�aÒÐ�E %¤+Ñ.Ô.°1°g°,Ò>Ð>ÝˆÝ
ð  ]µ5¸¼Ñ3EÔ3Eð  ]ð  ]ÐZ[ÐY[Ð_bÑYbÐWcð  ]ð  ]Ðgnð  ]ð  ]Ð~ð  ]ð  ]ð  DKð  ]ð  ]ð  ]ñ  ^ô  ^ð  ^r#   c                 óÎ  — | j         \  }}| dz   }t          j        ||ft          j        | j        ¬¦  «        }t          j        ||ft          j        | j        ¬¦  «        }d}t          t          j        ||¦  «        |f         |  	                    t          j
        ¦  «        |||||                      d¦  «        |                      d¦  «        |||dz  dk    d¬¦  «         ||fS )zÄ(FP16 copy of BF16 x (M, K), FP32 scales (M, ceil(K / 128))): x = q * scale per row and 128-element block, exactly for values within
    2^28 of their block's maximum (Turing's FP16 tensor cores).r   ©rh   Údevicer   r   r
   r   )Ú	num_warps)r“   ÚtorchÚemptyr9   rœ   r   r^   Útritonri   Úviewr&   Ústride)rN   r6   r˜   rT   rO   rP   Úbms          r!   Úprepare_a16r¤   ®   s  € ð Œ7�D€A€qØ��c‘	ˆl€GÝŒ�Q˜�F¥%¤-¸¼ÐAÑAÔA€AÝ	Œ�a˜�\­¬¸q¼xÐ	HÑ	HÔ	H€BØ	€BÝ�œ Q¨Ñ+Ô+¨WÐ5Ô6°q·v²v½e¼lÑ7KÔ7KÈQÐPRÐTUÐWXÐZ[×ZbÒZbÐcdÑZeÔZeÐgh×goÒgoÐpqÑgrÔgrÐt{Ð}ð  BCð  FIñ  BIð  MNò  BNð  Z[ð  \ñ  \ô  \ð  \Øˆbˆ5€Lr#   c                 óP   — t           j                             | ¦  «        dk     rdndS )N©r   r   r   r   )rž   ÚcudaÚget_device_capability)rœ   s    r!   Ú_generationr©   º   s&   € Ý”Z×5Ò5°fÑ=Ô=ÀÒFÐFˆ6ˆ6ÈFÐRr#   c                 óh  — |\  }}	}
}}| j         \  }}|j         d         }t          j                             | j        ¦  «        }|j        t          j        k    o|dk     }| j        t          j        k    r9|dk     s$t          j	         
                    dd¦  «        dk    rt          nt          }n8|dk    r+t          j	         
                    dd¦  «        dk    rt          nt          }|t          k    r|n|                     t          j        ¦  «        }|r@|                     t          j        ¦  «        }|�|                     t          j        ¦  «        }| dz   }t#          t%          j        ||¦  «        t%          j        ||	¦  «        z  f         | ||||||||||                      d¦  «        |                      d¦  «        |                     d¦  «        |                     d¦  «        |                     d¦  «        |                     d¦  «        ||d u|d u|||dz  dk    ||	|
d	||¬
¦  «         d S )Nr   r¦   )r   é	   ÚWAN2GP_FP8_TRITON_DECODEÚ r    r   r
   r   )r�   Ú
num_stages)r“   rž   r§   r¨   rœ   rh   Úbfloat16r9   ÚosÚenvironÚgetÚDECODE_BITS_F16ÚDECODE_NATIVEÚDECODE_BITS_F32r¡   Úuint8r&   r�   r    ri   r¢   )rŠ   rP   Úweightr]   Úwmoderr   rs   Úconfigr£   ÚbnÚbkÚwarpsÚstagesr6   r˜   r—   Ú
capabilityÚbits_ioÚdecoderp   rT   s                        r!   Ú_launchrÁ   ¾   s=  € Ø &Ñ€BˆˆB��vØŒ7�D€A€qØŒ�QŒ€AÝ”×1Ò1°!´(Ñ;Ô;€JØŒi�5œ>Ò)ÐA¨j¸6Ò.A€GØ„w•%”-ÒÐØ$.°Ò$7Ð$7½2¼:¿>º>ÐJdÐfhÑ;iÔ;iÐmsÒ;sÐ;s•�õ  zGˆˆà",°Ò"6Ð"6½2¼:¿>º>ÐJdÐfhÑ;iÔ;iÐmsÒ;sÐ;s•�õ  zIˆØ�MÒ)Ð)ˆˆ¨v¯{ª{½5¼;Ñ/GÔ/G€AØð +Ø�hŠh•u”|Ñ$Ô$ˆØÐØ—9’9�Uœ\Ñ*Ô*ˆDØ��c‘	ˆl€GÝ•6”;˜q "Ñ%Ô%­¬°A°rÑ(:Ô(:Ñ:Ð<Ô=Ø	ˆ1ˆb�%˜˜s A q¨!¨W°a·h²h¸q±k´kÀ1Ç8Â8ÈAÁ;Ä;ÐPQ×PXÒPXÐYZÑP[ÔP[Ð]^×]eÒ]eÐfgÑ]hÔ]hÐjm×jtÒjtÐuvÑjwÔjwÐy|÷  zDò  zDð  EFñ  zGô  zGØˆr˜ˆ~˜t¨4Ð/°¸À!ÀcÁ'ÈQÂ,ÐPRÐTVÐXZÐ\]Ðinð  |BðCñ Cô Cð Cð Cð Cr#   c                  ó
  — t           €v	 t          j        t                               d¬¦  «        ¦  «        } |                      d¦  «        t          k    r| d         ni a n# t          t          t          f$ r i a Y nw xY wt           S )Núutf-8©ÚencodingÚversionÚconfigs)
Ú_tunedÚjsonÚloadsÚ_CACHE_PATHÚ	read_textr²   Ú_CACHE_VERSIONÚOSErrorr–   ÚKeyError)Úpayloads    r!   Ú_load_cacherÑ   Ó   s„   € å€~ð	Ý”j¥×!6Ò!6ÀÐ!6Ñ!HÔ!HÑIÔIˆGØ+2¯;ª;°yÑ+AÔ+AÅ^Ò+SÐ+S�W˜YÔ'Ð'ÐY[ˆFˆFøÝ�¥XÐ.ð 	ð 	ð 	ØˆFˆFˆFð	øøøå€Ms   ‰AA ÁA;Á:A;c                  ó  — 	 t           j                             dd¬¦  «         t                                t	          j        t          t          dœd¬¦  «        d¬¦  «         d S # t          $ r} t          d| › d	�¦  «         Y d } ~ d S d } ~ ww xY w)
NT)ÚparentsÚexist_ok)rÆ   rÇ   r
   )ÚindentrÃ   rÄ   z-[FP8] Triton block tuning results not saved (z).)
rË   ÚparentÚmkdirÚ
write_textrÉ   ÚdumpsrÍ   rÈ   rÎ   Úprint)Úexcs    r!   Ú_save_cacherÜ   Þ   s¬   € ðGÝÔ× Ò ¨¸Ð Ñ=Ô=Ð=Ý×Ò�tœzµnÕQWÐ*XÐ*XÐabÐcÑcÔcÐnuÐÑvÔvÐvÐvÐvøÝð Gð Gð GÝÐE¸cÐEÐEÐEÑFÔFÐFÐFÐFÐFÐFÐFÐFøøøøðGøøøs   ‚AA" Á"
B
Á,BÂB
c                 óv  —  | ¦   «          t           j                             d¬¦  «        t           j                             d¬¦  «        }}|                     ¦   «          t	          |¦  «        D ]} | ¦   «          Œ|                     ¦   «          |                     ¦   «          |                     |¦  «        |z  S )NT)Úenable_timing)rž   r§   ÚEventÚrecordrl   ÚsynchronizeÚelapsed_time)ÚlaunchÚrepsÚstartÚendÚ_s        r!   Ú_timerè   æ   s¢   € Ø
€F�H„H€HÝ”×!Ò!°Ð!Ñ5Ô5µu´z×7GÒ7GÐVZÐ7GÑ7[Ô7[ˆ3€EØ	‡L‚L�N„N€NÝ�4‰[Œ[ð ð ˆØˆ‰ŒˆˆØ‡J‚J�L„L€LØ‡O‚OÑÔÐØ×Ò˜cÑ"Ô" TÑ)Ð)r#   c	           	      óª  ‡‡— t          |j        d         d¦  «        }	|d|	…         |€dn	|d|	…         |||||d|	…         fŠi }
|D ]@Š	 t          ˆˆfd„d¦  «        |
‰<   Œ# t          j        j        j        t          f$ r Y Œ=w xY wt          |
¦  «        D ]-Št          |
‰         t          ˆˆfd„d¦  «        ¦  «        |
‰<   Œ.t          |
|
j	        ¬¦  «        }|d         |
v r#|
|         dz  |
|d                  k    r|d         }t          5  t          |¦  «        t          ¦   «         | <   t          ¦   «          ddd¦  «         n# 1 swxY w Y   |S )	zÂThe slot's configuration, timed on at most 2048 rows of the call (its own input and output), kept when at least 2% faster than
    the static one; configurations this GPU cannot run are skipped.r   i   Nc                  ó    •— t          g ‰ ¢‰‘R Ž S ©N©rÁ   ©Úbenchr¹   s   €€r!   ú<lambda>z_tune.<locals>.<lambda>ù   s   ø€ ­GÐ,C°UÐ,C¸FÐ,CÐ,CÐ,C€ r#   r   c                  ó    •— t          g ‰ ¢‰‘R Ž S rë   rì   rí   s   €€r!   rï   z_tune.<locals>.<lambda>ý   s   ø€ ½WÐ=TÀeÐ=TÈVÐ=TÐ=TÐ=T€ r#   )ÚkeygR¸…ëQð?)Úminr“   rè   r    ÚruntimeÚerrorsÚOutOfResourcesÚRuntimeErrorÚlistr²   Ú_lockrÑ   rÜ   )rñ   rŠ   rP   r·   r]   r¸   rr   rs   Ú
candidatesÚrowsÚtimingsÚbestrî   r¹   s               @@r!   Ú_tunerý   ñ   sÌ  øø€ õ ˆqŒw�qŒz˜4Ñ Ô €DØˆu�ˆuŒX˜r˜z�t�t¨r°%°4°%¬y¸&À%ÈÐPTÐVYÐZ_Ð[_ÐZ_ÔV`Ða€EØ€GØð ð ˆð	Ý#Ð$CÐ$CÐ$CÐ$CÐ$CÀQÑGÔGˆG�F‰OˆOøÝ”Ô%Ô4µlÐCð 	ð 	ð 	ØˆHð	øøøå�w‘-”-ð Zð ZˆÝ˜g fœo­uÐ5TÐ5TÐ5TÐ5TÐ5TÐVWÑ/XÔ/XÑYÔYˆ�‰ˆÝˆw˜GœKÐ(Ñ(Ô(€DØ�!„}˜ÐÐ G¨D¤M°DÑ$8¸7À:ÈaÄ=Ô;QÒ$QÐ$QØ˜!Œ}ˆÝ	ð ð Ý! $™ZœZ�‰Œ�cÑÝ‰Œˆðð ð ñ ô ð ð ð ð ð ð øøøð ð ð ð ð €Ks$   Á
A"Á" BÂBÄ-EÅEÅEc                 ó  — | j         \  }}|j         d         }t          |||¦  «        }|j        t          j        k    r|                     ¦   «         r|n%|                     ¦   «                              ¦   «         }|€#t          j        ||f| j        | j	        ¬¦  «        n|}|�*|j        | j        k    r| 
                    | j        ¦  «        }t          | j	        ¦  «        }	| j        t          j        k    r|	dk    rt          | ¦  «        \  }
}n| d}}
|dk    rdnd}t                               | j	        j        ¦  «        }|€8t          j                             | j	        ¦  «        x}t          | j	        j        <   |› d|› d|› d|› d|› dt'          | j        ¦  «        d	d…         › dt)          |du¦  «        › �}t+          ¦   «                              |¦  «        }|€…t,          |	         |         }t.          oI|dk    s|d
k    o=t          j                             ¦   «          ot          j                             ¦   «          }|rt7          ||
|||||||¦	  «	        n|d         }t9          |
||||||t;          |¦  «        ¦  «         |S )zŸx (M, K) BF16/FP16 @ (weight (N, K) E4M3 * scale per 128 x 128 block or per channel and 128-element block of K).T + bias, in x's
    dtype; K a multiple of 16.r   Nr›   r   r   r   r   ú|é   i   )r“   r™   rh   rž   r   Úis_contiguousr3   Ú
contiguousrŸ   rœ   r   r©   r¯   r¤   Ú_device_namesr²   Úindexr§   Úget_device_nameÚstrÚintrÑ   Ú_CANDIDATESÚ_TUNINGÚcompilerÚis_compilingÚis_current_stream_capturingrý   rÁ   r’   )rN   r·   r]   rr   rs   r6   r˜   r—   r¸   Ú
generationrŠ   rP   Úslotrœ   rñ   r¹   rù   Útunes                     r!   Úlinearr    s‹  € ð Œ7�D€A€qØŒ�QŒ€AÝ�u˜a Ñ#Ô#€EØ”[¥E¤MÒ1Ð1°e×6IÒ6IÑ6KÔ6KÐ1ˆEˆEÐQV×Q\ÒQ\ÑQ^ÔQ^×QiÒQiÑQkÔQk€EØADÀ�%Œ+�q˜!�f A¤G°A´HÐ
=Ñ
=Ô
=Ð
=ÐRU€CØÐ˜DœJ¨!¬'Ò1Ð1Ø�wŠw�q”wÑÔˆÝ˜QœXÑ&Ô&€JØ„w•%”.Ò Ð  Z°6Ò%9Ð%9Ý˜A‘”‰ˆˆ2ˆ2à�4ˆ2ˆØ˜3’h�hˆ7ˆ7 G€DÝ×Ò˜qœxœ~Ñ.Ô.€FØ€~Ý16´×1KÒ1KÈAÌHÑ1UÔ1UÐUˆ•˜qœxœ~Ñ.ØÐ
WÐ
W�dÐ
WÐ
W˜QÐ
WÐ
W Ð
WÐ
W UÐ
WÐ
W­S°´©\¬\¸!¸"¸"Ô-=Ð
WÐ
WÅÀDÐPTÐDTÑ@UÔ@UÐ
WÐ
W€CÝ‰]Œ]×Ò˜sÑ#Ô#€FØ€~Ý  Ô,¨TÔ2ˆ
Ýð  Q˜D GšOÐ8¨q°Dªyð  QÅ%Ä.×B]ÒB]ÑB_ÔB_Ð>_ð  QÕhmÔhr÷  iOò  iOñ  iQô  iQð  eQˆØSWÐj•�s˜A˜r 6¨5°%¸¸sÀJÑOÔOÐOÐ]gÐhiÔ]jˆÝˆAˆr�6˜5 %¨¨sµE¸&±M´MÑBÔBÐBØ€Jr#   )Ú!install_triton_compilation_loggerÚlog_compilation_tuning)NN)1Ú__doc__rÉ   r°   Ú	threadingÚtimeÚpathlibr   rž   r    Útriton.languageÚlanguager   rÍ   r±   r²   ÚhomeÚ
expanduserrË   r	  r´   rµ   r³   r”   r•   r  ÚLockrø   rÈ   r  Újitr"   r)   r7   r:   Ú	constexprr^   r�   r™   r¤   r©   rÁ   rÑ   rÜ   rè   rý   r  Ú%shared.kernels.triton_compilation_logr  r  © r#   r!   ú<module>r      sÞ  ðð^ð ^ð  €€€Ø 	€	€	€	Ø Ð Ð Ð Ø €€€Ø Ð Ð Ð Ð Ð à €€€Ø €€€Ø Ð Ð Ð Ð Ð à€Øˆd�2”:—>’>Ð"AÀ9À4Ä9Á;Ä;ÐQZÑCZÐ]gÑCgð  kHñ  DHñ  Iô  Iñ  Jô  J÷  Uò  Uñ  Wô  W€Ø
Œ*�.Š.Ð5°sÑ
;Ô
;¸sÒ
B€Ø29Ñ /€ˆ ØÑ €ˆ(ð lÐkÐkð Xð  Xð  XðYð YàUÐUÐUð ]ð  ]ð  ]ð^ð ^ðð €ð 	ˆ	ŒÑÔ€Ø	€Ø€ð „ðCð Cñ „ðCð „ðCð Cñ „ðCð „ð4ð 4ñ „ð4ð „ðPð Pñ „ðPð €„˜s˜eÐ$Ñ$Ô$ðh¸r¼|ð hÐUWÔUað hð hð hñ %Ô$ðhð" €„˜s˜eÐ$Ñ$Ô$ð@aØœð@aØ/1¬|ð@aØGIÄ|ð@aØ]_Ô]ið@aØtvô  uAð@aàœ,ð@aà,.¬Lð@aà>@¼lð@aàPRÔP\ð@aàgiÔgsð@að @að @añ %Ô$ð@aðF^ð ^ð ^ð	ð 	ð 	ðSð Sð SðCð Cð Cð*ð ð ðGð Gð Gð*ð *ð *ðð ð ð,ð ð ð ð: lÐ kÐ kÐ kÐ kÐ kÐ kÐ kØ !Ð !Ñ #Ô #Ð #ØÐ˜uÑ%Ô%€€€r#   