U
    ?h                      @   s   d dl Z d dlmZ d dlmZmZ d dlmZmZm	Z	 d dl
mZ d dlmZmZmZ d dl
mZ e	ee	ekZedG d	d
 d
ejZedG dd dejZedG dd dejZdZdZdZedkre  dS )    N)ir)nvvmruntime)c_size_tc_uint64sizeof)unittest)	LibDevice	NvvmErrorNVVM)skip_on_cudasimz(NVVM Driver unsupported in the simulatorc                   @   sh   e Zd Zdd Zdd Zdd Zdd Zd	d
 Zdd Zdd Z	dd Z
dd Zedddd ZdS )TestNvvmDriverc                 C   s(   t   }t| }t  j}tj||dS )N)data_layoutmetadata)r   Zget_ir_versionmetadata_nvvm70r   nvvmir_genericformat)selfversionsr   r    r   [/var/www/html/venv/lib/python3.8/site-packages/numba/cuda/tests/cudadrv/test_nvvm_driver.py
get_nvvmir   s    
zTestNvvmDriver.get_nvvmirc                 C   s8   |   }t|d}| d|k | d|k d S )Nutf8simpleave)r   r   llvm_to_ptxdecode
assertTrue)r   nvvmirptxr   r   r   test_nvvm_compile_simple   s    z'TestNvvmDriver.test_nvvm_compile_simplec                 C   sH   t  dk r| d |  }tj|dd dd}| |d d d d S )N)      z,-gen-lto unavailable in this toolkit version   Z
compute_52)optZgen_ltoarch   s   CN)r   get_versionZskipTestr   r   r   assertEqual)r   r   Zltoirr   r   r    test_nvvm_compile_nullary_option   s
    
z/TestNvvmDriver.test_nvvm_compile_nullary_optionc              	   C   s.   d}|  t| tjddd W 5 Q R X d S )Nz*-made-up-option=2 is an unsupported option    )Zmade_up_option)assertRaisesRegexr
   r   r   )r   msgr   r   r   test_nvvm_bad_option,   s    z#TestNvvmDriver.test_nvvm_bad_optionc                 C   s   t d}d|_t| t t  t dg}t j||dd}t 	|
d}|  t| t j|_tt|d}| d|k tr| d|k n| d	|k d S )
Ntest_nvvm_from_llvmnvptx64-nvidia-cuda    Zmycudakernelnameentryr   z.address_size 64z.address_size 32)r   Moduletripler   add_ir_versionFunctionTypeVoidTypeZIntTypeFunction	IRBuilderappend_basic_blockret_voidset_cuda_kernelr   r   r   strr   r   is64bit)r   mftykernelZbldrr   r   r   r   r/   3   s    



z"TestNvvmDriver.test_nvvm_from_llvmc              	   C   sN   t d}d|_t j|_t| | td t	t
| W 5 Q R X d S )NZtest_bad_irzunknown-unknown-unknownzInvalid target triple)r   r5   r6   r   r   r   r7   r,   r
   r   r?   )r   rA   r   r   r   test_nvvm_ir_verify_failE   s    


z'TestNvvmDriver.test_nvvm_ir_verify_failc                 C   sZ   dj | }|  }tj||ddddd}| dj | | | d| | d| d S )	Nzcompute_{0}{1}   r   )r%   ZftzZ	prec_sqrtZprec_divr   z.target sm_{0}{1}r   r   )r   r   r   r   r   assertIn)r   r%   Z
compute_xxr   r   r   r   r   _test_nvvm_supportM   s    
z!TestNvvmDriver._test_nvvm_supportc                 C   s   t  D ]}| j|d qdS )z"Test supported CC by NVVM
        r%   N)r   get_supported_ccsrG   )r   r%   r   r   r   test_nvvm_supportV   s    z TestNvvmDriver.test_nvvm_supportc              	   C   s   t d}d|_t j|_t| t t  g }t j	||dd}t 
|d}|  t| |jd tjdd}tt| W 5 Q R X | t|d	 | d
t|d  d S )Ntest_nvvm_warningr0   Zinlinekernelr2   r4   ZnoinlineT)recordrE   zoverriding noinline attributer   )r   r5   r6   r   r   r   r7   r8   r9   r:   r;   r<   r=   r>   
attributesaddwarningscatch_warningsr   r?   r(   lenrF   )r   rA   rB   rC   Zbuilderwr   r   r   rK   \   s    



z TestNvvmDriver.test_nvvm_warningTzNo new CC unknown to NVVM yetc              
   C   sL   g }|D ]>}dj | }| t}| j|d W 5 Q R X | ||j qdS )z>Test unsupported CC to help track the feature support
        z-arch=compute_{0}{1}rH   N)r   ZassertRaisesr
   rG   rF   r-   )r   Zfuture_archsr%   patZraisesr   r   r   test_nvvm_future_supportq   s    
z'TestNvvmDriver.test_nvvm_future_supportN)__name__
__module____qualname__r   r    r)   r.   r/   rD   rG   rJ   rK   r   ZskipIfrT   r   r   r   r   r      s   	
r   c                   @   s   e Zd Zdd ZdS )TestArchOptionc                 C   s   |  tddd |  tddd |  tddd t }|D ]}|  tj| d|  qH|  tddd|d	   d S )
Nr"   r#   Z
compute_53   Z
compute_75zcompute_%d%di  r   )r(   r   Zget_arch_optionrI   )r   Zsupported_ccr%   r   r   r   test_get_arch_option   s    
z#TestArchOption.test_get_arch_optionN)rU   rV   rW   r[   r   r   r   r   rX      s   rX   c                   @   s   e Zd Zdd ZdS )TestLibDevicec                 C   s    t  }| |jd d d d S )Nr&   s   BC)r	   r(   bc)r   Z	libdevicer   r   r   test_libdevice_load   s    z!TestLibDevice.test_libdevice_loadN)rU   rV   rW   r^   r   r   r   r   r\      s   r\   a-  target triple="nvptx64-nvidia-cuda"
target datalayout = "{data_layout}"

define i32 @ave(i32 %a, i32 %b) {{
entry:
%add = add nsw i32 %a, %b
%div = sdiv i32 %add, 2
ret i32 %div
}}

define void @simple(i32* %data) {{
entry:
%0 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
%1 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
%mul = mul i32 %0, %1
%2 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
%add = add i32 %mul, %2
%call = call i32 @ave(i32 %add, i32 %add)
%idxprom = sext i32 %add to i64
%arrayidx = getelementptr inbounds i32, i32* %data, i64 %idxprom
store i32 %call, i32* %arrayidx, align 4
ret void
}}

declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() nounwind readnone

declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x() nounwind readnone

declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() nounwind readnone

{metadata}
z
!nvvmir.version = !{!1}
!1 = !{i32 %s, i32 %s, i32 %s, i32 %s}

!nvvm.annotations = !{!2}
!2 = !{void (i32*)* @simple, !"kernel", i32 1}
z\
!nvvm.annotations = !{!1}
!1 = metadata !{void (i32*)* @simple, metadata !"kernel", i32 1}
__main__)rO   Zllvmliter   Znumba.cuda.cudadrvr   r   ctypesr   r   r   Znumba.cuda.testingr   Znumba.cuda.cudadrv.nvvmr	   r
   r   r   r@   ZTestCaser   rX   r\   r   r   Zmetadata_nvvm34rU   mainr   r   r   r   <module>   s$   r#	