o 3/i3@s&ddlmZmZmZmZmZddlmZddlZddl m Z ddlm Z m Z m Z mZddlmZmZmZddlmZddlmZddlZddlmZdd lmZdd lmZd Zedd ddd de jde j de j dee j de!ddf ddZ"edddde jde j de j de de ddf ddZ#edddde j de j fddZ$edddde jde j de j e j%Bde j fddZ&edddd e j'd!e j%de j de j fd"d#Z(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)d0e*d-e*d%e*d&e!de jf d1d2Z/e j0d3e j d4e de j fd5d6Z1d7Z2d8Z3d9e2Z4e ddddd:e j d;e*dZ5e ddddd:e j d?e d@e d;e*dj?dTeej)de j+fdUdVZ@ ddWejAe*dXee j)dYe!dZee*de j+f d[d\ZB ddWejAe*dXee j)dYe!dZee*de j+f d]d^ZC dd_e jDd`e j de dWe*dYe!dZee*deee je j ffdadbZE dd_e jDd`e j de dWe*dYe!deee je j ff dcddZF dd_e jDdee j%dYe!de jfdfdgZG dd_e jDdee j%d`ee j de dWe*dYe!deee je j e j ffdhdiZH dd_e jDdje j de dWe*deee je j ff dkdlZI dd_e jDdje j de dWe*dme!deee je j ff dndoZJeddddpe j8dqe j8dre*e BfdsdtZKeddddue j+de j8fdvdwZLed,ddddxdye j8dze j8d{e j8d|e d}ee d~e*ddfddZM dde j de j de!defddZNe ddddde j+de jOde jPde j de j de!de!defddZQd+edejRjSfddZTe j0de j'de j dje j de j de de defddZUe j0de j'de j dje j de j de de defddZVe j0 ,ddue j+de j dje j de j de de*d~e*defddZWdS))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,]/lsinfo/ai/hellotax_ai/llm_service/venv_vllm/lib/python3.10/site-packages/quack/copy_utils.pycvt_copys  "r.rrseedtidxc Cst|jtjr|jtjjksJddlm}ddl m }t ||j } | } || ||||d} | || | j|j | }tj|||||ddS)zGLike cvt_copy but uses stochastic rounding for FP32 -> BF16 conversion.r)convert_f32_to_bf16_sr) TensorSSAr/N)rrrr r!r"r#quack.roundingr2cutlass.cute.tensorr3r%r$r'r&shaper)) rrrr0r1rrr2r3r+src_vecraw_vecr,r,r- sr_cvt_copy,s  r9cCs*tj||j||d}tj||||d|SNr/)rr%r$ autovec_copy)rrrrr,r,r-load_s2rDsr< dst_shapecCsJtt|tj rtj||j||d}n|}tj||||||d|Sr:)r rrTensormake_rmem_tensorr$r)r)rrr=rrrr,r,r-load_s2r_retileKs  r@thr_copyr6cCs>t|}tj||j|j||d}tj|||||d|Sr:)rmake_identity_tensorr? partition_Dr6r$r))rAr6rrrcDstrr,r,r-load_t2r]s rEdtypenum_copy_elemsis_asynccCs:ttd||j}|rtntj}tj|||dS)Nnum_bits_per_copy) r minwidthr CopyG2SOprnvgpuCopyUniversalOpmake_copy_atom)rFrGrHrr num_copy_bitscopy_opr,r,r- get_copy_atomgsrT)rrHrrc Ks>|jdd}t|j||}tj|||f|||d|dS)Nrr)r6rTr$rr)) rrrrHrrr*rG copy_atomr,r,r-r)ps "r) num_threadsc CsR||j}|r tntj}tj|||d}t|}t|}t|||S)NrJ) rMr rNrrOrPrQ make_layoutmake_tiled_copy_tv) rFrWrGrHrRrSrU thr_layout val_layoutr,r,r- tiled_copy_1ds   r\threads_per_rowc Csr||j}|r tntj}tj|||d}||dks Jtj|||fdd}td|f} t ||| S)NrJr)rVr)orderrV) rMr rNrrOrPrQmake_ordered_layoutrXrY) rFr]rWrGrHrRrSrUrZr[r,r,r- tiled_copy_2ds  r`tAcAlimitc 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)NrrcrertrVru)rrvrhry) rprrrqrrrrrvbig_int offset_val offset_tuple index_tupler,r,r-offset_ragged_tensor s rptr_intbmscCs(d|>d}|||>}|||@|?AS)NrVr,)rrrrbit_mskyyy_mskr,r,r- swizzle_int's  rptrcCs8|jj}t||j|j|j}tj|j ||j |j dS)N) assumed_align) type swizzle_typertointnum_bitsnum_base num_shiftrmake_ptrrFr! alignment)rswzrr,r,r- swizzle_ptr-srtensorc Csh|j}|jj}|jjj}t|j|j |j }t |dt |dt d||}t tj|j|jd|S)Nr)rF)layoutr$rMrrrr make_swizzlerrr recast_layoutmake_composed_layoutrz recast_ptr)routerrMrinner new_layoutr,r,r-&as_position_independent_swizzle_tensor3s rcC$tt||j|t|jSrt)rrzrrCrrrrArr,r,r- partition_D_position_independentArcCrrt)rrzr partition_Srrrrr,r,r- partition_S_position_independentJrrlayout_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 rsr/) rri cutlass_dsl NumericMeta TypeError is_m_major_crMrrQr LdMatrix8x8x16bOprOrP)rrrr is_m_majorr,r,r-sm90_get_smem_load_opSs  rarchr$ transposemajor_mode_sizecC|t|dkp |jdkrtjtj||sdnd|jdS|dus'|ddkr)dn |ddkr1dnd}ttj||d |S NZrrerVrJrrsr)r num_matrices)r rMrrQrOrPr StMatrix8x8x16bOprr$rrrr,r,r-get_smem_store_atomt rcCrr)r rMrrQrOrPr rrr,r,r-get_smem_load_atomrr tiled_mmasCc 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_fnrt) r$rrmake_tiled_copy_C get_slicer rCrr>rr) rrr1rrposition_independentrrFrUrArr,rr-get_smem_store_Cs        rc s|j}t|||}t|||}t| r||nt||t|||} t| ||} | t |j ddj ddt t ffdd } | |fS)Nresrc_idxcs6t|durnddd|f}t|fdi|SNr=)r r@)rr src_tensor tRS_shapetSR_sCrr,r-rs z get_smem_load_C..copy_fnrt) r$rrrrr rrrrBr6rr) rrr1rrrrFrUrA copy_atom_RS thr_copy_RSrr,rr-get_smem_load_Cs        repi_tilecCs<ttj||dddkrdnddtj}t||}|S)NrVrrrsre)r)rrQr rriFloat16make_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)Nrerrcs$t|ddd|ffi|dSrtr.rrrrr,r-rs$z#get_smem_store_epi..copy_fn)r r$rirrrrmake_tiled_copy_SrrCrr6rrBr?op acc_dtyper>r)rrrr1rrrrFrrUrAsC_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 rCrr>r) rrr1rrrFrrUrArr,rr-get_smem_store_As       rwith_dst_tensorc s|j}|jjtjjk}t|||}t|| |} t | r'| |nt | || |jdddtffdd } dtdtjffdd } |sU| | fS| | fS)Nrercs"tddd|ffdi|Srr@)rrrtSR_sArr,r-r(sz get_smem_load_A..copy_fnrcs tddd|f|fi|Srtr)rrr)rrr,r-copy_fn_w_dst_tensor-s z-get_smem_load_A..copy_fn_w_dst_tensor)r$rrrrrrrrrr rrpartition_shape_Ar6rr>) rrr1rrrrFrrUrArrr,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_f323s  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-r Ls r )num_ctamulticast_maskrr tma_desc_ptr dst_smem_ptr mbarrier_ptrcol_idx row_indicesr c 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 ... ) rsz,gather4 requires exactly 4 row indices, got cSsg|]}t|qSr,)rr).0row_idxr,r,r- z$tma_gather4_load..r/rVNzmulticast 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];rreruzr,l,r,r,r,r,r,rTF)rrrrr) len ValueErrorrrrSm100MmaPeerBitMaskr rrrr)rrrrrr rrrcol_valrow_vals desc_addrdst_addr mbar_addrmulticast_mask_valptxr,r,r-tma_gather4_loadasF @     r"rr single_stagec stt||s dnd}tt||sdnd}t|d|t|d|dtjffdd }dtjffdd }t| rG|S|S)NrVr tma_bar_ptrcspttj}tj tj|d|fd|ffd|i|WddS1s1wYdSNmbar_ptrrrQr CopyBulkG2SOpr$r elect_oner))rrr$ratomrr*rr,r- copy_bulks   "z+cpasync_bulk_get_copy_fn..copy_bulkcs`ttj}tjtj|fd|i|WddS1s)wYdSr%r')r$rr*r+r,r-copy_bulk_single_stages  "z8cpasync_bulk_get_copy_fn..copy_bulk_single_stage)r rrv group_modesr )rrr#r*group_rank_srcgroup_rank_dstr,r-r,r+r-cpasync_bulk_get_copy_fns r1r* 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)NrVrr/cs8tjd|fd|ffi|||ddSr:rr))rrrrrr*rr*rr,r-copy_tmasz!tma_get_copy_fn..copy_tmacs(tjfi|||ddSr:r5)rrrr6r,r-copy_tma_single_stages(z.tma_get_copy_fn..copy_tma_single_stage)r rrrr r!r"smemrvr tma_partitionr.r4r)r*r2r3rrr4r#rrr* src_is_smem smem_tensor gmem_tensorgroup_rank_smemgroup_rank_gmemrgr7r8r,r6r-tma_get_copy_fns2      rApipelinecsdtjjffdd }|S)Nproducer_statecs"d||j|d|dS)N)rrr$r,)indexproducer_get_barrier)rrCrr)rBr,r-r#s z%tma_producer_copy_fn..copy_fn)rirB PipelineState)r)rBrr,rFr-tma_producer_copy_fn"srH 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) NrrcrVre)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) NrVTrNrre up_to_rankNNrrc)rrrr) r rr?r rirangerjr6 tiled_divide append_onesrhr)) rrrtApA_k limit_k_curkmA_currmA_rowki cols_per_threadelems_per_loadis_even_m_smemrMmA_km_idxt0AcAratApA_mtAsArI tile_shape_mkr,r-rYs& &z%gather_m_get_copy_fn..copy_fnF)rrhrCr6r.slice_tiler_mnr rLrBrrr?r rirTrlogical_dividebool) rIrJrrKrLrMcArows_per_threadrrrr,r]r-gather_m_get_copy_fn.s6       *rncsJd\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)NNrrcrVrureTrNFrrcsd}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)NrVTrNr) r rr?r rirT consumer_waitrr sync_warpr)consumer_release) a_prefetch_pipelinerra_prefetch_consumer_staterrWrXrY sAIdx_currqr)r^rMsAIdxrcrarfr,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)NrerPrVrS)r r prepend_onesrirjr6r)) rrr|rrqrW tApA_k_predrYr)ratAmArdrerIr,r-rs z%gather_k_get_copy_fn..copy_fnrg)r r!rr"gmemr9rhrCr.rir6rLrBrrr?r rirTthr_idxrj flat_dividerkrr>)rIrJrrKrLrMr`r_rlrmrthreads_per_colr1rsr{rr,) r^rrrMrzrcrarrdrerIrfr-gather_k_get_copy_fnrs^       . rrzwarp_idx num_warpscstj|dgd}t|d||ddksJ|}ttjtjtddt|td} | |} | |} | |t | t |} t t | |ddtjffd d } | S) NrrcrRrsrIrJ)r r$c s|}tjtjdgdddD]3fddtdD}dd|fj}tj||||Wdn1s?wYqdS)NrVrcTrNcsg|]}|fqSr,r,)rv)r tSR_rAIdxr,r-rrz=gather_m_get_tma_copy_fn..copy_fn..rs)rirTrrhrrr))rrr$rrrrrtile_Ktma_gather4_load_fn)rr-rs z)gather_m_get_tma_copy_fn..copy_fn)rrhrYrQrOrPrrXrrr<r rr"r )rrJrrzrrr tile_M cta_group copy_AIdx_s2rwarp_copy_AIdx_s2r tSR_sAIdxrrr,rr-gather_m_get_tma_copy_fns"    rrg)rVF)rF)FN)FFN)FF)rV)Xtypingrrrrr functoolsrri cutlass.cuterrr r r cutlass.cute.nvgpur r rcutlass.cute.nvgpu.tcgen05.mmarcutlass.cutlass_dslrcutlass.pipelinecutlass._mlir.dialectsr cutlass._mlirrrrr TiledCopyr>rkr.r9r<Shaper@ThrCopyrENumericintCopyAtomrTr)r\r`jitrorwrxr{rrrr rrcorerrutils LayoutEnumr ConstexprrrTiledMmarrrrrrrr r"r1CoordLayoutrArB PipelineAsyncrHrnrrr,r,r,r-sP                  D !      #            r    0  C o