§
    OŠtj–  ã                   óä   — d dl Z d dlmZ d dlmZ d	d„Ze j        d„ ¦   «         Ze j        d	dej	        fd„¦   «         Z
e j        d	dej	        fd„¦   «         Ze j        d	dej	        fd„¦   «         ZdS )
é    N)ÚTensorDescriptorc                 óþ  ‡ — t          |¦  «        }t          ‰ j        ¦  «        }t          |¦  «        }|dk     r||z  }d|cxk    r	|dz
  k     sn J d¦   «         ‚|dk    s
J d¦   «         ‚t          |¦  «        |k    s
J d¦   «         ‚d}d}||         |k    s
J d	¦   «         ‚|||<   ‰                      |¦  «        }d
|z
  |gˆ fd„t	          |¦  «        D ¦   «         z   }||g|z   }	ddg|z   }
t          ‰ |	||
¦  «        S )a§  
    Given a 2- or 3-dimensional tensor T, this creates a 'ragged descriptor'
    which behaves like a concatenation (along the first axis) of subarrays
    of potentially unequal size.

    The load_ragged and store_ragged device functions can be used to read
    and write from subarrays T[slice_off : slice_off + slice_size]
    with hardware bounds-checking preventing any sort of leakage outside
    the subarray.
    r   é   zlast dimension cannot be raggedé   z<read-write ragged descriptors must have at most 3 dimensionsz1block shape must have same length as tensor shapei  ÿé   @z#number of rows may not exceed 2**30l        c                 ó:   •— g | ]}‰                      |¦  «        ‘ŒS © )Ústride)Ú.0ÚiÚTs     €úU/var/www/html/CA-Chatbot/venv/lib/python3.11/site-packages/triton/tools/ragged_tma.pyú
<listcomp>z,create_ragged_descriptor.<locals>.<listcomp>)   s#   ø€ Ð:\Ð:\Ð:\È1¸1¿8º8ÀA¹;¼;Ð:\Ð:\Ð:\ó    )ÚlistÚshapeÚlenr
   Úranger   )r   Úblock_shapeÚ
ragged_dimÚtensor_shapeÚrankÚmax_intÚbillionÚragged_strideÚ
tma_strideÚ	tma_shapeÚ	box_shapes   `          r   Úcreate_ragged_descriptorr      sQ  ø€ õ �{Ñ#Ô#€KÝ˜œ‘=”=€LÝˆ|ÑÔ€Dà�A‚~€~Ø�dÑˆ
à�
Ð%Ð%Ò%Ð%˜T A™XÒ%Ð%Ð%Ð%Ð%Ð'HÑ%Ô%Ð%Ø�1Š9ˆ9ˆ9ÐT‰9Œ9ˆ9åˆ{ÑÔ˜tÒ#Ð#Ð#Ð%XÑ#Ô#Ð#à€GØ€Gà˜
Ô# wÒ.Ð.Ð.Ð0UÑ.Ô.Ð.Ø&€L�ÑØ—H’H˜ZÑ(Ô(€Mð ˜-Ñ'¨Ð7Ð:\Ð:\Ð:\Ð:\ÕPUÐVZÑP[ÔP[Ð:\Ñ:\Ô:\Ñ\€JØ˜7Ð# lÑ2€IØ�Q�˜+Ñ%€Iå˜A˜y¨*°iÑ@Ô@Ð@r   c                 ó*   — d}||z
  |z   }| |z   }|||fS )z;
    Helper function for load_ragged and store_ragged.
    r   r	   )Ú	slice_offÚ
slice_sizeÚrowr   ÚxÚys         r   Úto_ragged_indicesr&   0   s.   € ð €GØ�*Ñ˜sÑ"€AØ�JÑ€Aà�A�qˆ=Ðr   r   c                 ó^  — t          j        t          | j        ¦  «        t          |¦  «        dz   k    d¦  «         t	          ||||         ¦  «        \  }}}|                      ||g|d|…         z   |gz   ||dz   d…         z   ¦  «        }t          j        ||j        dd…         ¦  «        }|S )zó
    Read from a subarray T[slice_off : slice_off + slice_size] with
    hardware bounds-checking, where reading outside the subarray gives zeros.

    Coords should be an appropriately-sized list of integers, just like in
    TMA.load().
    é   z*TMA must be a read-write ragged descriptorNr   )ÚtlÚstatic_assertr   r   r&   ÚloadÚreshape)	ÚTMAr!   r"   Úcoordsr   Úc0Úc1Úc2Údatas	            r   Úload_raggedr3   =   s¨   € õ Ô•S˜œ‘^”^¥s¨6¡{¤{°Q¡Ò6Ð8dÑeÔeÐeå" 9¨j¸&ÀÔ:LÑMÔM�J€BˆˆBØ�8Š8�R˜�H˜v k z kÔ2Ñ2°b°TÑ9¸FÀ:ÐPQÁ>À?À?Ô<SÑSÑTÔT€DÝŒ:�d˜DœJ q r rœNÑ+Ô+€DØ€Kr   c                 óä   — t          ||||         ¦  «        \  }}}t          j        |ddg|j        z   ¦  «        }|                      ||g|d|…         z   |gz   ||dz   d…         z   |¦  «         dS )zÿ
    Write to a subarray T[slice_off : slice_off + slice_size] with
    hardware bounds-checking, where writes outside the subarray are masked
    correctly.

    Coords should be an appropriately-sized list of integers, just like in
    TMA.store().
    r   N)r&   r)   r,   r   Ústore©	r-   r!   r"   r.   r2   r   r/   r0   r1   s	            r   Ústore_raggedr7   O   s�   € õ # 9¨j¸&ÀÔ:LÑMÔM�J€BˆˆBÝŒ:�d˜Q ˜F T¤ZÑ/Ñ0Ô0€DØ‡I‚Iˆr�2ˆh˜   Ô,Ñ,°¨tÑ3°f¸ZÈ!¹^¸_¸_Ô6MÑMÈtÑTÔTÐTÐTÐTr   c                 óä   — t          ||||         ¦  «        \  }}}t          j        |ddg|j        z   ¦  «        }|                      ||g|d|…         z   |gz   ||dz   d…         z   |¦  «         dS )a	  
    Atomic add into a subarray T[slice_off : slice_off + slice_size] with
    hardware bounds-checking, where adds outside the subarray are masked
    correctly.

    Coords should be an appropriately-sized list of integers, just like in
    TMA.atomic_add().
    r   N)r&   r)   r,   r   Ú
atomic_addr6   s	            r   Úatomic_add_raggedr:   _   s‚   € õ # 9¨j¸&ÀÔ:LÑMÔM�J€BˆˆBÝŒ:�d˜Q ˜F T¤ZÑ/Ñ0Ô0€DØ‡N‚N�B˜�8˜f [ j [Ô1Ñ1°R°DÑ8¸6À*ÈqÁ.À/À/Ô;RÑRÐTXÑYÔYÐYÐYÐYr   )r   )ÚtritonÚtriton.languageÚlanguager)   Útriton.tools.tensor_descriptorr   r   Újitr&   Ú	constexprr3   r7   r:   r	   r   r   ú<module>rA      s  ðØ €€€Ø Ð Ð Ð Ð Ð Ø ;Ð ;Ð ;Ð ;Ð ;Ð ;ð
%Að %Að %Að %AðP „ð	ð 	ñ „ð	ð „ðð ÀÄð ð ð ñ „ðð" „ðUð UÀrÄ|ð Uð Uð Uñ „ðUð „ðZð ZÈBÌLð Zð Zð Zñ „ðZð Zð Zr   