o
    8ήc                     @   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}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| t| t|}	t|	}
| dt|
v  d S )	Ncu_rsqrtentryzrsqrt.approx.f32 $0, $1;z=f,fT)side_effectr   zrsqrt.approx.f32)r   Module__name__FunctionTypeVoidTypePointerType	FloatTypeFunction	IRBuilderappend_basic_block	InlineAsmloadargscallstoreret_voidr   fix_data_layoutset_cuda_kernelstrllvm_to_ptx
assertTrue)selfmodfntyfnbldrrsqrt_approx_fnty	inlineasmvalresnvvmirptx r*   O/tmp/pip-target-vg8gfxp4/lib/python/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    