^j:6ddlmZddlmZmZmZddlmZddlm cm cm cm Z ddlmZmZddlmZmZer ddlmZddlmZgd Zed Gd d e j.ZeGdde j2Ze d ddZe d ddZe d ddZedddZe d d dZe d d!dZ e d" d#dZ!e d" d$dZ"y)%) annotations)ListTuple TYPE_CHECKING) dataclassN)PaddedSharedLayoutSwizzledSharedLayout)builtin_unwrap_if_constexpr)ir)shared_memory_descriptor) async_load async_waitmake_tensor_descriptortensor_descriptortensor_descriptor_typeprefetch async_scatterT)eqcbeZdZUdZded<ded<ded<ded<dd Zdd Zdd Zdd Zdd Z y)rz!The type for a tensor descriptor.zttgl.block_type block_typezttgl.tuple_type shape_type strides_type)PaddedSharedLayout | SwizzledSharedLayoutlayoutc<d|jd|jdS)Nztensor_descriptor)rrselfs }/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/amd/gfx1250/tdm.py__str__ztensor_descriptor_type.__str__s #DOO#4Bt{{m1EEc||}|dz }|jj||\}}|jj||\}}t||||}||fS)N)r _unflatten_irrr)rhandlescursorhandleshapestridesvalues r r%z$tensor_descriptor_type._unflatten_irsd! 55gvF v++99'6J!&%$?f}r"c|jjj}|j|jj |||j j |SN)r element_ty is_int_signed!get_tensor_descriptor_layout_typeto_irr_to_ir)rbuilder is_signeds r r2ztensor_descriptor_type._to_ir'sTOO..<<> 88 OO ! !' *  KK  w '  r"c|j|j||jj|||jj||yr-)appendr2r_flatten_ir_typesr)rr3outs r r7z(tensor_descriptor_type._flatten_ir_types/sA 4;;w'( ))'37 ++GS9r"c d|jjd|jjd|jjd|jjd S)NTD_)rmanglerrrrs r r<ztensor_descriptor_type.mangle4sbDOO**,-Qt/E/E/G.H$J[J[JbJbJdIeefgkgrgrgygyg{f||~r"N)returnstr)r&List[ir.value]r'intr=zTuple[tensor_descriptor, int])r3 ir.builderr=zir.type)r3rAr8z List[ir.type]r=None) __name__ __module__ __qualname____doc____annotations__r!r%r2r7r<r"r rrs8+!! 55F : @r"rceZdZUdZded<ded<ded<ded<dd Zed Zed Zed Z ed Z y)rz4A descriptor representing a tensor in global memory.zir.valuer(z ttgl.tupler)r*rtypec|j|j|jj||jj|yr-)r6r(r) _flatten_irr*)rr&s r rLztensor_descriptor._flatten_irAs6t{{# w'   )r"c.|jjSr-)rJrrs r rztensor_descriptor.block_typeFsyy###r"cB|jjjSr-)rJrr)rs r block_shapeztensor_descriptor.block_shapeJsyy##)))r"cB|jjjSr-)rJrr.rs r dtypeztensor_descriptor.dtypeNsyy##...r"c.|jjSr-)rJrrs r rztensor_descriptor.layoutRsyyr"N)r&r?r=rB) rCrDrErFrGrLpropertyrrOrQrrHr"r rr8sr>     * $$**//  r"rct|}d|cxkrdks nJd|dt||k(sJd|dt|t||k(sJd|dt|t|jtjsJd t |}t|t tfsJd t|tr|jdk(sJd |j}|j|d }|j|d } tj|}tj|}tj|jj|} t| |j|j|} |j!d} |j"j%| j'|j"||| | } t)| ||| S)aMake a tensor descriptor object. Args: base (tensor): base pointer of the tensor in global memory. shape (List[int]): shape of the tensor. strides (List[int]): strides of the tensor. block_shape (List[int]): block shape of the tensor. layout (PaddedSharedLayout | SwizzledSharedLayout): the layout of the tensor in shared memory. Returns: tensor_descriptor: the created tensor descriptor object r$z Expected 1 <= ndim <= 5 but got z dimensionsz Expected z strides but got zExpected block_shape to have z dimensions but got zExpected base to be a pointerzBExpected layout to be a PaddedSharedLayout or SwizzledSharedLayoutz3Expected max_phase to be 1 for SwizzledSharedLayoutF require_i64Tzero)len isinstancerQttgl pointer_typer rr max_phaser(_convert_to_ir_valuestuplerrJr.r_str_to_padding_optionr3create_make_tensor_descriptorr2r)baser)r*rOr _semanticndim base_handle shape_handlesstride_handlesrrJpaddingr(s r rrWs u:D >>O=dV;OO> w<4 R9TF2CCL>!RR  { t #m'DTFJ^_bcj_k^l%mm # djj$"3"3 4U6UU 4 !& )F f13GH IMLM I&./1$[&[[$++K33Eu3MM44W$4ON JJu Ejj!G!5!5{CJ !*ejj',, OD..v6G    < @r"ct|j}|dk(s Jd|dt|j}|dk(s Jd|d|j|gdd}t |}| |j nt jj} |jj|j |j ||j | y) auGather data from non-contiguous rows in global memory to shared memory asynchronously. This operation uses TDM gather mode to read data from non-contiguous rows in global memory. Unlike async_load which reads from contiguous rows, gather allows reading from arbitrary rows specified by the src_row_indices tensor. The dtype of src_row_indices determines the index size: - int16: up to 16 rows can be gathered per TDM instruction - int32: up to 8 rows can be gathered per TDM instruction If more rows are needed, multiple TDM instructions will be automatically issued. Args: desc (tensor_descriptor): the source tensor descriptor. Must be 2D. src_row_indices (tensor): 1D tensor of row indices (int16 or int32) in the source tensor. src_col_offset (int or tensor): the starting column offset in the source tensor for all gathered rows. dst (shared_memory_descriptor): the shared memory destination to store gathered data. Must be 2D. mbarrier (shared_memory_descriptor, optional): The barrier object to signal "arrive" on. r{z)TDM gather only supports 2D tensors, got r|zTDM gather dst must be 2D, got FrVrN) rYrOr)r^r r(r[r r+r3create_async_tdm_gather) r~src_row_indicessrc_col_offsetdstrprcrddst_ndimsrc_col_offset_handlerss r async_gatherrs, t D 19IA$qII9399~H q=G;H:QGG=&;;^rs"--!77`RQ   d!@T^^!@!@H     < Y](;$F(;Ui(;#L(;ar(; (;V LPM!9MUYM M( EI v2 vNR v  v  = = GK"@4"@PT"@ "@J FJ"?3"?OS"? "?J _c26gg;?g g$ rvEI -1 NY   r"