L iETddlmZddlmZddlmZddlmcmZddl m Z ddl m Z m Z mZmZmZmZmZmZmZddlmZddlmZmZmZmZmZmZmZdd lm Z er dd l!m"Z"dd lm#Z#ed GddZ$ed GddZ%d!dZ& d"dZ'ed GddZ(e(ddejRe(ddejTgZ+ d#dZ,ed$dZ-ed GddZ.ed$dZ/ d%dZ0ed$d Z1y)&) annotations) dataclass) TYPE_CHECKINGN) translate) BaseCTypeBindingCTypeExpr NamedCTypeopmath_tscalar_tStructuredImplSignatureVectorizedCType)with_native_function)ArgumentBaseTyBaseType DispatchKeyNativeFunctionsGroup ScalarTypeUfuncKey) OrderedSet)Sequence)UfunctorBindingsT)frozenc\eZdZUded<ded<ded<ddZddZdd Zdd Zdd Zdd Z y )UfunctorSignaturergz int | Nonescalar_tensor_idxstrnamecbtj|j|jtS)N)rr )ufuncufunctor_argumentsrrr selfs Y/mnt/ssd/data/python-lab/Trading/venv/lib/python3.12/site-packages/torchgen/dest/ufunc.py argumentszUfunctorSignature.argumentsIs&'' FFd&<&z0UfunctorSignature.decl_fields..Xs%F1AFF81QVVHA.F&()joinr/r%s r' decl_fieldszUfunctorSignature.decl_fieldsWsyyF FFFr)cdjd|jjD}djd|jjD}|jd|d|dS)N, c3<K|]}|jywr1declr9as r'r;z5UfunctorSignature.inline_defn_ctor..[sE!QVVXEc3TK|] }|jd|jd"yw)z_()Nr!rDs r'r;z5UfunctorSignature.inline_defn_ctor..^s%Sxr!&&3Sr<(z) : z {})r=r(r,r!)r&args_strinit_strs r'inline_defn_ctorz"UfunctorSignature.inline_defn_ctorZsc99Et~~/?/D/DEE99ST^^=M=R=RSS))AhZtH:U;;r)cdjd|jjD}|jj d|dS)Nr@c3<K|]}|jywr1rBrDs r'r;z/UfunctorSignature.decl_apply..bsF!QVVXFrFz operator()(z) const)r=r(applyr2cpp_type)r&rKs r' decl_applyzUfunctorSignature.decl_applyasI99Ft~~/?/E/EFF##%..01hZwOOr)N)returnrrSz list[Binding])rSr rSr ) __name__ __module__ __qualname____annotations__r(r/r2r>rMrRr)r'rrCs4!! I G# G<Pr)rc<eZdZUded<ded<ded<d dZd dZy ) UfuncSignaturerrr r!r compute_tcXtj|j|jS)N)r])r#ufunc_argumentsrr]r%s r'r(zUfuncSignature.argumentsls$$TVVt~~FFr)c |jddjdt||jDdS)NrJr@c34K|]}|jywr1exprrDs r'r;z&UfuncSignature.call..ps'Y1'YrHr!r=rr(r&ctxs r'callzUfuncSignature.callos8))Adii'Y #t~~GW8X'YYZZ[\\r)NrT)rgzSequence[Binding | Expr]rSr )rVrWrXrYr(rhrZr)r'r\r\fs IG]r)r\ctd|jjjjD}|dk(S)Nc3VK|]!}|jjsd#yw)N)r8is_tensor_likerDs r'r;z.s$qvv?T?T?Vs)))sum functionalfuncr( flat_non_out)r num_tensorss r')eligible_for_binary_scalar_specializationrss<<<$$..;;K ! r)ci}g}|jj}tjdtjdtj di}t |r0tjtjtj g}nEtj g}tjtjfD]}||vrJd|d|D]}||vrHt|||||j}||jD]}||j|i|<Pd} t} tjtjfD]D} | |vr| || j} n| || jk(sJd| || jz} F| J|d| } t|||| }| D]}||j|i|<t|d| tt  } |j#|j%j&z}|j)d |jd |j+d |j-d |j/d| j1|d |dj3|fS)Nrkrz cannot use z on non-binary function)rr!z0ScalarOnly and Generic must have same ufunc namer+ufunc::r!r]z% template struct z3 { using opmath_t = at::opmath_type; z z __device__ z { return z ; } }; r4)outufunc_inner_looprCUDAFunctorOnSelfCUDAFunctorOnOther CUDAFunctorrsrr!supported_dtypes setdefaultr ScalarOnlyGenericr\rr r/r(rPappendr>rMrRrhr=)r ufunctor_sigs ufunctorsloopsscalar_tensor_idx_lookupkeysk ufunctor_sigdtype ufunc_namer|lkr! ufunc_sig apply_ctxs r'compute_ufunc_cuda_functorsrsJLMI EE " "E""A##Qd  13  & &  ' '    $$%,,h.I.IJ LAE> K[3J#K K> L 8  :,%=a%@uQx}}Lq22 F9E ((3A6 F  3=<&&(8(89 ;B!"2Y^^ "U2Y^^3F3 b : : :  ;%%%Aj\"( !9!!<4 & BE5AM $ $UB / 2 B# gj\*i6I !'')L,B,B,D,J,JJ   "#$%%' () NN9 % &'  Y8 t $))I. ..r)c,eZdZUded<ded<ded<y) BinaryScalarSpecializationConfigint scalar_idxr ctor_tensorr ufunc_keyN)rVrWrXrYrZr)r'rrsOr)rr&)rrrrkotherc d}|dz }tD]}|j|vr||j}|jdz}t|}|j t d|dt |jttdjdt||jjD} |d |d |jd | d |d z }|tj }djdt||jjD} |d|jd| dz }|S)Nz+using opmath_t = at::opmath_type;zif (false) {} rkziter.scalar_value(rH)rcr8r@c34K|]}|jywr1rbrDs r'r;z0compute_ufunc_cuda_dtype_body.. s, AFF, rdzelse if (iter.is_cpu_scalar(z)) { z ufunctor(z); iter.remove_operand(z"); gpu_kernel(iter, ufunctor); }c34K|]}|jywr1rbrDs r'r;z0compute_ufunc_cuda_dtype_body..s((rdz else { gpu_kernel(iter, z (z )); } )!BinaryScalarSpecializationConfigsrrlistrr r rrr r=rr(r,r!rr{) rr inner_loops parent_ctxbodyconfigrrrgufunctor_ctor_exprs_strs r'compute_ufunc_cuda_dtype_bodyrs 9D D3   ; . "6#3#34 &&* %)$4 3J>r)ct|jjjjjdS)N_kernelrr%s r'rzStubSignature.kernel_name\s1dff'',,116678@@r)ct|jjjjjdS)N_fnrr%s r' type_namezStubSignature.type_name`s1dff'',,116678<.is8RA8RrdrH)r(r=)r&cpp_argss r'r8zStubSignature.typegs.>>#.tyy8R8R/R.SSTUUr)c<d|jd|jdS)NzDECLARE_DISPATCH(r@rH)rr!r%s r'rzStubSignature.dispatch_declks "4>>"2"TYYKqAAr)c"d|jdS)NzDEFINE_DISPATCH(rHrIr%s r' dispatch_defnzStubSignature.dispatch_defnns!$))A..r)cpd|jddjd|jDdS)Nzvoid z(TensorIteratorBase& iter, r@c3<K|]}|jywr1)rrDs r'r;z,StubSignature.kernel_defn..rsMq[\affhMqrFrH)rr=r(r%s r'rzStubSignature.kernel_defnqs:t''((CDIIMq`d`n`n`pMqDqCrrsttr)cBd|jd|jS)Nzusing  = )rr8r%s r'rzStubSignature.type_defnts 's499;-88r)c |jddjdt||jDdS)Nz(device_type(), *this, r@c34K|]}|jywr1rbrDs r'r;z%StubSignature.call..ys=oaff=ordrHrerfs r'rhzStubSignature.callxs=))3DII=oiX[]a]k]k]mNn=o4o3ppqrrr)c |jddjdt||jDdS)Nz(*this, r@c34K|]}|jywr1rbrDs r'r;z,StubSignature.direct_call..}s5gaff5grdrH)rr=rr(rfs r'rzStubSignature.direct_call|s?""#8DII5giPSUYUcUcUeFf5g,g+hhijjr)NrUrT)rgSequence[Binding]rSr )rVrWrXrYpropertyr!rrr(r8rrrrrhrrZr)r'rrTsq ??AA==,VB/u9skr)rc Bt|}t|tj|tj }d|j d|jd|jd|jd|j|jd S)Nr4rz; rr) rrr#rrCPUrrrrrhr()rrrs r'compute_ufunc_cpursQH !!U%6%6q+//%J KC      ==!"#  r)c, tj|vsJ|d|j|jtjtjhksJ|tj}d}tj|vr|tj}g}g |D]}t |j t r1|j jttjk7rN|jd|jd|jd jtd|jt|jjt!t"||D]}t |j t r1|j jttjk7rN|jd|jd|jd jtd |jt|jjt%t!t"g}g} |j&j(j*j,D]} | jj/s| jttj0k(sJ|jt3| jt| jt!t"| || jt3| jt| jt%t!t"| d fd } d j5|} |hd | d dj5d|Dd|j7| |ddj5d| Dd|j7| | d Sd | ddj5d|Dd|j7| |dS)Nr@zauto _s_rz.to();_s_zauto _v_z$ = at::vec::Vectorized(_s_z);_v_)r!nctypeargumentcPg}|j|j||Sr1)extend)r.rrgs r'with_ctxz.compute_ufunc_cpu_dtype_body..with_ctxs#"$    r)r4z cpu_kernel_vec(iter, [=](c3<K|]}|jywr1rBr9r.s r'r;z/compute_ufunc_cpu_dtype_body..3a3rFz ) { return z ; }, [=](c3<K|]}|jywr1rBrs r'r;z/compute_ufunc_cpu_dtype_body..s0a0rFz; } ); z cpu_kernel(iter, [=](c3<K|]}|jywr1rBrs r'r;z/compute_ufunc_cpu_dtype_body..rrF)r.rrSzlist[Expr | Binding])r CPUScalarr CPUVector isinstancerrr8rrScalarrr!r r rrr rrorpr(rqrlTensorrr=rh)rrrr scalar_loopvec_looprr.scalar_bindings vec_bindingsrErbody_strrgs @r'compute_ufunc_cpu_dtype_bodyrs    ,Nr+:J:J:L9M.NN ,    ("4"4h6H6H!I II Ih001KH[(x112 D C Y ajj( + 8 MMD 1   hqvvhc!&&1ABC 4#affX 188==)HBU(VWX Y A!**h/AJJOOx H5 KK166("FqvvhbQ  JJ!&&Nqxx}}oi>Q.RS  "OL \\   ( ( 5 5vv$$& vv&--0000 VV!!&&)H*=>       %affoi>Q.RS ( yyH   yy3?334LAQAQRZ[jRkAl@mn yy0<001hmmHUaLb>c=de     yy3?334LAQAQRZ[jRkAl@mn r)crt|}|jj}i}tjtj fD]4}g}||vr|j |tj|vr1|tjur|j tjtj|vr|j tj|D]}||jD]}|tjurtt}n1|tj urttt}nt|j|i} || vsmt|d||j || |<7g} |j#D]7\}} | j d|dt%||| |j'd9dj)| } d|j+d|j d | d |j-d |j/d |j d |j0dS)Nrurvrrrr4z namespace { rrz# ); } } // anonymous namespace rz REGISTER_DISPATCH(rz) )rrwrxrrrrr~rr|rr rAssertionErrorr}r\r!rrr(r=rrrr) rrr ufunc_sigsrlksrrr]rrrs r'compute_ufunc_cpu_kernelrsGQH EE " "ECEJ  ("4"4 5 : JJqM   % 'A1C1C,C JJx** +   u $ JJx'' ( Br33 *** )( 3I(,,, / (0C DI((#-#8#8#C ,,*8'%)..)9 :i+$Q'  4K#-#3#3#5  ""')!!U,P>P>RSTU    ii ,O ,,4MM?;   ==/X%9%9$:; r))rrrSbool)rrrSz?tuple[dict[ScalarType, dict[UfuncKey, UfunctorSignature]], str]) rrrrrz!dict[UfuncKey, UfunctorSignature]rrrSr )rrrSr ) rrrrrzdict[UfuncKey, UfuncSignature]rrrSr )2 __future__r dataclassesrtypingrtorchgen.api.ufuncapir#torchgen.api.translatertorchgen.api.typesrrr r r r r rrtorchgen.contextrtorchgen.modelrrrrrrrtorchgen.utilsrcollections.abcrrrr\rsrrrzryrrrrrrrrZr)r'rs"! "",   2&(3@ $PPPD $ ] ] ]<P/P/DP/f $%-- %,, %!,, ,3," ,  ,^&&` $(k(k(kV  ZZ Z0Z" Z  Zz;;r)