o c۷i4@sddlmZmZmZmZmZddlmZddlZddl m Z ddlm Z m Z m Z mZddlmZmZmZddlmZddlmZddlZddlmZdd lmZdd lmZd Zedd ddd de jde j de j dee j de!ddf ddZ"edddde j de j fddZ#edddde jde j de j e j$Bde j fddZ%edddde j&de j$de j de j fddZ'e ddddd eej(d!e)d"e!de j*fd#d$Z+edd ddd%de j de j dee j d"e!ddf d&d'Z, dd eej(d)e)d!e)d"e!de jf d*d+Z- ( dd eej(d,e)d)e)d!e)d"e!de jf d-d.Z.e j/d/e j d0e de j fd1d2Z0d3Z1d4Z2d5e1Z3e ddddd6e j d7e)d8e!de j fd9d:Z4e ddddd6e j d;e dZ5d?e d@e)dAe)dBe)de f dCdDZ6dEe j7fdFdGZ8dHe j de j fdIdJZ9de j:j&dHe j de j fdKdLZ;de j:j&dHe j de j fdMdNZdPeej(de j*fdQdRZ? ddSej@e)dTee j(dUe!dVee)de j*f dWdXZA ddSej@e)dTee j(dUe!dVee)de j*f dYdZZB dd[e jCd\e j d]e dSe)dUe!dVee)deee je j ffd^d_ZD dd[e jCd\e j d]e dSe)dUe!deee je j ff d`daZE dd[e jCdbe j$dUe!de jfdcddZF dd[e jCdbe j$d\ee j d]e dSe)dUe!deee je j e j ffdedfZG dd[e jCdge j d]e dSe)deee je j ff dhdiZH dd[e jCdge j d]e dSe)dje!deee je j ff dkdlZIeddddme j7dne j7doe)e BfdpdqZJeddddre j*de j7fdsdtZKed(ddddudve j7dwe j7dxe j7dye dzee d{e)ddfd|d}ZL dd~e j de j de!defddZMe ddddde j*de jNde jOd~e j de j de!de!defddZPd'edejQjRfddZSe j/de j&de j dge j de j de de defddZTe j/de j&de j dge j de j de de defddZUe j/ (ddre j*de j dge j de j de de)d{e)defddZVdS))OptionalTypeTupleCallableSequence)partialN)Int32Int16Boolean const_expr)cpasyncwarp warpgroup)CtaGroup) dsl_user_op)llvm)ir) cute_nvgpul}F)predretilelocip tiled_copysrcdstrrreturnc Kst|jtjr|jtjjksJt|j|jkr,t ||j}| | |j|}t|r5| |}tj|||f|||d|dS)Nrrr) isinstanceiteratorcutePointermemspace AddressSpacermemr element_typemake_rmem_tensor_likestoreloadtorcopy) rrrrrrrkwargssrc_cvtr,F/home/ubuntu/vllm_env/lib/python3.10/site-packages/quack/copy_utils.pycvt_copys  "r.rrcCs*tj||j||d}tj||||d|SNr/)rr%r$ autovec_copy)rrrrr,r,r-load_s2r,sr2 dst_shapecCsJtt|tj rtj||j||d}n|}tj||||||d|Sr0)r rrTensormake_rmem_tensorr$r)r)rrr3rrrr,r,r-load_s2r_retile3s  r6thr_copyshapecCs>t|}tj||j|j||d}tj|||||d|Sr0)rmake_identity_tensorr5 partition_Dr8r$r))r7r8rrrcDstrr,r,r-load_t2rEs r<dtypenum_copy_elemsis_asynccCs:ttd||j}|rtntj}tj|||dS)Nnum_bits_per_copy) r minwidthr CopyG2SOprnvgpuCopyUniversalOpmake_copy_atom)r=r>r?rr num_copy_bitscopy_opr,r,r- get_copy_atomOsrK)rr?rrc Ks>|jdd}t|j||}tj|||f|||d|dS)Nrr)r8rKr$rr)) rrrr?rrr*r> copy_atomr,r,r-r)Xs "r) num_threadsc CsR||j}|r tntj}tj|||d}t|}t|}t|||S)NrA) rDr rErrFrGrH make_layoutmake_tiled_copy_tv) r=rNr>r?rIrJrL thr_layout val_layoutr,r,r- tiled_copy_1dhs   rSthreads_per_rowc Csr||j}|r tntj}tj|||d}||dks Jtj|||fdd}td|f} t ||| S)NrAr)rMr)orderrM) rDr rErrFrGrHmake_ordered_layoutrOrP) r=rTrNr>r?rIrJrLrQrRr,r,r- tiled_copy_2dss  rWtAcAlimitc Csttjtj|ddgdtj|dgdtj|dgdftj|dgdddfdt}t|jdD]!}t|jdD]}t|d|fd|fd|||d|f<q||dks#Jd|| fd||d} d|d||f} n#||dksFJd|| fd||d} d|d|||f} t| || S)NrrZr\rkrMrl)rrmr_rp) rgryrzrhrirrrmbig_int offset_val offset_tuple index_tupler,r,r-offset_ragged_tensors rptr_intbmscCs(d|>d}|||>}|||@|?AS)NrMr,)rrrrbit_mskyyy_mskr,r,r- swizzle_ints  rptrcCs8|jj}t||j|j|j}tj|j ||j |j dS)N) assumed_align) type swizzle_typertointnum_bitsnum_base num_shiftrmake_ptrr=r! alignment)rswzrr,r,r- swizzle_ptrsrtensorc Csh|j}|jj}|jjj}t|j|j |j }t |dt |dt d||}t tj|j|jd|S)Nr)r=)layoutr$rDrrrr make_swizzlerrr recast_layoutmake_composed_layoutrq recast_ptr)routerrDrinner new_layoutr,r,r-&as_position_independent_swizzle_tensors rcC$tt||j|t|jSrk)rrqrr:rrrr7rr,r,r- partition_D_position_independent)rcCrrk)rrqr partition_Srrrrr,r,r- partition_S_position_independent2rrlayout_c elem_ty_ccCs`t|tjjstd||}|jdkr$tjt |d|||dStjtj |||dS)a Selects the largest vectorized smem load atom available subject to constraint of gmem layout. Parameters: ----------- layout_c : LayoutEnum The layout enum of the output tensor D. elem_ty_c : Type[Numeric] The element type for output tensor D. Returns: -------- Either SmemLoadMatrix or SimtSyncCopy, based on the input parameters. z%elem_ty_c must be a Numeric, but got rjr/) rr` cutlass_dsl NumericMeta TypeError is_m_major_crDrrHr LdMatrix8x8x16bOprFrG)rrrr is_m_majorr,r,r-sm90_get_smem_load_op;s  rarchr$ transposemajor_mode_sizecC|t|dkp |jdkrtjtj||sdnd|jdS|dus'|ddkr)dn |ddkr1dnd}ttj||d |S NZrr\rMrArrjr)r num_matrices)r rDrrHrFrGr StMatrix8x8x16bOprr$rrrr,r,r-get_smem_store_atom\ rcCrr)r rDrrHrFrGr rrr,r,r-get_smem_load_atomtrr tiled_mmasCtidxc sv|j}t||||d}t|||} t| r!| |nt| |ddtjdt t ffdd } | | fS)N)rrdst_idxcs<t|durnddd|f}t||fddi|dSNrT)r r.)rr new_kwargs dst_tensortRS_sCrr,r-copy_fns z!get_smem_store_C..copy_fnrk) r$rrmake_tiled_copy_C get_slicer r:rr4rr) rrrrrposition_independentrr=rLr7rr,rr-get_smem_store_Cs        rc s|j}t|||}t|||}t| r||nt||t|||} t| ||} | t |j ddj ddt t ffdd } | |fS)Nr\src_idxcs6t|durnddd|f}t|fdi|SNr3)r r6)rr src_tensor tRS_shapetSR_sCrr,r-rs z get_smem_load_C..copy_fnrk) r$rrrrr rrrr9r8rr) rrrrrrr=rLr7 copy_atom_RS thr_copy_RSrr,rr-get_smem_load_Cs        repi_tilecCs<ttj||dddkrdnddtj}t||}|S)NrMrrrjr\)r)rrHr rr`Float16make_tiled_copy_C_atom)rrr copy_atom_Ctiled_copy_C_atomr,r,r-epilog_smem_copy_atoms   rcst|dur |jntj}t||}t|||} t| ||} dt|dur:t| r5| |nt | ||durE|j ddn|} | t | j } t| |jj} dtjdtffdd }t|durm|nd| | fS)Nr\rrcs$t|ddd|ffi|dSrkr.rrrrr,r-rs$z#get_smem_store_epi..copy_fn)r r$r`rrrrmake_tiled_copy_Srr:rr8rr9r5op acc_dtyper4r)rrrrrrrr=rrLr7sC_shape tRS_rC_shapetRS_rCrr,rr-get_smem_store_epis        rsAc s||j}|jjtjjk}t|||}t|| |}t | r'| |nt ||dtj dtffdd } | |fS)Nrrcs(t|ddd|ffddi|dSrrrtRS_sArr,r-rs(z!get_smem_store_A..copy_fn)r$r a_major_moderOperandMajorModeMNrrmake_tiled_copy_Arr r:rr4r) rrrrrr=rrLr7rr,rr-get_smem_store_As       rwith_dst_tensorc s|j}|jjtjjk}t|||}t|| |} t | r'| |nt | || |jdddtffdd } dtdtjffdd } |sU| | fS| | fS)Nr\rcs"tddd|ffdi|Srr6)rrrtSR_sArr,r-rsz get_smem_load_A..copy_fnrcs tddd|f|fi|Srkr)rrr)rrr,r-copy_fn_w_dst_tensors z-get_smem_load_A..copy_fn_w_dst_tensor)r$rrrrrrrrrr rrpartition_shape_Ar8rr4) rrrrrrr=rrLr7rrr,rr-get_smem_load_As      rsmem_ptrgmem_ptr store_bytesc CsB|j||d}tjd|j|t|gddddtjjddS)Nr/zJcp.reduce.async.bulk.global.shared::cta.bulk_group.add.f32 [$0], [$1], $2;zl,r,rTF)has_side_effectsis_align_stack asm_dialect)rir_valuer inline_asmllvm_ptrr AsmDialectAD_ATT)rrrrr smem_ptr_i32r,r,r-cpasync_reduce_bulk_add_f32s  rtma_atomcCs2tj|jj||d}tjd}tj||||dS)aw Get the address of the TMA descriptor embedded in a TMA Copy Atom. Extracts the constant memory address of the TMA descriptor for use with custom PTX instructions. :param tma_atom: TMA Copy Atom from make_tiled_tma_atom :return: Pointer to TMA descriptor in constant memory Example: >>> desc_ptr = get_tma_descriptor_address(tma_atom) r/z@!cute.ptr>)_cute_nvgpu_iratom_make_exec_tma_traitvaluerrparseget_tma_desc_addr)rrr exec_atomtma_desc_ptr_typer,r,r-r4s r)num_ctamulticast_maskrr tma_desc_ptr dst_smem_ptr mbarrier_ptrcol_idx row_indicesrc Cst|dkrtdt|t|} dd|D} |j||d} |j||d} |j||d} |dkr=| t@} | } d}|durMt|}|dusUJdd |d }tjd| | | | d | d| d | d | g|dddtj j ||d dS)a! Perform TMA gather4 load from global memory to shared memory. Issues PTX instruction: cp.async.bulk.tensor.2d.shared::cta.global.tile::gather4.mbarrier::complete_tx::bytes [dstMem], [tensorMap, {col_idx, row0, row1, row2, row3}], [smem_bar]; This loads 4 rows (specified by row_indices) from a 2D tensor at the given column index into shared memory, using the TMA descriptor. :param tma_desc_ptr: Pointer to TMA descriptor in constant memory (128-byte aligned) :type tma_desc_ptr: Pointer :param dst_smem_ptr: Destination address in shared memory :type dst_smem_ptr: Pointer :param mbarrier_ptr: Pointer to mbarrier in shared memory for completion tracking :type mbarrier_ptr: Pointer :param col_idx: Column index :type col_idx: Int32 :param row_indices: Sequence of exactly 4 row indices :type row_indices: Sequence[Int32] :param num_cta: Number of CTAs participating (default: 1) :type num_cta: int :param multicast_mask: Optional multicast mask :type multicast_mask: Int16 Requirements: - row_indices must contain exactly 4 elements - Compute capability >= SM_100 (Blackwell) - TMA descriptor must be properly initialized for 2D tensor Example: >>> from cutlass.cute.nvgpu import cpasync >>> from cutlass.cute import core >>> >>> # Create TMA descriptor >>> tma_atom, tma_tensor = cpasync.make_tiled_tma_atom(...) >>> tma_desc_ptr = get_tma_descriptor_address(tma_atom) >>> >>> # Compute indices (typically from kernel logic) >>> col_idx = core.get(...) or 5 # Int32 value >>> row_indices = [core.get(...) for _ in range(4)] # 4 Int32 values >>> >>> # Gather 4 rows at computed column >>> tma_gather4_load( ... tma_desc_ptr=tma_desc_ptr, ... dst_smem_ptr=smem_ptr, ... mbarrier_ptr=barrier_ptr, ... col_idx=col_idx, ... row_indices=row_indices ... ) rjz,gather4 requires exactly 4 row indices, got cSsg|]}t|qSr,)rr).0row_idxr,r,r- z$tma_gather4_load..r/rMNzmulticast is not supported yetzacp.async.bulk.tensor.2d.shared::cta.global.tile::gather4.mbarrier::complete_tx::bytes.cta_group::z( [$0], [$1, {$2, $3, $4, $5, $6}], [$7];rr\rlzr,l,r,r,r,r,r,rTF)rrrrr) len ValueErrorrrrSm100MmaPeerBitMaskr rrrr)rrr r r rrrrcol_valrow_vals desc_addrdst_addr mbar_addrmulticast_mask_valptxr,r,r-tma_gather4_loadIsF @     rrr single_stagec stt||s dnd}tt||sdnd}t|d|t|d|dtjffdd }dtjffdd }t| rG|S|S)NrMr tma_bar_ptrcspttj}tj tj|d|fd|ffd|i|WddS1s1wYdSNmbar_ptrrrHr CopyBulkG2SOpr$r elect_oner))rrrratomrr*rr,r- copy_bulks   "z+cpasync_bulk_get_copy_fn..copy_bulkcs`ttj}tjtj|fd|i|WddS1s)wYdSrr)rrr"r#r,r-copy_bulk_single_stages  "z8cpasync_bulk_get_copy_fn..copy_bulk_single_stage)r rrm group_modesr )rrrr*group_rank_srcgroup_rank_dstr$r%r,r#r-cpasync_bulk_get_copy_fns r)r" cta_coord cta_layout filter_zerosc s*tt|jtjo|jtjjk} | r||fn||f\} } tt| |s%dnd} tt| |s2dnd}t j ||t | d| t | d|||d\}}t|r\t |}t |}| rb||fn||f\t dddfdd }t dddfdd }t| r|||fS|||fS)NrMrr/cs8tjd|fd|ffi|||ddSr0rr))rrrrrr"rr*rr,r-copy_tmasz!tma_get_copy_fn..copy_tmacs(tjfi|||ddSr0r-)rrrr.r,r-copy_tma_single_stages(z.tma_get_copy_fn..copy_tma_single_stage)r rrrr r!r"smemrmr tma_partitionr&r,r)r"r*r+rrr,rrrr* src_is_smem smem_tensor gmem_tensorgroup_rank_smemgroup_rank_gmemrgr/r0r,r.r-tma_get_copy_fns2      r9pipelinecsdtjjffdd }|S)Nproducer_statecs"d||j|d|dS)N)rrrr,)indexproducer_get_barrier)rr;rr)r:r,r-r s z%tma_producer_copy_fn..copy_fn)r`r: PipelineState)r)r:rr,r>r-tma_producer_copy_fn sr@ thr_copy_AmAgsAIdxlimit_mlimit_kc stj|dgdtj|dgdf | jddksJtt ddd d jdjdkt rAt| d}t jddt } | d ||dd}ddttjjdgd}ttjjdgdt |t tj|ddD]}d|dfd|k|<qt |ttj|ddD]}d|dfd} |r|| |<qd|<qt|d dfd d tf f d d } | S) NrrZrMr\)NNrNT unroll_fullFrc s d}t|r+tt}| d}tjddD]}dd|fd|k||<qdd|ff}tjdD]G}ttj ||dfdddfd}tsZ|rtj jdgddksgJd d} tj |d| f d|f|f|d q;dS) NrMTrFrr\ up_to_rankNNrrZ)rrrr) r rr5r r`rangerar8 tiled_divide append_onesr_r)) rrrtApA_k limit_k_curkmA_currmA_rowki cols_per_threadelems_per_loadis_even_m_smemrEmA_km_idxt0AcArXtApA_mtAsArA tile_shape_mkr,r-rAs& &z%gather_m_get_copy_fn..copy_fnF)rr_r:r8r&slice_tiler_mnr rCr9rrr5r r`rLrlogical_dividebool) rArBrrCrDrEcArows_per_threadrr rr,rUr-gather_m_get_copy_fns6       *rfcsJd\t|jtjjkr|n |jtjjksJ|tj|dgdtj|dgdf |tdd d j dj dk}t| rOt | d}tj dd}t } | d ||dd}ddttjj dgd} ttjj dgdt| ttj| ddD]} d| dfd|k| <qt j dj |} j} tt||f|| fd| | dfdfdd td ttjtjff fd d } dd td ttjtjff fdd } ddttjtjfd tf fdd }|tdur"| fS|fS)N)NNrrZrMrlr\TrFFrrcsd}t|r+tt}|d}tjddD]}dd|fd|k||<qd|f}tt}tD]&}dd|fd}t| rS||||<q<||r^||||<q.prefetch_from_gmem_fnc sd}t|r+tt}|d}tjddD]}dd|fd|k||<q||d|f}tt} tD]}dd|fd} || | |<qAtj tj | |Wd| |fS1spwY| |fS)NrMTrFr) r rr5r r`rL consumer_waitrr sync_warpr!consumer_release) a_prefetch_pipelinerra_prefetch_consumer_staterrOrPrQ sAIdx_currir )rVrEsAIdxr[rXr^r,r-prefetch_from_smem_fns&        z3gather_k_get_copy_fn..prefetch_from_smem_fn k_idx_tApA_kc s|\}}d}t|rtj|dd}tjdD]4}tjdD])}|rLtjd|||fd||f|ft|durDdn|d|fdq#qdS)Nr\rHrMrK)r r prepend_onesr`rar8r)) rrrtrrirO tApA_k_predrQr)rXtAmAr\r]rAr,r-rs z%gather_k_get_copy_fn..copy_fnr_)r r!rr"gmemr1r_r:r&rar8rCr9rrr5r r`rLthr_idxrb flat_dividercrr4)rArBrrCrDrErXrWrdrerthreads_per_colrrkrsrr,) rVrjrErrr[rXrwr\r]rAr^r-gather_k_get_copy_fnZs^       . r|rrwarp_idx num_warpscstj|dgd}t|d||ddksJ|}ttjtjtddt|td} | |} | |} | |t | t |} t t | |ddtjffd d } | S) NrrZrJrjr@rA)rrc s|}tjtjdgdddD]3fddtdD}dd|fj}tj||||Wdn1s?wYqdS)NrMrZTrFcsg|]}|fqSr,r,)r v)r tSR_rAIdxr,r-rrz=gather_m_get_tma_copy_fn..copy_fn..rj)r`rLrr_rrr!)rrrr r rrrtile_Ktma_gather4_load_fn)rr-rs z)gather_m_get_tma_copy_fn..copy_fn)rr_rPrHrFrGrrOrrr2rrrr )rrBrrrr}r~rtile_M cta_group copy_AIdx_s2rwarp_copy_AIdx_s2r tSR_sAIdxrrr,rr-gather_m_get_tma_copy_fns"    rr_)rMF)rF)FN)FFN)FF)rM)Wtypingrrrrr functoolsrr` cutlass.cuterrr r r cutlass.cute.nvgpur r rcutlass.cute.nvgpu.tcgen05.mmarcutlass.cutlass_dslrcutlass.pipelinecutlass._mlir.dialectsr cutlass._mlirrrrr TiledCopyr4rcr.r2Shaper6ThrCopyr<NumericintCopyAtomrKr)rSrWjitrfrnrorrrxrrr rrcorerrutils LayoutEnumr ConstexprrrTiledMmarrrrrrrrrr)CoordLayoutr9r: PipelineAsyncr@rfr|rr,r,r,r-s0                 D !      #            r    0  C o