o ´^i_ ã@slddlZddlmZddlmZd dd„Zejdd„ƒZejd dej fdd „ƒZ ejd dej fd d „ƒZ dS) éN)ÚTensorDescriptorc sìt|ƒ}tˆjƒ}t|ƒ}|dkr||7}d|kr$|dks)Jdƒ‚Jdƒ‚|dks1Jdƒ‚t|ƒ|ks;Jdƒ‚d}d}|||ksIJd ƒ‚|||<ˆ |¡}d ||g‡fd d „t|ƒDƒ}||g|} ddg|} 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[batch_offset : batch_offset + batch_size] with hardware bounds-checking preventing any sort of leakage outside the subarray. rézlast dimension cannot be raggedéz)sz,create_ragged_descriptor..)ÚlistÚshapeÚlenrÚranger) r Ú block_shapeÚ ragged_dimÚ tensor_shapeÚrankÚmax_intÚbillionÚ ragged_strideÚ tma_strideÚ tma_shapeÚ box_shaperr r Úcreate_ragged_descriptors" ( "  rcCs"d}|||}||}|||fS)z; Helper function for load_ragged and store_ragged. rr)Ú batch_offsetÚ batch_sizeÚrowrÚxÚyrrr Úto_ragged_indices0s  r"rc Cs~t t|jƒt|ƒdkd¡t||||ƒ\}}}| ||g|d|…|g||dd…¡}t ||jdd…¡}|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_assertrrr"ÚloadÚreshape) ÚTMArrÚcoordsrÚc0Úc1Úc2Údatarrr Ú load_ragged=s  0r.c Csbt||||ƒ\}}}t |ddg|j¡}| ||g|d|…|g||dd…|¡dS)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(). rN)r"r$r'rÚstore) r(rrr)r-rr*r+r,rrr Ú store_raggedOs 6r0)r) ÚtritonÚtriton.languageÚlanguager$Útriton.tools.tensor_descriptorrrÚjitr"Ú constexprr.r0rrrr Ús   (