o
    ۾i                     @   sd   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r0e
  d	S 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|
v  d S )
Nznvptx64-nvidia-cudacu_rsqrtentryzrsqrt.approx.f32 $0, $1;z=f,fT)side_effectr   zrsqrt.approx.f32)r   Module__name__tripler   add_ir_versionFunctionTypeVoidTypePointerType	FloatTypeFunction	IRBuilderappend_basic_block	InlineAsmloadargscallstoreret_voidNVVMdata_layoutset_cuda_kernelstr
compile_ir
assertTrue)selfmodfntyfnbldrrsqrt_approx_fnty	inlineasmvalresnvvmirptx r-   \/home/ubuntu/.local/lib/python3.10/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)llvmliter   numba.cuda.cudadrvr   numba.cuda.testingr   r   r   r   r   mainr-   r-   r-   r.   <module>   s    