o ܹi_ @slddlZddlmZddlmZd ddZejddZejd dej fdd Z ejd dej fd d Z dS) N)TensorDescriptorc st|}tj}t|}|dkr||7}d|kr$|dks)JdJd|dks1Jdt||ks;Jdd}d}|||ksIJd |||<|}d ||gfd 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. rzlast dimension cannot be raggedz)sz,create_ragged_descriptor..)listshapelenrranger) r block_shape ragged_dim tensor_shaperankmax_intbillion 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_sizerowrxyrrr to_ragged_indices0s  r"rc Cs~tt|jt|dkdt||||\}}}|||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"loadreshape) TMArrcoordsrc0c1c2datarrr 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'rstore) r(rrr)r-rr*r+r,rrr store_raggedOs 6r0)r) tritontriton.languagelanguager$triton.tools.tensor_descriptorrrjitr" constexprr.r0rrrr s   (