Ë
    ÛéñiÆ  ã                   ó  — 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y)
é    N)ÚTensorDescriptorc                 óÜ  — t        |«      }t        | j                  «      }t        |«      }|dk  r||z  }d|cxk  r|dz
  k  sJ d«       ‚ J d«       ‚|dk  sJ d«       ‚t        |«      |k(  sJ d«       ‚d}d}||   |k  sJ d	«       ‚|||<   | j                  |«      }d
|z
  |gt	        |«      D �cg c]  }| j                  |«      ‘Œ c}z   }	||g|z   }
ddg|z   }t        | |
|	|«      S c c}w )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[batch_offset : batch_offset + batch_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        )ÚlistÚshapeÚlenÚstrideÚranger   )ÚTÚblock_shapeÚ
ragged_dimÚtensor_shapeÚrankÚmax_intÚbillionÚragged_strideÚiÚ
tma_strideÚ	tma_shapeÚ	box_shapes               úY/var/www/pod-logistic/pod-ai/venv/lib/python3.12/site-packages/triton/tools/ragged_tma.pyÚcreate_ragged_descriptorr      s)  € ô �{Ó#€KÜ˜Ÿ™“=€LÜˆ|Ó€Dà�A‚~Ø�dÑˆ
à�
Ô%˜T A™XÒ%ÐHÐ'HÓHÑ%ÐHÐ'HÓHÐ%Ø�1Š9ÐTÐTÓTˆ9äˆ{Ó˜tÒ#ÐXÐ%XÓXÐ#à€GØ€Gà˜
Ñ# wÒ.ÐUÐ0UÓUÐ.Ø&€L�ÑØ—H‘H˜ZÓ(€Mð ˜-Ñ'¨Ð7ÔPUÐVZÓP[Ö:\È1¸1¿8¹8ÀA½;Ò:\Ñ\€JØ˜7Ð# lÑ2€IØ�Q�˜+Ñ%€Iä˜A˜y¨*°iÓ@Ð@ùò	 ;]s   Â0C)c                 ó*   — d}||z
  |z   }| |z   }|||fS )z;
    Helper function for load_ragged and store_ragged.
    r   © )Úbatch_offsetÚ
batch_sizeÚrowr   ÚxÚys         r   Úto_ragged_indicesr"   0   s.   € ð €GØ�*Ñ˜sÑ"€AØ�zÑ!€Aà�A�qˆ=Ðó    r   c                 ó0  — t        j                  t        | j                  «      t        |«      dz   k(  d«       t	        ||||   «      \  }}}| j                  ||g|d| z   |gz   ||dz   d z   «      }t        j                  ||j                  dd «      }|S )zù
    Read from a subarray T[batch_offset : batch_offset + batch_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_raggedr0   =   s˜   € ô ×Ñ”S˜Ÿ™“^¤s¨6£{°Q¡Ñ6Ð8dÔeä" <°¸VÀJÑ=OÓP�J€BˆˆBØ�8‰8�R˜�H˜v k zÐ2Ñ2°b°TÑ9¸FÀ:ÐPQÁ>À?Ð<SÑSÓT€DÜ�:‰:�d˜DŸJ™J q r˜NÓ+€DØ€Kr#   c                 óÄ   — t        ||||   «      \  }}}t        j                  |ddg|j                  z   «      }| j	                  ||g|d| z   |gz   ||dz   d z   |«       y)a  
    Write to a subarray T[batch_offset : batch_offset + batch_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+   r/   r   r,   r-   r.   s	            r   Ústore_raggedr4   O   so   € ô # <°¸VÀJÑ=OÓP�J€BˆˆBÜ�:‰:�d˜Q ˜F T§Z¡ZÑ/Ó0€DØ‡I�Iˆr�2ˆh˜  Ð,Ñ,°¨tÑ3°f¸ZÈ!¹^¸_Ð6MÑMÈtÕTr#   c                 óÄ   — t        ||||   «      \  }}}t        j                  |ddg|j                  z   «      }| j	                  ||g|d| z   |gz   ||dz   d z   |«       y)a  
    Atomic add into a subarray T[batch_offset : batch_offset + batch_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_addr3   s	            r   Úatomic_add_raggedr7   _   sp   € ô # <°¸VÀJÑ=OÓP�J€BˆˆBÜ�:‰:�d˜Q ˜F T§Z¡ZÑ/Ó0€DØ‡N�N�B˜�8˜f [ jÐ1Ñ1°R°DÑ8¸6À*ÈqÁ.À/Ð;RÑRÐTXÕYr#   )r   )ÚtritonÚtriton.languageÚlanguager&   Útriton.tools.tensor_descriptorr   r   Újitr"   Ú	constexprr0   r4   r7   r   r#   r   ú<module>r>      sœ   ðÛ Ý Ý ;ó
%AðP ‡�ñ	ó ð	ð ‡�ñÀ2Ç<Á<ò ó ðð" ‡�ñUÈ"Ï,É,ò Uó ðUð ‡�ñZÈrÏ|É|ò Zó ñZr#   