o
    ۾iT                     @   s   d dl Z d dlmZ d dlmZmZ d dlmZ d dlm	Z	m
Z
mZ d dlm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edkr[e  dS dS )    N)ir)nvvmruntime)unittest)	LibDevice	NvvmErrorNVVM)skip_on_cudasimz(NVVM Driver unsupported in the simulatorc                   @   s\   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dd ZdS )TestNvvmDriverc                 C   s    t   }t  j}tj||dS )N)data_layoutv)r   get_ir_versionr   nvvmir_genericformat)selfversionsr    r   ]/home/ubuntu/.local/lib/python3.10/site-packages/numba/cuda/tests/cudadrv/test_nvvm_driver.py
get_nvvmir   s   
zTestNvvmDriver.get_nvvmirc                 C   s8   |   }t|d}| d|v  | d|v  d S )Nutf8simpleave)r   r   
compile_ir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   
compute_52)optgen_ltoarch   s   CN)r   get_versionskipTestr   r   r   assertEqual)r   r   ltoirr   r   r    test_nvvm_compile_nullary_option   s
   
z/TestNvvmDriver.test_nvvm_compile_nullary_optionc                 C   sD   d}|  t| tjddd W d    d S 1 sw   Y  d S )Nz*-made-up-option=2 is an unsupported option    )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|v  | d|v  d S )	Ntest_nvvm_from_llvmnvptx64-nvidia-cuda    mycudakernelnameentryr   z.address_size 64)r   Moduletripler   add_ir_versionFunctionTypeVoidTypeIntTypeFunction	IRBuilderappend_basic_blockret_voidset_cuda_kernelr   r   r   strr   r   )r   mftykernelbldrr   r   r   r   r1   .   s   



z"TestNvvmDriver.test_nvvm_from_llvmc                 C   s   t d}d|_t j|_t| t t  t 	dg}t j
||dd}t |d}|  t| dd t| D }d	}| t|d
| |d }| d| | d| | d| d S )Ntest_used_listr2   r3   r4   r5   r7   c                 S   s   g | ]}d |v r|qS )z	llvm.usedr   ).0liner   r   r   
<listcomp>L   s    z1TestNvvmDriver.test_used_list.<locals>.<listcomp>z'Expected exactly one @"llvm.used" array   r   zappending globalzsection "llvm.metadata")r   r8   r9   r   r   r   r:   r;   r<   r=   r>   r?   r@   rA   rB   rC   
splitlinesr(   lenassertIn)r   rD   rE   rF   rG   
used_linesr/   	used_liner   r   r   rH   =   s    



zTestNvvmDriver.test_used_listc                 C   sd   t d}d|_t j|_t| | td t	t
| W d    d S 1 s+w   Y  d S )Ntest_bad_irzunknown-unknown-unknownzInvalid target triple)r   r8   r9   r   r   r   r:   r.   r   r   rC   )r   rD   r   r   r   test_nvvm_ir_verify_failY   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}rL   r   )r$   ftz	prec_sqrtprec_divr   z.target sm_{0}{1}r   r   )r   r   r   r   r   rO   )r   r$   
compute_xxr   r   r   r   r   _test_nvvm_supporta   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_ccsrX   )r   r$   r   r   r   test_nvvm_supportj   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 d    n1 sQw   Y  | t|d	 | d
t|d  d S )Ntest_nvvm_warningr2   inlinekernelr5   r7   noinlineT)recordrL   zoverriding noinline attributer   )r   r8   r9   r   r   r   r:   r;   r<   r>   r?   r@   rA   rB   
attributesaddwarningscatch_warningsr   rC   r(   rN   rO   )r   rD   rE   rF   builderwr   r   r   r[   p   s   



z TestNvvmDriver.test_nvvm_warningN)__name__
__module____qualname__r   r   r*   r0   r1   rH   rS   rX   rZ   r[   r   r   r   r   r
   
   s    	r
   c                   @      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|  q$|  tddd|d	   d S )
Nr   r    
compute_53   
compute_75zcompute_%d%di  r   )r(   r   get_arch_optionrY   )r   supported_ccr$   r   r   r   test_get_arch_option   s   
z#TestArchOption.test_get_arch_optionN)re   rf   rg   rp   r   r   r   r   ri          ri   c                   @   rh   )TestLibDevicec                 C   s    t  }| |jd d d d S )Nr%   s   BC)r   r(   bc)r   	libdevicer   r   r   test_libdevice_load   s   z!TestLibDevice.test_libdevice_loadN)re   rf   rg   ru   r   r   r   r   rr      rq   rr   a3  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

!nvvmir.version = !{{!1}}
!1 = !{{i32 {v[0]}, i32 {v[1]}, i32 {v[2]}, i32 {v[3]}}}

!nvvm.annotations = !{{!2}}
!2 = !{{void (i32*)* @simple, !"kernel", i32 1}}

@"llvm.used" = appending global [1 x i8*] [i8* bitcast (void (i32*)* @simple to i8*)], section "llvm.metadata"
__main__)ra   llvmliter   numba.cuda.cudadrvr   r   numba.cuda.testingr   numba.cuda.cudadrv.nvvmr   r   r   r	   TestCaser
   ri   rr   r   re   mainr   r   r   r   <module>   s     {)