§
     ’Þj"  ã                   ó2  — d Z ddlZddlmZ ddgZ ej        ee¬¦  «        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dl	Z	ddl
mZ  eddd¦  «        Zg d¢ZdZdZd„ ZddlmZ  e¦   «          dS )aå  GEMM for MMGP row views: output columns [start, start + N) of one projection of a fused weight that interleaves its projections
head by head, rows (heads, slots, head_dim): out[:, n] = x @ W[row(n)].T (+ b[row(n)]), row(n) = (n // HEAD) * HEAD_STRIDE + slot offset
+ n % HEAD. A tile of output columns lies within one head (BLOCK_N divides HEAD), so its weight rows are contiguous; tiles sharing rows
of x run together for L2 reuse. FP32 accumulation, one rounding to the output dtype.é    NÚMÚN)Údo_not_specializeÚdo_not_specialize_on_alignmentÚHEADÚHEAD_STRIDEÚHAS_BIASÚEVEN_KÚBLOCK_MÚBLOCK_NÚBLOCK_KÚGROUP_Mc           
      ó�  — t          j        d¦  «        }t          j        ||¦  «        ||z  }}||z  }||z  |z  }t          j        ||z
  |¦  «        }|||z  |z  z   }||z  |z  }|||z  z   }||z  |z  |z   ||z  z   }||z  t          j        d|¦  «        z                        t           j        ¦  «        }t          j        d|¦  «        }t          j        d|¦  «        } | |d d …d f         |	z  z   | d d d …f         z   }!|||z                        t           j        ¦  «        d d d …f         |
z  z   | d d …d f         z   }"|d d …d f         |k     }#t          j        ||ft           j        ¬¦  «        }$t          d||¦  «        D ]£}%|r,t          j
        |!|#d¬¦  «        }&t          j
        |"¦  «        }'nSt          j
        |!|#|%| d d d …f         z   |k     z  d¬¦  «        }&t          j
        |"|%| d d …d f         z   |k     d¬¦  «        }'t          j        |&|'|$¦  «        }$|!|z  }!|"|z  }"Œ¤|rD|$t          j
        ||z   |z   ¦  «                             t           j        ¦  «        d d d …f         z  }$||d d …d f         |z  z   ||z  |z   d d d …f         z   }(t          j        |(|$                     |j        j        ¦  «        |#¬¦  «         d S )Nr   )Údtypeg        )ÚmaskÚother)r   )ÚtlÚ
program_idÚcdivÚminimumÚarangeÚtoÚint64ÚzerosÚfloat32ÚrangeÚloadÚdotÚstorer   Ú
element_ty))ÚXÚWÚBÚOUTr   r   ÚKÚSTARTÚSLOT_OFFSETÚ	STRIDE_XMÚ	STRIDE_WNÚ	STRIDE_OMr   r   r	   r
   r   r   r   r   ÚpidÚtiles_mÚtiles_nÚ
group_sizeÚfirst_mÚrows_in_groupÚpid_mÚpid_nÚcolumnÚ
weight_rowÚrmÚrnÚrkÚx_ptrsÚw_ptrsÚ	row_validÚaccÚkÚaÚbÚout_ptrss)                                            úDC:\Users\johnn\OneDrive\Documents\AI\shared\kernels\rowmap_triton.pyÚ_rowmap_gemmrA      s  € õ Œ-˜Ñ
Ô
€CÝ”w˜q 'Ñ*Ô*¨A°©LˆW€GØ˜7Ñ"€JØ�ZÑ 'Ñ)€GÝ”J˜w¨Ñ0°'Ñ:Ô:€MØ�s˜ZÑ'¨=Ñ8Ñ8€EØ�:Ñ -Ñ/€EØ�U˜W‘_Ñ$€FØ˜D‘. KÑ/°+Ñ=ÀÈÁÑM€JØ
�'‰/�BœI a¨Ñ1Ô1Ñ
1×	5Ò	5µb´hÑ	?Ô	?€BÝ	Œ�1�gÑ	Ô	€BÝ	Œ�1�gÑ	Ô	€BØ��A�A�A�t�G”˜yÑ(Ñ(¨2¨d°A°A°A¨g¬;Ñ6€FØ�*˜r‘/×%Ò%¥b¤hÑ/Ô/°°a°a°a°Ô8¸9ÑDÑDÀrÈ!È!È!ÈTÈ'Ä{ÑR€FØ�1�1�1�d�7”˜a’€IÝ
Œ(�G˜WÐ%­R¬ZÐ
8Ñ
8Ô
8€CÝ�1�a˜Ñ!Ô!ð 	ð 	ˆØð 	EÝ”˜ Y°cÐ:Ñ:Ô:ˆAÝ”˜‘”ˆAˆAå”˜ Y°!°b¸¸q¸q¸q¸´k±/ÀAÒ2EÑ%FÈcÐRÑRÔRˆAÝ”˜ Q¨¨A¨A¨A¨t¨G¬¡_°qÒ%8ÀÐDÑDÔDˆAÝŒf�Q˜˜3ÑÔˆØ�'ÑˆØ�'ÑˆˆØð DØ�rŒw�q˜:‘~¨Ñ*Ñ+Ô+×.Ò.­r¬zÑ:Ô:¸4ÀÀÀ¸7ÔCÑCˆØ�R˜˜˜˜4˜”[ 9Ñ,Ñ,°¸±À"Ñ0DÀdÈAÈAÈAÀgÔ/NÑN€HÝ„HˆX�s—v’v˜cœiÔ2Ñ3Ô3¸)ÐDÑDÔDÐDÐDÐDó    )ÚLaunchTunerÚwan2gp_rowmap_linearé   z	Row Views))é€   rF   é    é   é   é   )rF   rF   rG   rH   rI   rI   )rF   rF   é@   rH   rI   rJ   )rF   rF   rK   rH   rH   rJ   )é   rF   rG   rH   rH   rJ   )rF   rK   rK   rH   rI   rI   )rK   rF   rK   rH   rI   rI   )rF   rK   rG   rH   rI   rI   i   i @  c                 óx  ‡ ‡‡‡‡‡‡‡‡‡— ‰ j         \  }Š|‰z
  Št          j        |‰f‰ j        ‰ j        ¬¦  «        }	ˆfd„t
          D ¦   «         }
|dk    r|	S ˆˆˆˆˆˆˆˆfd„Šˆˆˆ fd„}t                               d‰ j        › d‰› d‰› d	‰› d
‰du› �	‰ j        |
||t          k    ¦  «        } ‰‰ |	|¦  «         |	S )zËColumns [start, stop) (multiples of head_dim) of projection `slot` of `weight` (heads * slots * head_dim, in_features), for x
    (tokens, in_features) with contiguous rows; weight and bias in x's dtype.©r   Údevicec                 ó2   •— g | ]}‰|d          z  dk    ¯|‘ŒS )é   r   © )Ú.0ÚconfigÚhead_dims     €r@   ú
<listcomp>zrows.<locals>.<listcomp>B   s,   ø€ ÐMÐMÐM˜V°8¸fÀQ¼iÑ3GÈ1Ò3LÐ3L�&Ð3LÐ3LÐ3LrB   r   c                 ód  •— |\  }}}}}}t          j        | j        d         |¦  «        ‰|z  z  f}	t          |	         | ‰‰
€‰n‰
|| j        d         ‰‰‰‰‰z  |                      d¦  «        ‰                     d¦  «        |                     d¦  «        f‰‰‰z  ‰
d u‰|z  dk    ||||||dœ
Ž d S )Nr   )
r   r   r	   r
   r   r   r   r   Ú	num_warpsÚ
num_stages)Útritonr   ÚshaperA   Ústride)ÚsourceÚtargetrT   Úblock_mÚblock_nÚblock_kÚgroup_mÚwarpsÚstagesÚgridÚbiasÚcolumnsrU   ÚslotÚslotsÚstartÚweightÚwidths             €€€€€€€€r@   Úlaunchzrows.<locals>.launchF   s  ø€ Ø<BÑ9ˆ�˜' 7¨E°6Ý”˜FœL¨œO¨WÑ5Ô5¸ÀGÑ9KÑLÐNˆÝ�TÔ˜6 6°T°\¨6¨6ÀtÈVÐU[ÔUaÐbcÔUdÐfmÐotÐv{ð  ~Bð  EMñ  ~MØ!Ÿ=š=¨Ñ+Ô+¨V¯]ª]¸1Ñ-=Ô-=¸v¿}º}ÈQÑ?OÔ?Oð	PØV^ÐlqÐt|Ñl|Ø$(°Ð$4¸UÀW¹_ÐPQÒ=QÐ[bÐlsð  ~EØ#*°eÈð	Pð 	Pð 	Pð 	Pð 	Pð 	PrB   c                  ó�   •‡ ‡— ‰d t           …         Š t          j        ‰ j        d         ‰f‰j        ‰j        ¬¦  «        Šˆˆ ˆfd„S )Nr   rN   c                 ó   •—  ‰‰‰| ¦  «        S )NrR   )rT   rm   ÚsampleÚscratchs    €€€r@   ú<lambda>z/rows.<locals>.prepare_scratch.<locals>.<lambda>Q   s   ø€ ˜f˜f V¨W°fÑ=Ô=€ rB   )Ú_SCRATCH_ROWSÚtorchÚemptyr[   r   rO   )rp   rq   rg   rm   Úxs   @@€€€r@   Úprepare_scratchzrows.<locals>.prepare_scratchN   sO   øøø€ Ø�>•M�>Ô"ˆÝ”+˜vœ|¨Aœ°Ð8ÀÄÐPQÔPXÐYÑYÔYˆØ=Ð=Ð=Ð=Ð=Ð=Ð=rB   Úrowmapz|Kú|rv   z|biasN)	r[   rt   ru   r   rO   Ú_CONFIGSÚ_TUNERrT   Ú
_TUNE_ROWS)rv   rk   rf   rh   ri   rU   rj   ÚstopÚtokensÚoutÚ
candidatesrw   rT   rg   rm   rl   s   ```````      @@@r@   Úrowsr�   <   sL  øøøøøøøøøø€ ð ”G�M€FˆEØ�U‰l€GÝ
Œ+�v˜wÐ'¨q¬w¸q¼xÐ
HÑ
HÔ
H€CØMÐMÐMÐM¥xÐMÑMÔM€JØ�‚{€{Øˆ
ðPð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð Pð>ð >ð >ð >ð >ð >ð >õ
 �]Š]˜8¨¬Ð%dÐ%d°5Ð%dÐ%d¸8Ð%dÐ%dÀeÐ%dÐ%dÐRVÐ^bÐRbÐ%dÐ%dÐfgÔfnÐpzð  }LØ!¥ZÒ/ñ1ô 1€Fà
€Fˆ1ˆc�6ÑÔÐØ€JrB   )Ú!install_triton_compilation_logger)Ú__doc__rZ   Útriton.languageÚlanguager   Ú_RUNTIMEÚjitÚ	constexprrA   rt   Úshared.kernels.triton_tunerrC   r{   rz   r|   rs   r�   Ú%shared.kernels.triton_compilation_logr‚   rR   rB   r@   ú<module>r‹      s_  ððXð Xð €€€Ø Ð Ð Ð Ð Ð à�ˆ:€ð €„˜hÀxÐPÑPÔPð EÐceÔcoð  Eð  Aô  Kð  EØœ<ð EØ13´ð EØHJÌð EØ_aÔ_kð EØvxô  wCð Eàœ,ð Eð  Eð  Eñ QÔPð EðH €€€Ø 3Ð 3Ð 3Ð 3Ð 3Ð 3à	ˆÐ+¨Q°Ñ	<Ô	<€ðTð Tð T€à€
Ø€ðð ð ð: TÐ SÐ SÐ SÐ SÐ SØ !Ð !Ñ #Ô #Ð #Ð #Ð #rB   