o i@s ddlmZddlmZmZmZddlmZddlm m m m Z ddlmZmZddlmZmZerCddlmZddlmZgd Zed d Gd d d e jZeGddde jZe d*d+ddZe d,d-d$d%Ze d*d.d&d'Zed/d0d(d)ZdS)1) annotations)ListTuple TYPE_CHECKING) dataclassN)PaddedSharedLayoutSwizzledSharedLayout)builtin_unwrap_if_constexpr)ir)shared_memory_descriptor) async_load async_waitmake_tensor_descriptortensor_descriptortensor_descriptor_typeT)eqc@sdeZdZUdZded<ded<ded<ded<d!d d Zd"ddZd#ddZd$ddZd!ddZ d S)%rz!The type for a tensor descriptor.zttgl.block_type block_typezttgl.tuple_type shape_type strides_type)PaddedSharedLayout | SwizzledSharedLayoutlayoutreturnstrcCsd|jd|jdS)Nztensor_descriptor)rrselfre/home/kuhnn/.local/lib/python3.10/site-packages/triton/experimental/gluon/language/amd/gfx1250/tdm.py__str__sztensor_descriptor_type.__str__handlesList[ir.value]cursorintTuple[tensor_descriptor, int]cCsJ||}|d7}|j||\}}|j||\}}t||||}||fS)N)r _unflatten_irrr)rr r"handleshapestridesvaluerrrr&s z$tensor_descriptor_type._unflatten_irbuilder ir.builderir.typecCs*|jj}||j|||j|SN)r element_ty is_int_signed!get_tensor_descriptor_layout_typeto_irr_to_ir)rr+ is_signedrrrr3$s   ztensor_descriptor_type._to_irout List[ir.type]NonecCs0||||j|||j||dSr.)appendr3r_flatten_ir_typesr)rr+r5rrrr9,sz(tensor_descriptor_type._flatten_ir_typesc Cs6d|jd|jd|jd|jd S)NTD_)rmanglerrrrrrrr<1s6ztensor_descriptor_type.mangleN)rr)r r!r"r#rr$)r+r,rr-)r+r,r5r6rr7) __name__ __module__ __qualname____doc____annotations__rr&r3r9r<rrrrrs     rc@sleZdZUdZded<ded<ded<ded<dd dZeddZeddZeddZ eddZ dS)rz4A descriptor representing a tensor in global memory.zir.valuer'z ttgl.tupler(r)rtyper r!rr7cCs(||j|j||j|dSr.)r8r'r( _flatten_irr))rr rrrrC>s  ztensor_descriptor._flatten_ircC|jjSr.)rBrrrrrrCztensor_descriptor.block_typecC |jjjSr.)rBrr(rrrr block_shapeG ztensor_descriptor.block_shapecCrFr.)rBrr/rrrrdtypeKrHztensor_descriptor.dtypecCrDr.)rBrrrrrrOrEztensor_descriptor.layoutN)r r!rr7) r=r>r?r@rArCpropertyrrGrIrrrrrr5s     rbase ttgl.tensorr("List[ttgl.constexpr | ttgl.tensor]r)rGList[ttgl.constexpr]rrrcCsRt|}d|krdksnJd|dt||ks)Jd|dt|t||ks;Jd|dt|t|jtjsFJd t|}t|ttfsUJd t|trc|jdkscJd |j }|j |d d }|j |dd } t |}t |}t |j j|} t| |j |j |} |d} |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 isinstancerIttgl pointer_typer rr max_phaser'_convert_to_ir_valuestuplerrBr/r_str_to_padding_optionr+create_make_tensor_descriptorr3r)rKr(r)rGr _semanticndim base_handle shape_handlesstride_handlesrrBpaddingr'rrrrTs.&$$    rsrcoffsetsdestr predboolmbarrierr7c CsZ|j|dd}||}|j}t|}|dur|jntj}|j|j||j||dS)a-Load a block of tensor specified in tensor descriptor from global memory to shared memory asynchronously. Args: src (tensor_descriptor): the source tensor descriptor. offsets (List[int]): the offsets from the base pointer in the tensor descriptor. dest (shared_memory_descriptor): the shared memory destination to store the loaded data. pred (bool, optional): Predicate to enable or disable the load. Defaults to True. mbarrier (shared_memory_descriptor, optional): The barrier object to signal "arrive" on. FrPN) rX to_tensorr'r rUr r*r+%create_async_tdm_copy_global_to_local) rbrcrdrergr\offset_handles pred_handlembarrier_handlerrrr s r cCs&|j|dd}|j|j||jdS)akStore a block of tensor specified in tensor descriptor from shared memory to global memory asynchronously. Args: dest (tensor_descriptor): the destination tensor descriptor. offsets (List[int]): the offsets from the base pointer in the tensor descriptor. src (shared_memory_descriptor): the shared memory source to load the data. FrPN)rXr+%create_async_tdm_copy_local_to_globalr')rdrcrbr\rjrrr async_stores rncCst|}|j|dS)zWait for the outstanding asynchronous tensor operations to complete. Args: num_outstanding (int): number of outstanding async tensor operations to wait for. N)r r+create_async_tdm_wait)num_outstandingr\rrrrsrr.) rKrLr(rMr)rMrGrNrrrr)TNN) rbrrcrMrdr rerfrgr rr7)rdrrcrMrbr rr7)rN)rr7) __future__rtypingrrr dataclassesr(triton.experimental.gluon.language._core experimentalgluonlanguage_corerU+triton.experimental.gluon.language._layoutsrrr r triton._Cr r __all__ base_typer base_valuerrr rnrrrrrs2    $+