a
    hE                     @  s  d dl mZ d dlmZ d dl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mZmZmZmZmZmZ d d	lm Z  erd d
l!m"Z" d dlm#Z# eddG dd dZ$eddG dd dZ%dddddZ&dd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+dd d!d"d#d$d%d&Z,edd#dd'd(Z-eddG d)d* d*Z.edd#dd+d,Z/dd d-d"d#d$d.d/Z0edd#dd0d1Z1dS )2    )annotations)	dataclass)TYPE_CHECKINGN)	translate)		BaseCTypeBindingCTypeExpr
NamedCTypeopmath_tscalar_tStructuredImplSignatureVectorizedCType)with_native_function)ArgumentBaseTyBaseTypeDispatchKeyNativeFunctionsGroup
ScalarTypeUfuncKey)
OrderedSet)Sequence)UfunctorBindingsT)frozenc                   @  sz   e Zd ZU ded< ded< ded< ddd	d
ZddddZddddZddddZddddZddddZ	dS )UfunctorSignaturer   gz
int | Nonescalar_tensor_idxstrnamer   returnc                 C  s   t j| j| jtdS )N)r   r   )ufuncZufunctor_argumentsr   r   r   self r%   A/var/www/auris/lib/python3.9/site-packages/torchgen/dest/ufunc.py	argumentsI   s    
zUfunctorSignature.argumentslist[Binding]c                 C  s   dd |   jD S )Nc                 S  s   g | ]}| |j d qS )_)renamer   .0br%   r%   r&   
<listcomp>P       z,UfunctorSignature.fields.<locals>.<listcomp>)r'   ctorr#   r%   r%   r&   fieldsN   s    zUfunctorSignature.fieldsr   c                 C  s   t tS N)r   r   r#   r%   r%   r&   returns_typeR   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>X   r/   z0UfunctorSignature.decl_fields.<locals>.<genexpr>)joinr1   r#   r%   r%   r&   decl_fieldsW   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   [   r/   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   ^   r/   (z) : z {})r:   r'   r0   r   )r$   args_strZinit_strr%   r%   r&   inline_defn_ctorZ   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   b   r/   z/UfunctorSignature.decl_apply.<locals>.<genexpr>z operator()(z) const)r:   r'   applyr3   Zcpp_type)r$   rD   r%   r%   r&   
decl_applya   s    zUfunctorSignature.decl_applyN)
__name__
__module____qualname____annotations__r'   r1   r3   r;   rE   rG   r%   r%   r%   r&   r   C   s   
r   c                   @  sD   e Zd ZU ded< ded< ded< ddd	d
ZdddddZdS )UfuncSignaturer   r   r   r   r   	compute_tr(   r    c                 C  s   t j| j| jdS )N)rM   )r"   Zufunc_argumentsr   rM   r#   r%   r%   r&   r'   l   s    zUfuncSignature.argumentszSequence[Binding | Expr]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   p   r/   z&UfuncSignature.call.<locals>.<genexpr>rA   r   r:   r   r'   r$   rO   r%   r%   r&   callo   s    zUfuncSignature.callN)rH   rI   rJ   rK   r'   rT   r%   r%   r%   r&   rL   f   s
   
rL   r   bool)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]   z?tuple[dict[ScalarType, dict[UfuncKey, UfunctorSignature]], str]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 ]}||vsTJ d| dqT|D ]X}||v rt| || || jd}|| j	D ]}||
|i |< qqvd }	t }
tjtjfD ]H}||vrq|	d u r|| j}	n|	|| jksJ d|
|| j	O }
q|	d us,J | d|	 }t| || |d}|
D ]}||
|i |< qPt| d|	 ttd	}| | j }|d
|j d|  d|  d|  d|| d qv|d|fS )NrV   r   zcannot use z on non-binary function)r   r   z0ScalarOnly and Generic must have same ufunc namer)   ufunc::r   rM   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]   r   r   supported_dtypes
setdefaultr   
ScalarOnlyGenericrL   r   r   r1   r'   rF   appendr;   rE   rG   rT   r:   )r   ufunctor_sigs	ufunctorsloopsZscalar_tensor_idx_lookupkeyskufunctor_sigdtypeZ
ufunc_namere   lkr   Z	ufunc_sigZ	apply_ctxr%   r%   r&   compute_ufunc_cuda_functors   sv    
	
rr   c                   @  s&   e Zd ZU ded< ded< ded< dS ) BinaryScalarSpecializationConfigint
scalar_idxr   ctor_tensorr   	ufunc_keyN)rH   rI   rJ   rK   r%   r%   r%   r&   rs      s   
rs   r$   )ru   rv   rw   rV   otherr   z!dict[UfuncKey, UfunctorSignature]Sequence[Binding]r   )r   rp   inner_loops
parent_ctxr!   c           
      C  s   d}|d7 }t D ]}|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) {}
rV   ziter.scalar_value<opmath_t>(rA   )rQ   r7   r<   c                 s  s   | ]}|j V  qd S r2   rP   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   rP   r?   r%   r%   r&   r9     s   z
else {
  gpu_kernel(iter, z<scalar_t>(z
));
}
    )!BinaryScalarSpecializationConfigsrw   ru   listri   r	   r
   rv   r   r   r:   r   r'   r0   r   r   rd   )
r   rp   rz   r{   bodyconfigro   ru   rO   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)

 {
  ;
}
)rr   r   r"   kernel_namer   ZCUDAitemsri   r   r'   r:   StubSignature	type_defndispatch_declkernel_defnr   defndirect_call)	r   rj   rk   sigdtype_casesrp   inner_ufunc_sigsdtype_cases_strstub_sigr%   r%   r&   compute_ufunc_cuda#  sD    

r   c                   @  s   e Zd ZU ded< eddddZeddddZeddd	d
ZddddZddddZ	ddddZ
ddddZddddZddddZdddddZdddddZdS )r   r   r   r   r    c                 C  s   t | jjjjj dS )NZ_stubr   r   rZ   r[   r   r#   r%   r%   r&   r   X  s    zStubSignature.namec                 C  s   t | jjjjj dS )NZ_kernelr   r#   r%   r%   r&   r   \  s    zStubSignature.kernel_namec                 C  s   t | jjjjj dS )N_fnr   r#   r%   r%   r&   	type_name`  s    zStubSignature.type_namer(   c                 C  s   t | jS r2   )r"   Zstub_argumentsr   r#   r%   r%   r&   r'   d  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   i  r/   z%StubSignature.type.<locals>.<genexpr>rA   )r'   r:   )r$   Zcpp_argsr%   r%   r&   r7   g  s    zStubSignature.typec                 C  s   d| j  d| j dS )NzDECLARE_DISPATCH(r<   rA   )r   r   r#   r%   r%   r&   r   k  s    zStubSignature.dispatch_declc                 C  s   d| j  dS )NzDEFINE_DISPATCH(rA   rB   r#   r%   r%   r&   dispatch_defnn  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   r  r/   z,StubSignature.kernel_defn.<locals>.<genexpr>rA   )r   r:   r'   r#   r%   r%   r&   r   q  s    zStubSignature.kernel_defnc                 C  s   d| j  d|   S )Nzusing  = )r   r7   r#   r%   r%   r&   r   t  s    zStubSignature.type_defnry   rN   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   rP   r?   r%   r%   r&   r9   y  r/   z%StubSignature.call.<locals>.<genexpr>rA   rR   rS   r%   r%   r&   rT   x  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   rP   r?   r%   r%   r&   r9   }  r/   z,StubSignature.direct_call.<locals>.<genexpr>rA   )r   r:   r   r'   rS   r%   r%   r&   r   |  s    zStubSignature.direct_callN)rH   rI   rJ   rK   propertyr   r   r   r'   r7   r   r   r   r   rT   r   r%   r%   r%   r&   r   T  s   
r   c                 C  sZ   t | }t| t| tj}d|  d|  d|  d|	  d|
|  dS )Nr4   r   z;

r   r   )r   r   r"   r   r   ZCPUr   r   r   r   rT   r'   )r   r   r   r%   r%   r&   compute_ufunc_cpu  s    r   zdict[UfuncKey, UfuncSignature]c                   s  t j|v s J | d|  | t jt jhks8J |t j }d }t j|v 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 u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sJ |t|
jt|
jtt|
d
 |d ur`|	t|
jt|
jttt|
d
 q`ddd fdd}d|}|d ur`d| 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argumentry   zlist[Expr | Binding])r-   r!   c                   s   g }|   | |  |S r2   )extend)r-   rrO   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     r/   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     r/   z; }
);
z
cpu_kernel(iter,
  [=](c                 s  s   | ]}|  V  qd S r2   r=   r+   r%   r%   r&   r9     r/   )r   	CPUScalarrm   	CPUVector
isinstancer   r   r7   r   r   ZScalarri   r   r	   r
   r   r   r   r   rZ   r[   r'   r\   rW   ZTensorr   r:   rT   )r   rp   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  sv  t | }| jj}i }tjtjfD ]}g }||v r:|| tj|v rZ|tju rZ|tj tj|v rp|tj |D ]v}|| j	D ]f}|tju rt
t}n|tju rtt
t}nt||i }	||	v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   r`   ra   r   r   r   ri   rg   rh   re   r   r   r   AssertionErrorrf   rL   r   r   r   r'   r:   r   r   r   r   )r   r   rl   Z
ufunc_sigsrn   Zlksrq   rp   rM   r   r   r   r%   r%   r&   compute_ufunc_cpu_kernel  s`    






r   )2
__future__r   Zdataclassesr   typingr   Ztorchgen.api.ufuncapir"   Ztorchgen.api.translater   Ztorchgen.api.typesr   r   r   r	   r
   r   r   r   r   Ztorchgen.contextr   Ztorchgen.modelr   r   r   r   r   r   r   Ztorchgen.utilsr   collections.abcr   r   r   rL   r]   rr   rs   rc   rb   r|   r   r   r   r   r   r   r%   r%   r%   r&   <module>   sP   ,$	 "S/0+]