o
    ٗi                     @   sN  d dl Z d dlZd dlZd dlZd dlZd dlmZ d dl	Z	d dl
mZ d dlmZmZ d dlmZmZ g dZg dZee Zg dZee Zedg Zd	d
gZdge dg e dg Zdd Zdd Zdd Zdd Zdd Zd&dee fddZd'dej deeej!f fddZ"de#fddZ$d d! Z%d"d# Z&e	j'j(e&  d$d%Z)dS )(    N)RandomState)OptionalUnion)TensorWrapperreinterpret)int8int16int32int64)uint8uint16uint32uint64)float16float32float64bfloat16float8_e4m3fnfloat8_e5m2boolr   c                   C   s   t jdddkS )NTRITON_INTERPRET01)osenvironget r   r   V/home/ubuntu/transcripts/venv/lib/python3.10/site-packages/triton/_internal_testing.pyis_interpreter   s   r   c                   C   s   t  rd S tjjj S N)r   tritonruntimedriveractiveget_current_targetr   r   r   r   r$      s   r$   c                  C      t  } | d u r	dS | jdkS )NFcudar$   backendtargetr   r   r   is_cuda!      r+   c                  C   r%   )NFhipr'   r)   r   r   r   is_hip&   r,   r.   c                  C   s   t  } | d u r	dS t| jS )N )r$   strarchr)   r   r   r   get_arch+   r,   r2   rsc                 C   s8  t | tr| f} |du rtdd}|tt v rOttt|}|du r&|jnt	||j}|du r3|j	nt||j	}tt|}|j
||| |d}d||dk< |S |rad|v ra|j
dd	| tjd}|S |tv ro|dd| |S |d
kr|dd| ddtd@ dS |dv r|dd| dkS td| )zp
    Override `rs` if you're calling this function twice and don't want the same
    result for both calls.
    N   )seed)dtype   r   float8   (   r   r   r   l      )r   int1bool_g        zUnknown dtype )
isinstanceintr   
int_dtypesuint_dtypesnpiinfogetattrminmaxrandintr   float_dtypesnormalastypeviewr   RuntimeError)shape	dtype_strr3   lowhighrB   r6   xr   r   r   numpy_random0   s,   


*rQ   rP   returnc                 C   s   | j j}|tv r"|d}| tt|}ttj	||dtt
|S |r5d|v r5ttj	| |dtt
|S |dkrF|dkrFtj	| |d S tj	| |dS )z
    Note: We need dst_type because the type of x can be different from dst_type.
          For example: x is of type `float32`, dst_type is `bfloat16`.
          If dst_type is None, we infer dst_type from x.
    u)devicer8   r   r   )r6   namer@   lstriprI   rC   rA   r   torchtensortlr   )rP   rT   dst_typetsigned_type_namex_signedr   r   r   	to_tritonN   s   
r^   c                 C   sL   t | tjjr
| jS t | tjrtdt| }|	dS t
dt|  )Nz^torch\.(\w+)$r7   znot a triton or torch dtype: )r=   r    languager6   rU   rW   rematchr0   group	TypeErrortype)r6   mr   r   r   torch_dtype_namea   s   
rf   c                 C   sl   t | tr| j  ttt| j	S t | t
jr/| j	t
ju r)|    S |   S td|  )Nz Not a triton-compatible tensor: )r=   r   basecpunumpyrI   rC   rA   rf   r6   rW   Tensorr   float
ValueError)rP   r   r   r   to_numpyl   s   
 rm   c                   C   s   t  otj d dkS )Nr   	   )r+   rW   r&   get_device_capabilityr   r   r   r   supports_tmaw   s   rp   z.Requires TMA support (NVIDIA Hopper or higher))reason)NNNr   )*r   r`   ri   rA   rW   r    triton.languager_   rY   pytestnumpy.randomr   typingr   r   triton.runtime.jitr   r   r?   r@   integral_dtypesrG   dtypesdtypes_with_bfloat16torch_float8_dtypestorch_dtypesr   r$   r+   r.   r2   rQ   ndarrayrj   r^   r0   rf   rm   rp   markskipifrequires_tmar   r   r   r   <module>   s:    
 