+
    ‹°³j–  ã                   óÞ   € ^ RI t ^ RIHt ^ RIHt RR lt] P                  R 4       t] P                  RR R ll4       t	] P                  RR R ll4       t
] P                  RR	 R
 ll4       tR# )é    N)ÚTensorDescriptorc                ó"  € \        V4      p\        V P                  4      p\        V4      pV^ 8  d	   W$,          p^ Tu;8:  d   V^,
          8  g   Q R4       h Q R4       hV^8:  g   Q R4       h\        V4      V8X  g   Q R4       hRpRpW2,          V8:  g   Q R4       hWcV&   V P                  V4      pRV,
          V.\	        V4       Uu. uF  q€P                  V4      NK  	  up,           p	WU.V,           p
^^.V,           p\        W
W›4      # u upi )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.
zlast dimension cannot be raggedz<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   &&&         ÚX/home/luca/screen-checking/.venv/lib/python3.14/site-packages/triton/tools/ragged_tma.pyÚcreate_ragged_descriptorr      s!  € ô �{Ó#€KÜ˜Ÿ™“=€LÜˆ|Ó€Dà�A„~ØÕˆ
à�
Ö%˜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Ð.Ø&�ÑØ—H‘H˜ZÓ(€Mð ˜-Õ'¨Ð7ÔPUÐVZÔP[Ó:\ÑP[È1¿8¹8ÀA¾;ÑP[Ñ:\Õ\€JØÐ# lÕ2€IØ�Q�˜+Õ%€Iä˜A¨*Ó@Ð@ùò	 ;]s   Ã	Dc                ó<   € RpW1,
          V,           pW,           pW5V3# )z3
Helper function for load_ragged and store_ragged.
r   © )Ú	slice_offÚ
slice_sizeÚrowr   ÚxÚys   &&&   r   Úto_ragged_indicesr    0   s(   € ð €GØÕ˜sÕ"€AØÕ€Aà�qˆ=Ðó    c                ó8   € V ^8„  d   QhR\         P                  /# ©é   r   ©ÚtlÚ	constexpr)Úformats   "r   Ú__annotate__r)   >   s   € ÷ ñ ÄÇÁñ r!   c                ó`  € \         P                  ! \        V P                  4      \        V4      ^,           8H  R4       \	        WW4,          4      w  rVpV P                  WV.VRV ,           V.,           W4^,           R ,           4      p\         P                  ! WˆP                  R,          4      pV# )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 descriptorN:r$   NN)r&   Ústatic_assertr   r   r    ÚloadÚreshape)	ÚTMAr   r   Úcoordsr   Úc0Úc1Úc2Údatas	   &&&&&    r   Úload_raggedr4   =   sŠ   € ô ×Ò”S˜Ÿ™“^¤s¨6£{°Q¥Ñ6Ð8dÔeä" 9¸&Õ:LÓM�J€BˆBØ�8‰8�R�H˜v k zÐ2Õ2°b°TÕ9¸FÐPQÅ>À?Ð<SÕSÓT€DÜ�:Š:�dŸJ™J r�NÓ+€DØ€Kr!   c                ó8   € V ^8„  d   QhR\         P                  /# r#   r%   )r(   s   "r   r)   r)   P   s   € ÷ Uñ UÄrÇ|Á|ñ Ur!   c                óî   € \        WW5,          4      w  rgp\        P                  ! V^^.VP                  ,           4      pV P	                  Wg.VRV ,           V.,           W5^,           R ,           V4       R# )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().
N)r    r&   r-   r   Ústore©	r.   r   r   r/   r3   r   r0   r1   r2   s	   &&&&&&   r   Ústore_raggedr9   O   se   € ô # 9¸&Õ:LÓM�J€BˆBÜ�:Š:�d˜Q ˜F T§Z¡ZÕ/Ó0€DØ‡I�Iˆrˆh˜  Ð,Õ,°¨tÕ3°fÈ!½^¸_Ð6MÕMÈtÖTr!   c                ó8   € V ^8„  d   QhR\         P                  /# r#   r%   )r(   s   "r   r)   r)   `   s   € ÷ Zñ ZÌBÏLÉLñ Zr!   c                óî   € \        WW5,          4      w  rgp\        P                  ! V^^.VP                  ,           4      pV P	                  Wg.VRV ,           V.,           W5^,           R ,           V4       R# )zñ
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().
N)r    r&   r-   r   Ú
atomic_addr8   s	   &&&&&&   r   Úatomic_add_raggedr=   _   sf   € ô # 9¸&Õ:LÓM�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    r4   r9   r=   r   r!   r   Ú<module>rC      sy   ðÛ Ý Ý ;ô
%AðP ‡�ñ	ó ð	ð ‡�öó ðð" ‡�öUó ðUð ‡�öZó òZr!   