o c۷i'@sddlZddlmZmZmZddlZddlmZddlmZm Z m Z ddl m Z m Z ddlmZmZmZe ddddejdejd ejfd d ZejdeejBd efd d Ze ddddejde d e fddZe ddddeeBe BejBdejdejdejjd df ddZe ddddee Bdee Bdee Bdee Bdejdejdejjd dfddZe ddddeeefdeeefd efddZe ddddeeBd efd d!Ze ddddeeBd e fd"d#Z ejd$ejd%eejd&ej!d dfd'd(Z"e ddddeded ejfd)d*Z#e dddd+ejd eeeffd,d-Z$ejd9de d.ee d e fd/d0Z%e dddde&e Bd1ejd e fd2d3Z'e dddde&e Bd1ejd e fd4d5Z(e  d:dejd6ejd dfd7d8Z)dS);N)OptionalTupleUnion)Float32Int32 const_expr)T dsl_user_op)llvmnvvmvectorlocipxcoordreturncCs|jtj||j||dS)Nr )iteratorcutecrd2idxlayout)rrrrrA/home/ubuntu/vllm_env/lib/python3.10/site-packages/quack/utils.py elem_pointersrcCs<tt|tjrtt|tddSt|tsJ|S)Nr)r isinstancerPointerr make_tensor make_layout)rrrrload_scalar_or_pointersrsmem_ptrpeer_cta_rank_in_clusterc Cs>|j||d}ttjt||gddddtjjdS)zMMap the given smem pointer to the address at another CTA rank in the cluster.r z$mapa.shared::cluster.u32 $0, $1, $2;z=r,r,rFhas_side_effectsis_align_stack asm_dialect) tointir_valuerr inline_asmri32 AsmDialectAD_ATT)r r!rr smem_ptr_i32rrrset_block_ranks r-valmbar_ptrc Cst||||d}t||||d}tt|trt|}t|tttjfs,Jdtdtdtjdit |}tdtdtjdit |} t j d||j||d|gd |d d | d d dt j j ddS)Nr z$val must be Float32, Int32, or Int64f32s32s64frlz6st.async.shared::cluster.mbarrier::complete_tx::bytes.z [$0], $1, [$2];zr,z,rTFr")r-r'rrfloatrrcutlassInt64typer r(r*r+) r.r r/r!rrremote_smem_ptr_i32remote_mbar_ptr_i32suffix constraintrrrstore_shared_remote/s0    r>val0val1val2val3cCst||||d} t||||d} t|ttfsJdt|tr&tnt} tdtdi| } tdtdi| } tjd| | | |j||d| |j||d| |j||d| |j||dgd| d| d | d | d | d | d d| d| d| d| ddtjjddS)Nr zval must be Float32, or Int32r0r1r3r4z { .reg .v4 .z abcd; mov.z abcd.x, $2; mov.z abcd.y, $3; mov.z abcd.z, $4; mov.zG abcd.w, $5; st.async.shared::cluster.mbarrier::complete_tx::bytes.v4.z [$0], abcd, [$1]; } zr,r,,TFr") r-r'rrrr r(r*r+)r?r@rArBr r/r!rrr:r;dtyper<r=rrrstore_shared_remote_x4OsR  rEabcCs|ttjjdkr%ttjtt|j ||dt|j ||d||dSttjt|j ||dt|j ||d||dS)N r ) r7r CUDA_VERSIONmajorrr fminrr0r')rFrGrrrrrrKs$ rKc Cs4ttjtt|j||dgddddtjjdS)Nr zsqrt.approx.f32 $0, $1;z=f,fFr")rr r(rr0r'r*r+rFrrrrrsqrtrMc Cs4ttjtt|j||dgddddtjjdS)Nr zcvt.rpi.ftz.s32.f32 $0, $1;z=r,fFr") rr r(rr)rr'r*r+rLrrrceilrNrOtXsXtXpX fill_valuec Cst|d}||t|jddD]5}t|jdD]*}t|dur=||d|fs}tjj ||dd}||kr.||7}q|S)Nrr)offsetmask_and_clamp) rrarchlane_idxr7rVintmathlog2 WARP_SIZEshuffle_sync_up)r.rlirm partial_sumrrrwarp_prefix_sums  rxgmem_ptrcCbddlm}|jdkr"|jdkr"tjttjj |j t | dStjtjj |j t | dSNr)rIrH )rfopptrrF)r}r~rF) r7rIrJminorr atomicrmwrr) AtomicOpKindINCllvm_ptrrr'rFryrrrIrrratomic_inc_i32 rcCrzr{) r7rIrJrr rrr)rADDrrr'rrrratomic_add_i32rrclc_response_ptrcCs"|j}|j}tj||||ddS)aN The clusterlaunchcontrol.try_cancel instruction requests atomically cancelling the launch of a cluster that has not started running yet. It asynchronously writes an opaque response to shared memory indicating whether the operation succeeded or failed. On success, the opaque response contains the ctaid of the first CTA of the canceled cluster. :param mbar_ptr: A pointer to the mbarrier address in SMEM :type mbar_ptr: Pointer :param clc_response_ptr: A pointer to the cluster launch control response address in SMEM :type clc_response_ptr: Pointer r N)rr clusterlaunchcontrol_try_cancel)r/rrr mbar_llvm_ptrclc_response_llvm_ptrrrrissue_clc_query_nomulticasts r)N)NN)*rrtypingrrrr7 cutlass.cuterrrrcutlass.cutlass_dslrr cutlass._mlir.dialectsr r r TensorCoordrrjitrr-r6r8Intr>rErKrMrONumericr\rgrkrxrqrrrrrrrs &    /0  &" &  &&