^j|+~ddlZddlmZmZmZddlZddlmZddlmZm Z m Z ddl m Z ddl mZmZmZddlmZmZeddddej(d ej*d ej,fd Zej0efd Zedddd ej,de d e fdZeddddeeze zej8zd ej,dej,dejj:d df dZeddddee zdee zdee zdee zd ej,dej,dejj:d dfdZeddddeeefdeeefd efdZ eddddeezd efdZ!eddddeezd e fdZ"ej0dej(deej(dejFd dfd Z$edddd!Z%eddddeded ej8fd"Z&edddd#ej8d eeeffd$Z'ej0d.de d%ee d e fd&Z(edddde)e zd'ej,d e fd(Z*edddde)e zd'ej,d e fd)Z+e d/dej,d*ej,d dfd+Z,edddd ej*d,ej(d ej(fd-Z-y)0N)OptionalTupleUnion)Float32Int32 const_expr)arith)llvmnvvmvector)T dsl_user_oplocipxcoordreturncb|jtj||j||zS)Nr)iteratorcutecrd2idxlayout)rrrrs d/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_vendor/quack/utils.py elem_pointerrs$ :: UAHH#"E EEctt|tjr2|tj|tj ddS|S)Nr)r isinstancerPointer make_tensor make_layout)rdtypes rload_scalar_or_pointerr$sB*Q -.T%%a)9)9!)<=a@AArsmem_ptrpeer_cta_rank_in_clusterc |j||j}ttjt j ||jgddddS)zMMap the given smem pointer to the address at another CTA rank in the cluster.rz$mapa.shared::cluster.u32 $0, $1, $2;z=r,r,rFhas_side_effectsis_align_stack)tointir_valuerr inline_asmr i32)r%r&rr smem_ptr_i32s rset_block_rankr0s_ >>cb>1:: ? 2 "    rvalmbar_ptrc.t||||j}t||||j}tt|tr t |}t|t t tjfsJdt dt dtjdit|}t dt dtjdit|} tjd||j|||gd |d d | d d dy)Nrz$val must be Float32, Int32, or Int64f32s32s64frlz6st.async.shared::cluster.mbarrier::complete_tx::bytes.z [$0], $1, [$2];zr,z,rTFr() r0r,rrfloatrrcutlassInt64typer r-) r1r%r2r&rrremote_smem_ptr_i32remote_mbar_ptr_i32suffix constraints rstore_shared_remoterB.s)*hj)*hj*S%()cl cGUGMM: ;c=cc ;ueUGMM5 A$s) LF3sGMM3?S JJOO cllsrl:r?r#r@rAs rstore_shared_remote_x4rIMsw)*hj)*hj dWe, -N/NN -!$0GeEueU +E 2F3s+E2JOO   $K SR 0 $K SR 0 $K SR 0 $K SR 0   HhhhhDDJ8L   zl!J from N scalar DSL values. Example: make_vector(cutlass.Uint32, v0, v1) -> <2 x i32> MLIR vector r)irr)positionrr) cutlass._mlirrelen mlir_type VectorTypegetr mlir_undef enumerater insertelementr,_arithconstantr r.) elem_typerrvaluesrenmlir_tyvec_tyvecivs r make_vectorrys ! F A!!G ]]  sG ,F //&cb 1C&! 1"" aL ! !cb ! 1 __QUUWaSR@    Jrc tjtjdtj|j |j f||}tj tjdtj |}tjtj|gdg||}|S)NrZrrrdynamic_positionstatic_positionrr) r from_elementsr r4r,bitcasti64r;r<extract)rJrKrr vec_f32x2 vec_i64x1ress r f32x2_to_i64rs$$ AEEGqzz|QZZ\:Iqxx15573Y?I --y2sPSXZ[ C Jrcc tjtjdtj|j f||}tj tjdtj |}ttj|gdg||}ttj|gdg||}||fS)NrrrZrr{) r r~r rr,rr4rr)rrrrrres0res1s r i64_to_f32x2rs$$QXXa%9AJJL?PSXZ[Iqxx15573Y?I y2sPSXZ[ D y2sPSXZ[ D :rlanec\t|durtjj}t j t tjtjjD]4}d|z}tjj||d}||k\s0||z }6|S)Nrr)offsetmask_and_clamp) rrarchlane_idxr;r]intmathlog2 WARP_SIZEshuffle_sync_up)r1rrwr partial_sums rwarp_prefix_sumrs$$,yy!!#  $ $S4993F3F)G%H Iaii//FST/U 6> ; C  Jrgmem_ptrcddlm}|jdk(rs|jdk(rdt j t jtjj|jt|jSt j tjj|jt|jSNr)rNrM )ropptrrJ)rrrJ) r;rNrOminorr atomicrmwr r. AtomicOpKindINCllvm_ptrrr,rJrrrrNs ratomic_inc_i32r$RL$6$6!$;~~D--11x7H7HERSHL]L]L_  ~~  $$(*;*;uQx?P?P?R  rcddlm}|jdk(rs|jdk(rdt j t jtjj|jt|jSt j tjj|jt|jSr) r;rNrOrr rr r.rADDrrr,rs ratomic_add_i32rrrclc_response_ptrcf|j}|j}tj||||y)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 rN)rr clusterlaunchcontrol_try_cancel)r2rrr mbar_llvm_ptrclc_response_llvm_ptrs rissue_clc_query_nomulticastr&s5$%%M,55((  rtensorc@t|jtjsJtj|j t ||j|j|jj}tj||jS)N) assumed_align) rrrr make_ptr element_typerr+memspace alignmentr!r)rrrrnew_ptrs rdomain_offset_alignedrBss foot|| 44 4mmVU#))+oo// G   GV]] 33r)N)NN).rtypingrrrr; cutlass.cuterrrrcutlass._mlir.dialectsr ror r r cutlass.cutlass_dslr rTensorCoordr rjitr$r0r:r<IntrBrIrPrSrUNumericrcryrrrrrrrrrrrs) ))..255. ;?DFDKKF FT\\F F$+  DHTll6;  "    5 7== 0llll#kkoo   <   + E/+ E/+ E/+ E/ + ll + ll +#kkoo+ + +\ DHTE%. !eE7N&;V] * $(T EGO g    $(T EGO e   Q4;;Qht{{&;QQZ^Q Q0 (, , 04G'--  *.4 GMM E'7BR