U
    ?h                     @   s`   d dl mZ d dlmZ d dlmZmZ d dlmZ edG dd deZe	dkr\e
  d	S )
    )ir)nvvm)unittestContextResettingTestCase)skip_on_cudasimz*Inline PTX cannot be used in the simulatorc                   @   s   e Zd Zdd ZdS )TestCudaInlineAsmc                 C   s   t t}d|_t| t t  t t 	 g}t 
||d}t |d}t t 	 t 	 g}t j|dddd}||jd }|||g}|||jd  |  t j|_t| t|}	t|	}
| d	t|
k d S )
Nznvptx64-nvidia-cudaZcu_rsqrtentryzrsqrt.approx.f32 $0, $1;z=f,fT)Zside_effectr   zrsqrt.approx.f32)r   Module__name__Ztripler   Zadd_ir_versionFunctionTypeZVoidTypeZPointerTypeZ	FloatTypeFunctionZ	IRBuilderZappend_basic_blockZ	InlineAsmloadargscallstoreZret_voidZNVVMZdata_layoutZset_cuda_kernelstrZllvm_to_ptxZ
assertTrue)selfmodZfntyfnZbldrZrsqrt_approx_fntyZ	inlineasmvalresZnvvmirptx r   Z/var/www/html/venv/lib/python3.8/site-packages/numba/cuda/tests/cudadrv/test_inline_ptx.pytest_inline_rsqrt
   s*    

 

z#TestCudaInlineAsm.test_inline_rsqrtN)r
   
__module____qualname__r   r   r   r   r   r      s   r   __main__N)Zllvmliter   Znumba.cuda.cudadrvr   Znumba.cuda.testingr   r   r   r   r
   mainr   r   r   r   <module>   s   