o
    HۂiJ                     @   s  d 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mZmZmZmZmZmZmZmZmZmZmZmZmZmZ ddlmZmZmZmZmZm Z m!Z!m"Z"m#Z#m$Z$m%Z%m&Z&m'Z'm(Z(m)Z)m*Z*m+Z+m,Z,m-Z-m.Z.m/Z/m0Z0m1Z1m2Z2m3Z3m4Z4m5Z5m6Z6m7Z7m8Z8m9Z9m:Z:m;Z;m<Z<m=Z=m>Z>m?Z?m@Z@mAZAmBZBmCZCmDZDmEZEmFZFmGZGmHZHmIZImJZJmKZKmLZLmMZMmNZNmOZOmPZPmQZQmRZRmSZSmTZTmUZUmVZVmWZWmXZXmYZYmZZZm[Z[m\Z\m]Z]m^Z^m_Z_m`Z`maZambZbmcZcmdZdmeZemfZfmgZgmhZhmiZimjZjmkZkmlZlmmZmmnZnmoZompZpmqZqmrZrmsZsmtZtmuZu ddlmvZvmwZwmxZxmyZymzZzm{Z{m|Z|m}Z}m~Z~mZmZmZmZmZmZmZmZ ddlmZmZmZmZmZmZmZmZmZmZ ddlmZ g d	Zd
d ZdS )zisort:skip_file   )math)extra)argmaxargminbitonic_mergecdivcumprodcumsumflip
interleavemaxminravel	reduce_orsigmoidsoftmaxsortsum	swizzle2dtopkxor_sumzeros
zeros_like)[PropagateNanTRITON_MAX_TENSOR_NUMELload_tensor_descriptorstore_tensor_descriptormake_tensor_descriptortensor_descriptortensor_descriptor_typeaddadvancearangeassociative_scanassume
atomic_add
atomic_and
atomic_cas
atomic_max
atomic_min	atomic_oratomic_xchg
atomic_xorbfloat16
block_type	broadcastbroadcast_tocatcastclamp	conditionconst	constexprconstexpr_typedebug_barrierdevice_assertdevice_printdot
dot_scaleddtypeexpand_dimsfloat16float32float64float8e4b15
float8e4nv
float8e4b8float8e5float8e5b16fullgather	histograminline_asm_elementwiseint1int16int32int64int8joinloadmake_block_ptrmap_elementwisemax_constancymax_contiguousmaximumminimummulmultiple_ofnum_programspermutepi32_tpointer_type
program_idrangereducereshapeslicesplitstatic_assertstatic_printstatic_rangestoresubtensortranstuple
tuple_typeuint16uint32uint64uint8viewvoidwhere)umulhiexpexp2fmaloglog2cosrsqrtsinsqrtsqrt_rnabsfdivdiv_rnerffloorceil)
pair_uniform_to_normalphiloxphilox_implrandrand4xrandint	randint4xrandnrandn4xuint_to_uniform_float)target_info)r   r   r   r   r   r   r   r    r!   r"   r   r   r#   r$   r%   r&   r'   r(   r)   r*   r+   r,   r-   r   r.   r/   r0   r1   r2   r   r   r3   r4   r5   r6   r7   rz   r   r	   r8   r9   r:   r   r;   r<   r=   r   ru   rv   r>   r   r   r
   r?   r@   rA   rB   rC   rD   rE   rF   r   rw   rG   rH   rI   rJ   r   rK   rL   rM   rN   rO   rP   rQ   rx   ry   rR   rS   r   r   rT   rU   rV   r   rW   rX   rY   rZ   r   r[   r   r   r\   r]   r^   r   r   r   r   r   r   r_   r   r`   r   ra   r{   rb   r   r|   r   r   rc   r}   r~   rd   re   rf   rg   rh   r   r   r   ri   r   rj   rk   rm   rn   ro   rp   r   rt   rq   rr   rs   r   r   r   c                    s:  ddl m} t| |r t| jdd }t fdd| D |S | d dkrG| dd  } d}| d d	kr<| dd  } d
}t|  }t||dS | 	dr| 
dd d}|j
ddd\}}|j
ddd\}	}dd |	d
dD }	|d}
t|
}t|d }t|	}ttg| }ttg| }t||	}|rddlm}m}m} ddlm} ddlm} t|
t|||d}
t|
|r|||||
S |||||
S t|||S | 	drt S i dtdtdtdtdtdt dt!dt"dt#d t$d!t%d"t&d#td$td%t$d&t'd't(t)t*t$d(}||  S ))N    )rk   _fieldsc                    s   g | ]}t | qS  )	str_to_ty).0xcr   [/home/ubuntu/maya3_transcribe/venv/lib/python3.10/site-packages/triton/language/__init__.py
<listcomp>  s    zstr_to_ty.<locals>.<listcomp>*r   FkT)
element_tyr5   
tensordesc<>[)maxsplit]c                 S   s   g | ]}t | qS r   )intstrip)r   sr   r   r   r   .  s    ,)NVMMASharedLayoutPaddedSharedLayoutSwizzledSharedLayout)r   r6   fp8e4nvfp8e4b8fp8e5fp8e5b16fp8e4b15fp16bf16fp32fp64i1i8i16i32i64u1u8u16)u32u64B)+builtinsrk   
isinstancetype__dict__getrl   r   r]   
startswithrc   rstriplstriplenrM   rN   r.   +triton.experimental.gluon.language._layoutsr   r   r   4triton.experimental.gluon.language.nvidia.hopper.tmar   2triton.experimental.gluon.language.amd.gfx1250.tdmevaldictr7   rC   rD   rE   rF   rB   r?   r-   r@   rA   rK   rO   rL   rp   rm   rn   ro   )namer   rk   fieldsr5   tyinnerr=   restblock_shapelayoutis_gluonndim
shape_typestride_typeblockr   r   r   nvidia_tensor_descriptor_typeamd_tensor_descriptor_typetysr   r   r   r     s   







	
r   N)__doc__ r   r   standardr   r   r   r   r   r	   r
   r   r   r   r   r   r   r   r   r   r   r   r   r   r   corer   r   r   r   r   r   r   r    r!   r"   r#   r$   r%   r&   r'   r(   r)   r*   r+   r,   r-   r.   r/   r0   r1   r2   r3   r4   r5   r6   r7   r8   r9   r:   r;   r<   r=   r>   r?   r@   rA   rB   rC   rD   rE   rF   rG   rH   rI   rJ   rK   rL   rM   rN   rO   rP   rQ   rR   rS   rT   rU   rV   rW   rX   rY   rZ   r[   r\   r]   r^   r_   r`   ra   rb   rc   rd   re   rf   rg   rh   ri   rj   rk   rl   rm   rn   ro   rp   rq   rr   rs   rt   ru   rv   rw   rx   ry   rz   r{   r|   r}   r~   r   r   r   r   r   r   randomr   r   r   r   r   r   r   r   r   r   r   __all__r   r   r   r   r   <module>   s    \v L]0 