o
    iE                  
   @   s  d dl mZ d dlmZmZmZmZmZmZ d dl	m
  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
dG dd dZ%ed
dG dd dZ&de de'fddZ(de deee!ee"e%f f e)f fddZ*ed
dG dd dZ+e+d de"j,de+dde"j-dgZ.de de!dee"e%f dee de)f
dd Z/ede de)fd!d"Z0ed
dG d#d$ d$Z1ede de)fd%d&Z2de de!dee"e&f dee de)f
d'd(Z3ede de)fd)d*Z4dS )+    )	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                   @   s   e Zd ZU eed< ee ed< eed< defddZ	de
e fddZdefd	d
ZdefddZdefddZdefddZdS )UfunctorSignaturegscalar_tensor_idxnamereturnc                 C   s   t j| j| jtdS )N)r    r   )ufuncZufunctor_argumentsr   r    r   self r&   b/var/www/html/eduruby.in/lip-sync/lip-sync-env/lib/python3.10/site-packages/torchgen/dest/ufunc.py	argumentsA   s   
zUfunctorSignature.argumentsc                 C   s   dd |   jD S )Nc                 S   s   g | ]}| |j d qS )_)renamer!   .0br&   r&   r'   
<listcomp>H   s    z,UfunctorSignature.fields.<locals>.<listcomp>)r(   ctorr$   r&   r&   r'   fieldsF   s   zUfunctorSignature.fieldsc                 C   s   t tS N)r
   r   r$   r&   r&   r'   returns_typeJ   s   zUfunctorSignature.returns_typec                 C   s   d dd |  D S )N
c                 s   s$    | ]}|j  d |j dV  qdS ) ;N)typer!   )r,   fr&   r&   r'   	<genexpr>P      " z0UfunctorSignature.decl_fields.<locals>.<genexpr>)joinr0   r$   r&   r&   r'   decl_fieldsO   s   zUfunctorSignature.decl_fieldsc                 C   sL   d dd |  jD }d dd |  jD }| j d| d| dS )N, c                 s       | ]}|  V  qd S r1   declr,   ar&   r&   r'   r8   S       z5UfunctorSignature.inline_defn_ctor.<locals>.<genexpr>c                 s   s$    | ]}|j  d |j  dV  qdS )z_()Nr!   r@   r&   r&   r'   r8   V   r9   (z) : z {})r:   r(   r/   r!   )r%   args_strZinit_strr&   r&   r'   inline_defn_ctorR   s   z"UfunctorSignature.inline_defn_ctorc                 C   s2   d dd |  jD }|    d| dS )Nr<   c                 s   r=   r1   r>   r@   r&   r&   r'   r8   Z   rB   z/UfunctorSignature.decl_apply.<locals>.<genexpr>z operator()(z) const)r:   r(   applyr2   Zcpp_type)r%   rF   r&   r&   r'   
decl_applyY   s   zUfunctorSignature.decl_applyN)__name__
__module____qualname__r   __annotations__r   intstrr   r(   r   r   r0   r   r2   r;   rG   rI   r&   r&   r&   r'   r   ;   s   
 r   c                   @   sV   e Zd ZU eed< eed< eed< dee fddZ	de
eeef  defdd	Zd
S )UfuncSignaturer   r!   	compute_tr"   c                 C   s   t j| j| jdS )N)rQ   )r#   Zufunc_argumentsr   rQ   r$   r&   r&   r'   r(   d   s   zUfuncSignature.argumentsctxc              	   C   ,   | j  dddd t||  D  dS )NrE   r<   c                 s       | ]}|j V  qd S r1   exprr@   r&   r&   r'   r8   h       z&UfuncSignature.call.<locals>.<genexpr>rC   r!   r:   r	   r(   r%   rR   r&   r&   r'   callg      ,zUfuncSignature.callN)rJ   rK   rL   r   rM   rO   r   r   r   r(   r   r   r   rZ   r&   r&   r&   r'   rP   ^   s   
 "rP   r   r"   c                 C   s"   t dd | jjjjD }|dkS )Nc                 s   s    | ]
}|j  rd V  qdS )   N)r6   is_tensor_liker@   r&   r&   r'   r8   ~   s    

z<eligible_for_binary_scalar_specialization.<locals>.<genexpr>   )sum
functionalfuncr(   flat_non_out)r   Znum_tensorsr&   r&   r'   )eligible_for_binary_scalar_specialization}   s   
rc   c                 C   s  i }g }| j j}tjdtjdtjd i}t| r tjtjtjg}ntjg}tjtjfD ]}||vs8J d| dq*|D ]}||v r^t| || || jd}|| j	D ]
}||
|i |< qRq;d }	t }
tjtjfD ]#}||vrpqi|	d u rz|| j}	n|	|| jksJ d|
|| j	O }
qi|	d usJ | d|	 }t| || |d}|
D ]
}||
|i |< qt| d|	 ttd	}| | j }|d
|j d|  d|  d|  d|| d q;|d|fS )Nr\   r   zcannot use z on non-binary function)r    r!   z0ScalarOnly and Generic must have same ufunc namer)   ufunc::r!   rQ   z%
template <typename scalar_t>
struct z3 {
  using opmath_t = at::opmath_type<scalar_t>;
  z
  z
  __device__ z {
    return z	;
  }
};
r3   )outufunc_inner_loopr   CUDAFunctorOnSelfCUDAFunctorOnOtherCUDAFunctorrc   r   r!   supported_dtypes
setdefaultr   
ScalarOnlyGenericrP   r
   r   r0   r(   rH   appendr;   rG   rI   rZ   r:   )r   ufunctor_sigs	ufunctorsloopsZscalar_tensor_idx_lookupkeyskufunctor_sigdtypeZ
ufunc_namerk   lkr!   Z	ufunc_sigZ	apply_ctxr&   r&   r'   compute_ufunc_cuda_functors   sv   	
rx   c                   @   s&   e Zd ZU eed< eed< eed< dS ) BinaryScalarSpecializationConfig
scalar_idxctor_tensor	ufunc_keyN)rJ   rK   rL   rN   rM   rO   r   r&   r&   r&   r'   ry      s   
 ry   r%   )rz   r{   r|   r\   otherrv   inner_loops
parent_ctxc           
      C   s   d}|d7 }t D ]K}|j|vrq||j }|jd }t|}|td| dt|jtt	d d
dd	 t|| jD }	|d
| d|j d|	 d| d	7 }q|tj }d
dd	 t|| jD }	|d|j d|	 d7 }|S )Nz+using opmath_t = at::opmath_type<scalar_t>;zif (false) {}
r\   ziter.scalar_value<opmath_t>(rC   )rV   r6   r<   c                 s   rT   r1   rU   r@   r&   r&   r'   r8         
z0compute_ufunc_cuda_dtype_body.<locals>.<genexpr>zelse if (iter.is_cpu_scalar(z)) {
  z<scalar_t> ufunctor(z);
  iter.remove_operand(z");
  gpu_kernel(iter, ufunctor);
}c                 s   rT   r1   rU   r@   r&   r&   r'   r8     r   z
else {
  gpu_kernel(iter, z<scalar_t>(z
));
}
    )!BinaryScalarSpecializationConfigsr|   rz   listro   r   r   r{   r
   r   r:   r	   r(   r/   r!   r   rj   )
r   rv   r~   r   bodyconfigru   rz   rR   Zufunctor_ctor_exprs_strr&   r&   r'   compute_ufunc_cuda_dtype_body   sH   







r   c           	      C   s   t | \}}t| t| tj}g }| D ]\}}|d| dt| |||	  d qd
|}t| }d| d|  d|  d|  d|j d	| d
|j d|j d|  d||	  dS )N"
AT_DISPATCH_CASE(at::ScalarType::,
  [&]() {
    
  }
)
r3   z

;
;

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

 {
  ;
}
)rx   r   r#   kernel_namer   CUDAitemsro   r   r(   r:   StubSignature	type_defndispatch_declkernel_defnr!   defndirect_call)	r   rp   rq   sigdtype_casesrv   inner_ufunc_sigsdtype_cases_strstub_sigr&   r&   r'   compute_ufunc_cuda  sD   

r   c                   @   s   e Zd ZU eed< edefddZedefddZedefddZ	de
e fd	d
ZdefddZdefddZdefddZdefddZdefddZdee defddZdee defddZdS )r   r   r"   c                 C      t | jjjjj dS )NZ_stubrO   r   r`   ra   r!   r$   r&   r&   r'   r!   P     zStubSignature.namec                 C   r   )NZ_kernelr   r$   r&   r&   r'   r   T  r   zStubSignature.kernel_namec                 C   r   )N_fnr   r$   r&   r&   r'   	type_nameX  r   zStubSignature.type_namec                 C   s   t | jS r1   )r#   Zstub_argumentsr   r$   r&   r&   r'   r(   \  s   zStubSignature.argumentsc                 C   s$   |   }dddd |D  dS )Nzvoid(*)(TensorIteratorBase&, r<   c                 s   rT   r1   )r6   r@   r&   r&   r'   r8   a  rW   z%StubSignature.type.<locals>.<genexpr>rC   )r(   r:   )r%   Zcpp_argsr&   r&   r'   r6   _  s   zStubSignature.typec                 C   s   d| j  d| j dS )NzDECLARE_DISPATCH(r<   rC   )r   r!   r$   r&   r&   r'   r   c     zStubSignature.dispatch_declc                 C   s   d| j  dS )NzDEFINE_DISPATCH(rC   rD   r$   r&   r&   r'   dispatch_defnf  s   zStubSignature.dispatch_defnc                 C   s(   d| j  dddd |  D  dS )Nzvoid z(TensorIteratorBase& iter, r<   c                 s   r=   r1   )r   r@   r&   r&   r'   r8   j  rB   z,StubSignature.kernel_defn.<locals>.<genexpr>rC   )r   r:   r(   r$   r&   r&   r'   r   i  s   (zStubSignature.kernel_defnc                 C   s   d| j  d|   S )Nzusing  = )r   r6   r$   r&   r&   r'   r   l  r   zStubSignature.type_defnrR   c              	   C   rS   )Nz(device_type(), *this, r<   c                 s   rT   r1   rU   r@   r&   r&   r'   r8   q  rW   z%StubSignature.call.<locals>.<genexpr>rC   rX   rY   r&   r&   r'   rZ   p  r[   zStubSignature.callc              	   C   rS   )Nz(*this, r<   c                 s   rT   r1   rU   r@   r&   r&   r'   r8   u  rW   z,StubSignature.direct_call.<locals>.<genexpr>rC   )r   r:   r	   r(   rY   r&   r&   r'   r   t  r[   zStubSignature.direct_callN)rJ   rK   rL   r   rM   propertyrO   r!   r   r   r   r   r(   r6   r   r   r   r   r   rZ   r   r&   r&   r&   r'   r   L  s    
 r   c                 C   sZ   t | }t| t| tj}d|  d|  d|  d|	  d|
|  dS )Nr3   r   r   r   r   )r   r   r#   r   r   ZCPUr   r   r   r   rZ   r(   )r   r   r   r&   r&   r'   compute_ufunc_cpux  s   r   c                    s  t j|v sJ | d|  | t jt jhksJ |t j }d }t j|v r-|t j }g }g  |D ]3}t|jtrE|jjtt	j
krEq3|d|j d|j d  td|j t|jjtt q3|d ur|D ]5}t|jtr|jjtt	j
krqm|d|j d|j d  td	|j t|jjttt qmg }g }	| jjjjD ]:}
|
j sq|
jtt	jksJ |t|
jt|
jtt|
d
 |d ur|	t|
jt|
jttt|
d
 qdtt dttttf  f fdd}d|}|d ur1d| dddd |D  d||| dddd |	D  d|||	 dS d| dddd |D  d||| dS )Nr<   zauto _s_r   z.to<scalar_t>();Z_s_zauto _v_z$ = at::vec::Vectorized<scalar_t>(_s_z);Z_v_)r!   nctypeargumentr-   r"   c                    s   g }|   | |  |S r1   )extend)r-   rrR   r&   r'   with_ctx  s   

z.compute_ufunc_cpu_dtype_body.<locals>.with_ctxr3   z
cpu_kernel_vec(iter,
  [=](c                 s   r=   r1   r>   r+   r&   r&   r'   r8     rB   z/compute_ufunc_cpu_dtype_body.<locals>.<genexpr>z) { return z; },
  [=](c                 s   r=   r1   r>   r+   r&   r&   r'   r8     rB   z; }
);
z
cpu_kernel(iter,
  [=](c                 s   r=   r1   r>   r+   r&   r&   r'   r8     rB   )r   	CPUScalarrs   	CPUVector
isinstancer   r   r6   r   r   ZScalarro   r!   r   r   r   r
   r   r   r`   ra   r(   rb   r]   ZTensorr   r   r   r   r:   rZ   )r   rv   r~   r   Zscalar_loopZvec_loopr   r-   Zscalar_bindingsZvec_bindingsrA   r   Zbody_strr&   r   r'   compute_ufunc_cpu_dtype_body  s    


	(
	
&

r   c                 C   sx  t | }| jj}i }tjtjfD ]g}g }||v r|| tj|v r-|tju r-|tj tj|v r8|tj |D ]<}|| j	D ]4}|tju rMt
t}n|tju rYtt
t}nt ||i }	||	vrut| d|| j |d|	|< qAq:qg }
| D ]\}}	|
d| dt| ||	|  d q~d|
}d|  d|j d	| d
|  d|  d|j d|j dS )Nrd   re   r   r   r   r3   z
namespace {

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

} // anonymous namespace

r   z;
REGISTER_DISPATCH(r   z);
)r   rf   rg   r   r   r   ro   rm   rn   rk   r
   r   r   AssertionErrorrl   rP   r!   r   r   r(   r:   r   r   r   r   )r   r   rr   Z
ufunc_sigsrt   Zlksrw   rv   rQ   r   r   r   r&   r&   r'   compute_ufunc_cpu_kernel  sf   







r   )5dataclassesr   typingr   r   r   r   r   r   Ztorchgen.api.ufuncapir#   Ztorchgen.api.translater	   Ztorchgen.api.typesr
   r   r   r   r   r   r   r   r   r   Ztorchgen.contextr   Ztorchgen.modelr   r   r   r   r   r   r   Ztorchgen.utilsr   r   rP   boolrc   rO   rx   ry   ri   rh   r   r   r   r   r   r   r   r&   r&   r&   r'   <module>   s|     ,$	"
S

/0+

]