
    '}hE              
          d dl mZ d dlmZmZmZmZmZmZ d dl	m
c mZ d dlmZ d dlmZmZmZmZmZmZmZmZmZ d dl	mZ d dlmZ d dlmZmZmZmZm Z m!Z!m"Z" d d	l#m$Z$  ed
       G d d             Z% ed
       G d d             Z&de de'fdZ(de deee!ee"e%f   f   e)f   fdZ* ed
       G d d             Z+ e+d de"jX                         e+dde"jZ                        gZ.de de!dee"e%f   dee   de)f
dZ/ede de)fd       Z0 ed
       G d d              Z1ede de)fd!       Z2de de!dee"e&f   dee   de)f
d"Z3ede de)fd#       Z4y)$    )	dataclass)DictListOptionalSequenceTupleUnionN)	translate)		BaseCTypeBindingCTypeExpr
NamedCTypeopmath_tscalar_tStructuredImplSignatureVectorizedCType)UfunctorBindings)with_native_function)ArgumentBaseTyBaseTypeDispatchKeyNativeFunctionsGroup
ScalarTypeUfuncKey)
OrderedSetT)frozenc                       e Zd ZU eed<   ee   ed<   eed<   defdZ	de
e   fdZdefdZdefdZdefd	Zdefd
Zy)UfunctorSignaturegscalar_tensor_idxnamereturnc                 b    t        j                  | j                  | j                  t              S )N)r"   r   )ufuncufunctor_argumentsr!   r"   r   selfs    R/var/www/html/test/engine/venv/lib/python3.12/site-packages/torchgen/dest/ufunc.py	argumentszUfunctorSignature.argumentsA   s&    ''FFd&<&<x
 	
    c                     | j                         j                  D cg c]   }|j                  |j                   d      " c}S c c}w )N_)r+   ctorrenamer#   )r)   bs     r*   fieldszUfunctorSignature.fieldsF   s6    040@0E0EF1AFF81&FFFs   %Ac                      t        t              S N)r   r   r(   s    r*   returns_typezUfunctorSignature.returns_typeJ   s     ""r,   c                 N    dj                  d | j                         D              S )N
c              3   T   K   | ]   }|j                    d |j                   d " yw) ;N)typer#   ).0fs     r*   	<genexpr>z0UfunctorSignature.decl_fields.<locals>.<genexpr>P   s%     F1AFF81QVVHA.F   &()joinr2   r(   s    r*   decl_fieldszUfunctorSignature.decl_fieldsO   s    yyFFFFr,   c                     dj                  d | j                         j                  D              }dj                  d | j                         j                  D              }| j                   d| d| dS )N, c              3   <   K   | ]  }|j                           y wr4   declr<   as     r*   r>   z5UfunctorSignature.inline_defn_ctor.<locals>.<genexpr>S   s     E!QVVXE   c              3   T   K   | ]   }|j                    d |j                    d " yw)z_()Nr#   rG   s     r*   r>   z5UfunctorSignature.inline_defn_ctor.<locals>.<genexpr>V   s%     Sxr!&&3Sr?   (z) : z {})r@   r+   r/   r#   )r)   args_strinit_strs      r*   inline_defn_ctorz"UfunctorSignature.inline_defn_ctorR   sc    99Et~~/?/D/DEE 99ST^^=M=R=RSS))AhZtH:U;;r,   c                     dj                  d | j                         j                  D              }| j                         j	                          d| dS )NrC   c              3   <   K   | ]  }|j                           y wr4   rE   rG   s     r*   r>   z/UfunctorSignature.decl_apply.<locals>.<genexpr>Z   s     F!QVVXFrI   z operator()(z) const)r@   r+   applyr5   cpp_type)r)   rN   s     r*   
decl_applyzUfunctorSignature.decl_applyY   sI    99Ft~~/?/E/EFF##%..01hZwOOr,   N)__name__
__module____qualname__r   __annotations__r   intstrr   r+   r   r   r2   r   r5   rA   rP   rU    r,   r*   r    r    ;   sj    }$
I
+ 

GW G#e #
GS G<# <PC Pr,   r    c                   ^    e Zd ZU eed<   eed<   eed<   dee   fdZ	de
eeef      defdZy)	UfuncSignaturer!   r#   	compute_tr$   c                 X    t        j                  | j                  | j                        S )N)r_   )r&   ufunc_argumentsr!   r_   r(   s    r*   r+   zUfuncSignature.argumentsd   s    $$TVVt~~FFr,   ctxc           
          | j                    ddj                  d t        || j                               D               dS )NrM   rC   c              3   4   K   | ]  }|j                     y wr4   exprrG   s     r*   r>   z&UfuncSignature.call.<locals>.<genexpr>h   s     'Y1'Y   rK   r#   r@   r
   r+   r)   rb   s     r*   callzUfuncSignature.callg   s8    ))Adii'Y	#t~~GW8X'YYZZ[\\r,   N)rV   rW   rX   r   rY   r[   r   r   r   r+   r   r	   r   rj   r\   r,   r*   r^   r^   ^   sI    
IG4= G]w}!56 ]3 ]r,   r^   r!   r$   c                     t        d | j                  j                  j                  j                  D              }|dk(  S )Nc              3   V   K   | ]!  }|j                   j                         sd  # yw)   N)r;   is_tensor_likerG   s     r*   r>   z<eligible_for_binary_scalar_specialization.<locals>.<genexpr>~   s$      qvv?T?T?Vs   ))   )sum
functionalfuncr+   flat_non_out)r!   num_tensorss     r*   )eligible_for_binary_scalar_specializationru   }   s<     <<$$..;; K !r,   c                    i }g }| j                   j                  }t        j                  dt        j                  dt        j
                  d i}t        |       r0t        j                  t        j                  t        j
                  g}nEt        j
                  g}t        j                  t        j                  fD ]  }||vrJ d| d        |D ]  }||v rHt        | ||   ||   j                        }||   j                  D ]  }||j                  |i       |<    Pd }	t               }
t        j                  t        j                  fD ]D  }||vr|	||   j                  }	n|	||   j                  k(  sJ d       |
||   j                  z  }
F |	J | d|	 }t        | ||   |      }|
D ]  }||j                  |i       |<    t        | d|	 t        t               	      }|j#                         |j%                         j&                  z   }|j)                  d
|j                   d|j+                          d|j-                          d|j/                          d|j1                  |       d        |dj3                  |      fS )Nrm   r   zcannot use z on non-binary function)r"   r#   z0ScalarOnly and Generic must have same ufunc namer.   ufunc::r#   r_   z%
template <typename scalar_t>
struct z3 {
  using opmath_t = at::opmath_type<scalar_t>;
  z
  z
  __device__ z {
    return z	;
  }
};
r7   )outufunc_inner_loopr   CUDAFunctorOnSelfCUDAFunctorOnOtherCUDAFunctorru   r    r#   supported_dtypes
setdefaultr   
ScalarOnlyGenericr^   r   r   r2   r+   rS   appendrA   rP   rU   rj   r@   )r!   ufunctor_sigs	ufunctorsloopsscalar_tensor_idx_lookupkeyskufunctor_sigdtype
ufunc_namer~   lkr#   	ufunc_sig	apply_ctxs                  r*   compute_ufunc_cuda_functorsr      s    JLMIEE""E""A##Qd 
 13&&''  
 $$%,,h.I.IJ 	LAE>K[3J#KK>	L 8
 :,%=a%@uQx}}L q22 F9E((3A6F 
3=<&&(8(89 
	;B!"2Y^^
 %)..0FEF0b	 : ::
	; %%%Aj\"(!9!!<4
 & 	BE5AM$$UB/2	B #gj\*i6I
	 !'')L,B,B,D,J,JJ	    "# $%%'( )NN9%& '
	
Y8
t $))I...r,   c                   ,    e Zd ZU eed<   eed<   eed<   y) BinaryScalarSpecializationConfig
scalar_idxctor_tensor	ufunc_keyN)rV   rW   rX   rZ   rY   r[   r   r\   r,   r*   r   r      s    Or,   r   r)   )r   r   r   rm   otherr   inner_loops
parent_ctxc                    d}|dz  }t         D ]  }|j                  |vr||j                     }|j                  dz   }t        |      }|j	                  t        d| dt        |j                  t        t                                 dj                  d t        ||j                         j                        D              }	|d	| d
|j                   d|	 d| d	z  } |t        j                      }dj                  d t        ||j                         j                        D              }	|d|j                   d|	 dz  }|S )Nz+using opmath_t = at::opmath_type<scalar_t>;zif (false) {}
rm   ziter.scalar_value<opmath_t>(rK   )rf   r;   rC   c              3   4   K   | ]  }|j                     y wr4   re   rG   s     r*   r>   z0compute_ufunc_cuda_dtype_body.<locals>.<genexpr>  s      ,
AFF,
rg   zelse if (iter.is_cpu_scalar(z)) {
  z<scalar_t> ufunctor(z);
  iter.remove_operand(z");
  gpu_kernel(iter, ufunctor);
}c              3   4   K   | ]  }|j                     y wr4   re   rG   s     r*   r>   z0compute_ufunc_cuda_dtype_body.<locals>.<genexpr>  s      ((rg   z
else {
  gpu_kernel(iter, z<scalar_t>(z
));
}
    )!BinaryScalarSpecializationConfigsr   r   listr   r   r   r   r   r   r@   r
   r+   r/   r#   r   r}   )
r!   r   r   r   bodyconfigr   r   rb   ufunctor_ctor_exprs_strs
             r*   compute_ufunc_cuda_dtype_bodyr      s    9DD3 ;."6#3#34&&*
 +/z*:

3J<qA 2 2Ih4GH	
 #')) ,
%c<+A+A+C+H+HI,
 #
 	 'L ))*A)B C!l # 	)6 x334L"ii (!*l.D.D.F.K.KL(  	  %%&k2I1J K D
 Kr,   c                 r   t        |       \  }}t        | t        j                  | t        j
                              }g }|j                         D ]7  \  }}|j                  d| dt        | |||j                                d       9 dj                  |      }t        |       }d| d|j                          d|j                          d|j                          d|j                   d	| d
|j                   d|j                   d|j!                          d|j#                  |j                                dS )N"
AT_DISPATCH_CASE(at::ScalarType::,
  [&]() {
    
  }
)
r7   z

;
;

. {
  AT_DISPATCH_SWITCH(iter.common_dtype(), "",
    z
  );
}
REGISTER_DISPATCH(, &z);

 {
  ;
}
)r   r   r&   kernel_namer   CUDAitemsr   r   r+   r@   StubSignature	type_defndispatch_declkernel_defnr#   defndirect_call)	r!   r   r   sigdtype_casesr   inner_ufunc_sigsdtype_cases_strstub_sigs	            r*   compute_ufunc_cudar     sh     ;1=M9 "!U%6%6q+:J:J%K
LCK#0#6#6#8 	
""' )"1e-=s}}OP Q	
	
 ii,OQH
 	  	  	  ,,/HH: 6  ==/X%9%9$: ; () * r,   c                       e Zd ZU eed<   edefd       Zedefd       Zedefd       Z	de
e   fdZdefdZdefdZdefd	Zdefd
ZdefdZdee   defdZdee   defdZy)r   r!   r$   c                     t        | j                  j                  j                  j                  j                         dS )N_stubr[   r!   rq   rr   r#   r(   s    r*   r#   zStubSignature.nameP  s1    dff'',,116678>>r,   c                     t        | j                  j                  j                  j                  j                         dS )N_kernelr   r(   s    r*   r   zStubSignature.kernel_nameT  s1    dff'',,116678@@r,   c                     t        | j                  j                  j                  j                  j                         dS )N_fnr   r(   s    r*   	type_namezStubSignature.type_nameX  s1    dff'',,116678<<r,   c                 @    t        j                  | j                        S r4   )r&   stub_argumentsr!   r(   s    r*   r+   zStubSignature.arguments\  s    ##DFF++r,   c                 Z    | j                         }ddj                  d |D               dS )Nzvoid(*)(TensorIteratorBase&, rC   c              3   4   K   | ]  }|j                     y wr4   )r;   rG   s     r*   r>   z%StubSignature.type.<locals>.<genexpr>a  s     8RA8Rrg   rK   )r+   r@   )r)   cpp_argss     r*   r;   zStubSignature.type_  s.    >>#.tyy8R8R/R.SSTUUr,   c                 <    d| j                    d| j                   dS )NzDECLARE_DISPATCH(rC   rK   )r   r#   r(   s    r*   r   zStubSignature.dispatch_declc  s     "4>>"2"TYYKqAAr,   c                 "    d| j                    dS )NzDEFINE_DISPATCH(rK   rL   r(   s    r*   dispatch_defnzStubSignature.dispatch_defnf  s    !$))A..r,   c                 p    d| j                    ddj                  d | j                         D               dS )Nzvoid z(TensorIteratorBase& iter, rC   c              3   <   K   | ]  }|j                           y wr4   )r   rG   s     r*   r>   z,StubSignature.kernel_defn.<locals>.<genexpr>j  s     Mq[\affhMqrI   rK   )r   r@   r+   r(   s    r*   r   zStubSignature.kernel_defni  s:    t''((CDIIMq`d`n`n`pMqDqCrrsttr,   c                 B    d| j                    d| j                          S )Nzusing  = )r   r;   r(   s    r*   r   zStubSignature.type_defnl  s     's499;-88r,   rb   c           
          | j                    ddj                  d t        || j                               D               dS )Nz(device_type(), *this, rC   c              3   4   K   | ]  }|j                     y wr4   re   rG   s     r*   r>   z%StubSignature.call.<locals>.<genexpr>q  s     =oaff=org   rK   rh   ri   s     r*   rj   zStubSignature.callp  s=    ))3DII=oiX[]a]k]k]mNn=o4o3ppqrrr,   c           
          | j                    ddj                  d t        || j                               D               dS )Nz(*this, rC   c              3   4   K   | ]  }|j                     y wr4   re   rG   s     r*   r>   z,StubSignature.direct_call.<locals>.<genexpr>u  s     5gaff5grg   rK   )r   r@   r
   r+   ri   s     r*   r   zStubSignature.direct_callt  s?    ""#8DII5giPSUYUcUcUeFf5g,g+hhijjr,   N)rV   rW   rX   r   rY   propertyr[   r#   r   r   r   r   r+   r;   r   r   r   r   r   rj   r   r\   r,   r*   r   r   L  s    ?c ? ? AS A A =3 = =,4= ,Vc VBs B/s /uS u93 9s) sc skx0 kS kr,   r   c                 B   t        |       }t        | t        j                  | t        j
                              }d|j                          d|j                          d|j                          d|j                          d|j                  |j                                dS )Nr7   r   r   r   r   )r   r   r&   r   r   CPUr   r   r   r   rj   r+   )r!   r   r   s      r*   compute_ufunc_cpur   x  s    QH
!!U%6%6q+//%J
KC	  	  	   ==!" # r,   c                 z   t         j                  |v sJ | d|j                                 |j                         t         j                  t         j                  hk  sJ |t         j                     }d }t         j                  |v r|t         j                     }g }g |D ]  }t	        |j
                  t              r1|j
                  j                  t        t        j                        k7  rN|j                  d|j                   d|j                   d       j                  t        d|j                   t        |j                  j                  t!        t"                                  ||D ]  }t	        |j
                  t              r1|j
                  j                  t        t        j                        k7  rN|j                  d|j                   d|j                   d       j                  t        d	|j                   t        |j                  j                  t%        t!        t"                                        g }g }	| j&                  j(                  j*                  j,                  D ]  }
|
j                  j/                         s|
j                  t        t        j0                        k(  sJ |j                  t3        |
j                  t        |
j                  t!        t"                    |

             ||	j                  t3        |
j                  t        |
j                  t%        t!        t"                          |

              dt4        t2           dt6        t8        t        t2        f      ffd}dj;                  |      }|hd| ddj;                  d |D               d|j=                   ||             ddj;                  d |	D               d|j=                   ||	             dS d| ddj;                  d |D               d|j=                   ||             dS )NrC   zauto _s_r   z.to<scalar_t>();_s_zauto _v_z$ = at::vec::Vectorized<scalar_t>(_s_z);_v_)r#   nctypeargumentr1   r$   c                 P    g }|j                         |j                  |        |S r4   )extend)r1   rrb   s     r*   with_ctxz.compute_ufunc_cpu_dtype_body.<locals>.with_ctx  s#    (*		r,   r7   z
cpu_kernel_vec(iter,
  [=](c              3   <   K   | ]  }|j                           y wr4   rE   r<   r1   s     r*   r>   z/compute_ufunc_cpu_dtype_body.<locals>.<genexpr>       3a3rI   z) { return z; },
  [=](c              3   <   K   | ]  }|j                           y wr4   rE   r   s     r*   r>   z/compute_ufunc_cpu_dtype_body.<locals>.<genexpr>  s     0a0rI   z; }
);
z
cpu_kernel(iter,
  [=](c              3   <   K   | ]  }|j                           y wr4   rE   r   s     r*   r>   z/compute_ufunc_cpu_dtype_body.<locals>.<genexpr>  r   rI   )r   	CPUScalarr   	CPUVector
isinstancer   r   r;   r   r   Scalarr   r#   r   r   r   r   r   r   rq   rr   r+   rs   rn   Tensorr   r   r   r	   r@   rj   )r!   r   r   r   scalar_loopvec_loopr   r1   scalar_bindingsvec_bindingsrH   r   body_strrb   s                @r*   compute_ufunc_cpu_dtype_bodyr     s    ,Nr+:J:J:L9M.NN,("4"4h6H6H!IIIIh001KH[(x112 D
C Yajj(+

8MMD
 1
 hqvvhc!&&1ABC

4#affX
188==)HBU(VWXY  	A!**h/AJJOOxH 5 KK166("FqvvhbQ JJ!&&Nqxx}}oi>Q.RS	" OL\\((55 vv$$&vv&--0000VV!!&&)H*=>	
 %affoi>Q.RS(HW% $uT7]/C*D  yyH	
 yy3?334LAQAQRZ[jRkAl@m nyy0<001hmmHUaLb>c=d e	 		
 yy3?334LAQAQRZ[jRkAl@m n 	r,   c                 z   t        |       }| j                  j                  }i }t        j                  t        j
                  fD ]8  }g }||v r|j                  |       t        j                  |v r1|t        j                  u r|j                  t        j                         t        j                  |v r|j                  t        j                         |D ]  }||   j                  D ]  }|t        j                  u rt        t              }n5|t        j
                  u rt        t        t                    }n
t               |j                  |i       }	||	vsqt        | d||   j                    |      |	|<     ; g }
|j#                         D ]7  \  }}	|
j                  d| dt%        | ||	|j'                                d       9 dj)                  |
      }d|j+                          d|j                    d	| d
|j-                          d|j/                          d|j                    d|j0                   dS )Nrw   rx   r   r   r   r7   z
namespace {

r   r   z#
  );
}

} // anonymous namespace

r   z;
REGISTER_DISPATCH(r   z);
)r   ry   rz   r   r   r   r   r   r   r~   r   r   r   AssertionErrorr   r^   r#   r   r   r+   r@   r   r   r   r   )r!   r   r   
ufunc_sigsr   lksr   r   r_   r   r   r   s               r*   compute_ufunc_cpu_kernelr     sJ   QH EE""ECEJ  ("4"45 :JJqM%'A1C1C,CJJx**+u$JJx''( 	Br33 *** )( 3I(,,, /	(0C DI(**#-#8#8#C ,,*8'%)..)9 :i+$Q'	4 K#-#3#3#5 	
""' )!!U,<h>P>P>RST U	
	
 ii,O 
  ,,4MM? ;  
  	  ==/X%9%9$: ; r,   )5dataclassesr   typingr   r   r   r   r   r	   torchgen.api.ufuncapir&   torchgen.api.translater
   torchgen.api.typesr   r   r   r   r   r   r   r   r   r   torchgen.contextr   torchgen.modelr   r   r   r   r   r   r   torchgen.utilsr   r    r^   boolru   r[   r   r   r|   r{   r   r   r   r   r   r   r   r\   r,   r*   <module>r      sB   ! ? ? " " ,
 
 
 0 1   &> $P P PD $	] 	] 	]<1E $ P/P/
4
D+<!<==>CDP/f $   %--
 %,,% !,,, h 112, !	,
 	,^ &. &3 & &` $(k (k (kV - #  ZZZ h./Z !	Z
 	Zz ; 4 ; ; ;r,   