U
    9%eE                     @   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&e e'dddZ(e eee!ee"e%f f e)f d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.e e!ee"e%f ee e)dddZ/ee e)dddZ0ed
dG d d! d!Z1ee e)dd"d#Z2e e!ee"e&f ee e)dd$d%Z3ee e)d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< edddZ	e
e dddZedd	d
ZedddZedddZed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'   R/var/www/html/Darija-Ai-API/env/lib/python3.8/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   s     z0UfunctorSignature.decl_fields.<locals>.<genexpr>)joinr1   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   s   | ]}|  V  qd S r2   declr-   ar'   r'   r(   r9   S   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(   r9   V   s     (z) : z {})r:   r)   r0   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   s   | ]}|  V  qd S r2   r=   r?   r'   r'   r(   r9   Z   s     z/UfunctorSignature.decl_apply.<locals>.<genexpr>z operator()(z) const)r:   r)   applyr3   Zcpp_type)r&   rD   r'   r'   r(   
decl_applyY   s    zUfunctorSignature.decl_applyN)__name__
__module____qualname__r   __annotations__r   intstrr   r)   r   r   r1   r   r3   r;   rE   rG   r'   r'   r'   r(   r   ;   s   
r   c                   @   sT   e Zd ZU eed< eed< eed< ee dddZ	e
eeef  eddd	Zd
S )UfuncSignaturer   r!   	compute_tr"   c                 C   s   t j| j| jdS )N)rO   )r$   Zufunc_argumentsr   rO   r%   r'   r'   r(   r)   d   s    zUfuncSignature.argumentsctxr#   c              	   C   s,   | j  dddd t||  D  dS )NrC   r<   c                 s   s   | ]}|j V  qd S r2   exprr?   r'   r'   r(   r9   h   s     z&UfuncSignature.call.<locals>.<genexpr>rA   r!   r:   r	   r)   r&   rQ   r'   r'   r(   callg   s    zUfuncSignature.callN)rH   rI   rJ   r   rK   rM   r   r   r   r)   r   r   r   rV   r'   r'   r'   r(   rN   ^   s
   
rN   )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)r7   is_tensor_liker?   r'   r'   r(   r9   ~   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    
r^   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}n2tjg}tjtjfD ]}||ksTtd| dqT|D ]X}||krt| || || j	d}|| j
D ]}|||i |< qqvd }	t }
tjtjfD ]H}||krq|	d kr|| j	}	n|	|| j	kstd|
|| j
O }
q|	d k	s,t| d|	 }t| || |d}|
D ]}|||i |< qPt| d|	 ttd	}| | j }|d
|j	 d|  d|  d|  d|| d qv|d|fS )NrW   r   zcannot use z on non-binary function)r    r!   z0ScalarOnly and Generic must have same ufunc namer*   ufunc::r!   rO   z%
template <typename scalar_t>
struct z3 {
  using opmath_t = at::opmath_type<scalar_t>;
  z
  z
  __device__ z {
    return z	;
  }
};
r4   )outufunc_inner_loopr   CUDAFunctorOnSelfCUDAFunctorOnOtherCUDAFunctorr^   AssertionErrorr   r!   supported_dtypes
setdefaultr   
ScalarOnlyGenericrN   r
   r   r1   r)   rF   appendr;   rE   rG   rV   r:   )r   ufunctor_sigs	ufunctorsloopsZscalar_tensor_idx_lookupkeyskufunctor_sigdtypeZ
ufunc_namerg   lkr!   Z	ufunc_sigZ	apply_ctxr'   r'   r(   compute_ufunc_cuda_functors   s       
  	    rt   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)rH   rI   rJ   rL   rK   rM   r   r'   r'   r'   r(   ru      s   
ru   r&   )rv   rw   rx   rW   other)r   rr   inner_loops
parent_ctxr#   c           
      C   s   d}|d7 }t D ]}|j|k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) {}
rW   ziter.scalar_value<opmath_t>(rA   )rS   r7   r<   c                 s   s   | ]}|j V  qd S r2   rR   r?   r'   r'   r(   r9     s    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   s   | ]}|j V  qd S r2   rR   r?   r'   r'   r(   r9     s    z
else {
  gpu_kernel(iter, z<scalar_t>(z
));
}
    )!BinaryScalarSpecializationConfigsrx   rv   listrk   r   r   rw   r
   r   r:   r	   r)   r0   r!   r   re   )
r   rr   rz   r{   bodyconfigrq   rv   rQ   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::,
  [&]() {
    
  }
)
r4   z

;
;

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

 {
  ;
}
)rt   r   r$   kernel_namer   CUDAitemsrk   r   r)   r:   StubSignature	type_defndispatch_declkernel_defnr!   defndirect_call)	r   rl   rm   sigdtype_casesrr   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edddZeedddZeedddZ	e
e dd	d
ZedddZedddZedddZedddZedddZee edddZee edddZdS )r   r   r"   c                 C   s   t | jjjjj dS )NZ_stubrM   r   r[   r\   r!   r%   r'   r'   r(   r!   P  s    zStubSignature.namec                 C   s   t | jjjjj dS )NZ_kernelr   r%   r'   r'   r(   r   T  s    zStubSignature.kernel_namec                 C   s   t | jjjjj dS )N_fnr   r%   r'   r'   r(   	type_nameX  s    zStubSignature.type_namec                 C   s   t | jS r2   )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   s   | ]}|j V  qd S r2   )r7   r?   r'   r'   r(   r9   a  s     z%StubSignature.type.<locals>.<genexpr>rA   )r)   r:   )r&   Zcpp_argsr'   r'   r(   r7   _  s    zStubSignature.typec                 C   s   d| j  d| j dS )NzDECLARE_DISPATCH(r<   rA   )r   r!   r%   r'   r'   r(   r   c  s    zStubSignature.dispatch_declc                 C   s   d| j  dS )NzDEFINE_DISPATCH(rA   rB   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   s   | ]}|  V  qd S r2   )r   r?   r'   r'   r(   r9   j  s     z,StubSignature.kernel_defn.<locals>.<genexpr>rA   )r   r:   r)   r%   r'   r'   r(   r   i  s    zStubSignature.kernel_defnc                 C   s   d| j  d|   S )Nzusing  = )r   r7   r%   r'   r'   r(   r   l  s    zStubSignature.type_defnrP   c              	   C   s,   | j  dddd t||  D  dS )Nz(device_type(), *this, r<   c                 s   s   | ]}|j V  qd S r2   rR   r?   r'   r'   r(   r9   q  s     z%StubSignature.call.<locals>.<genexpr>rA   rT   rU   r'   r'   r(   rV   p  s    zStubSignature.callc              	   C   s,   | j  dddd t||  D  dS )Nz(*this, r<   c                 s   s   | ]}|j V  qd S r2   rR   r?   r'   r'   r(   r9   u  s     z,StubSignature.direct_call.<locals>.<genexpr>rA   )r   r:   r	   r)   rU   r'   r'   r(   r   t  s    zStubSignature.direct_callN)rH   rI   rJ   r   rK   propertyrM   r!   r   r   r   r   r)   r7   r   r   r   r   r   rV   r   r'   r'   r'   r(   r   L  s   
r   c                 C   sZ   t | }t| t| tj}d|  d|  d|  d|	  d|
|  dS )Nr4   r   r   r   r   )r   r   r$   r   r   ZCPUr   r   r   r   rV   r)   )r   r   r   r'   r'   r(   compute_ufunc_cpux  s    r   c                    s  t j|ks t| d|  | t jt jhks8t|t j }d }t j|krZ|t j }g }g  |D ]f}t|jtr|jjt	t
jkrqf|d|j d|j d  td|j t|jjtt qf|d k	rL|D ]n}t|jtr|jjt	t
jkrq|d|j d|j d  td	|j t|jjttt qg }g }	| jjjjD ]~}
|
j stq`|
jt	t
jkst|t|
jt|
jtt|
d
 |d k	r`|	t|
jt|
jttt|
d
 q`tt ttttf  d fdd}d|}|d k	rpd| dddd |D  d||| dddd |	D  d|||	 dS d| dddd |D  d||| dS 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argument)r.   r#   c                    s   g }|   | |  |S r2   )extend)r.   rrQ   r'   r(   with_ctx  s    

z.compute_ufunc_cpu_dtype_body.<locals>.with_ctxr4   z
cpu_kernel_vec(iter,
  [=](c                 s   s   | ]}|  V  qd S r2   r=   r,   r'   r'   r(   r9     s     z/compute_ufunc_cpu_dtype_body.<locals>.<genexpr>z) { return z; },
  [=](c                 s   s   | ]}|  V  qd S r2   r=   r,   r'   r'   r(   r9     s     z; }
);
z
cpu_kernel(iter,
  [=](c                 s   s   | ]}|  V  qd S r2   r=   r,   r'   r'   r(   r9     s     ) r   	CPUScalarrf   ro   	CPUVector
isinstancer   r   r7   r   r   ZScalarrk   r!   r   r   r   r
   r   r   r[   r\   r)   r]   rX   ZTensorr   r   r   r   r:   rV   )r   rr   rz   r{   Zscalar_loopZvec_loopr~   r.   Zscalar_bindingsZvec_bindingsr@   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 }||kr:|| tj|krZ|tjkrZ|tj tj|krp|tj |D ]x}|| j	D ]h}|tjkrt
t}n|tjkrtt
t}nt ||i }	||	krt| d|| j |d|	|< qqtq g }
| D ].\}}	|
d| dt| ||	|  d qd|
}d|  d|j d	| d
|  d|  d|j d|j dS )Nr_   r`   r   r   r   r4   z
namespace {

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

} // anonymous namespace

r   z;
REGISTER_DISPATCH(r   z);
)r   ra   rb   r   r   r   rk   ri   rj   rg   r
   r   r   rf   rh   rN   r!   r   r   r)   r:   r   r   r   r   )r   r   rn   Z
ufunc_sigsrp   Zlksrs   rr   rO   r   r   r   r'   r'   r(   compute_ufunc_cpu_kernel  sd    




  

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   rN   boolr^   rM   rt   ru   rd   rc   r|   r   r   r   r   r   r   r'   r'   r'   r(   <module>   sb    ,$	"S
/0+
]