
    `iT                     `   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	ej                              Z ed           G d
 dej                              Z ed           G d dej                              ZdZedk    r ej                     dS dS )    N)ir)nvvmruntime)unittest)	LibDevice	NvvmErrorNVVM)skip_on_cudasimz(NVVM Driver unsupported in the simulatorc                   J    e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
d	 Zd
 ZdS )TestNvvmDriverc                     t                                                      }t                      j        }t                              ||          S )N)data_layoutv)r	   get_ir_versionr   nvvmir_genericformat)selfversionsr   s      }/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/numba/cuda/tests/cudadrv/test_nvvm_driver.py
get_nvvmirzTestNvvmDriver.get_nvvmir   s<    66((**ff($$$III    c                     |                                  }t          j        |                              d          }|                     d|v            |                     d|v            d S )Nutf8simpleave)r   r   
compile_irdecode
assertTrue)r   nvvmirptxs      r   test_nvvm_compile_simplez'TestNvvmDriver.test_nvvm_compile_simple   s`    ""of%%,,V44C(((%%%%%r   c                     t          j                    dk     r|                     d           |                                 }t	          j        |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   ltoirs      r    test_nvvm_compile_nullary_optionz/TestNvvmDriver.test_nvvm_compile_nullary_option   sy       7**MMHIII""At,OOO 	rr$788888r   c                     d}|                      t          |          5  t          j        dd           d d d            d S # 1 swxY w Y   d S )Nz*-made-up-option=2 is an unsupported option    )made_up_option)assertRaisesRegexr   r   r   )r   msgs     r   test_nvvm_bad_optionz#TestNvvmDriver.test_nvvm_bad_option'   s     ;##Is33 	2 	2OBq1111	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2 	2s   AA	Ac                    t          j        d          }d|_        t          j        |           t          j        t          j                    t          j        d          g          }t          j        ||d          }t          j	        |
                    d                    }|                                 t          j        |           t                      j        |_        t          j        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    s         r   r8   z"TestNvvmDriver.test_nvvm_from_llvm.   s   I+,,(Aobkmmbjnn-=>>Q.999|F55g>>??V$$$*oc!ff%%,,V44#-...*c122222r   c                    t          j        d          }d|_        t                      j        |_        t          j        |           t          j        t          j                    t          j	        d          g          }t          j
        ||d          }t          j        |                    d                    }|                                 t          j        |           d t          |                                          D             }d}|                     t%          |          d	|           |d
         }|                     d|           |                     d|           |                     d|           d S )Ntest_used_listr9   r:   r;   r<   r>   c                     g | ]}d |v |	S )z	llvm.used ).0lines     r   
<listcomp>z1TestNvvmDriver.test_used_list.<locals>.<listcomp>L   s,     . . .t$,, ,,,r   z'Expected exactly one @"llvm.used" array   r   zappending globalzsection "llvm.metadata")r   r?   r@   r	   r   r   rA   rB   rC   rD   rE   rF   rG   rH   rI   rJ   
splitlinesr-   lenassertIn)r   rK   rL   rM   rN   
used_linesr5   	used_lines           r   rP   zTestNvvmDriver.test_used_list=   sL   I&''(*A obkmmbjnn-=>>Q.999|F55g>>??V$$$. .s1vv'8'8':': . . .
7Z!S111qM	ni000()444/;;;;;r   c                 >   t          j        d          }d|_        t                      j        |_        t          j        |           |                     t          d          5  t          j	        t          |                     d d d            d S # 1 swxY w Y   d S )Ntest_bad_irzunknown-unknown-unknownzInvalid target triple)r   r?   r@   r	   r   r   rA   r4   r   r   rJ   )r   rK   s     r   test_nvvm_ir_verify_failz'TestNvvmDriver.test_nvvm_ir_verify_failY   s    Im$$,*A##I/FGG 	$ 	$OCFF###	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$s   #"BBBc                 .    dj         | }|                                 }t          j        ||ddd                              d          }|                      dj         | |           |                     d|           |                     d|           d S )	Nzcompute_{0}{1}rV   r   )r)   ftz	prec_sqrtprec_divr   z.target sm_{0}{1}r   r   )r   r   r   r   r   rY   )r   r)   
compute_xxr   r    s        r   _test_nvvm_supportz!TestNvvmDriver._test_nvvm_supporta   s    ,%,d3
""of:1'(* * **0&.. 	0)0$7===h$$$eS!!!!!r   c                 ^    t          j                    D ]}|                     |           dS )z"Test supported CC by NVVM
        )r)   N)r   get_supported_ccsrd   )r   r)   s     r   test_nvvm_supportz TestNvvmDriver.test_nvvm_supportj   s@     *,, 	/ 	/D###....	/ 	/r   c                    t          j        d          }d|_        t                      j        |_        t          j        |           t          j        t          j                    g           }t          j	        ||d          }t          j
        |                    d                    }|                                 t          j        |           |j                            d           t!          j        d          5 }t          j        t'          |                     d d d            n# 1 swxY w Y   |                     t+          |          d	           |                     d
t'          |d                              d S )Ntest_nvvm_warningr9   inlinekernelr<   r>   noinlineT)recordrV   zoverriding noinline attributer   )r   r?   r@   r	   r   r   rA   rB   rC   rE   rF   rG   rH   rI   
attributesaddwarningscatch_warningsr   rJ   r-   rX   rY   )r   rK   rL   rM   builderws         r   ri   z TestNvvmDriver.test_nvvm_warningp   s}   I)**(*AobkmmR00Q.999,v88AABBV$$$ 	j)))$D111 	$QOCFF###	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	$ 	Q###5s1Q4yyAAAAAs   "D11D58D5N)__name__
__module____qualname__r   r!   r/   r6   r8   rP   r^   rd   rg   ri   rR   r   r   r   r   
   s        J J J
& & &9 9 9 2 2 23 3 3< < <8$ $ $" " "/ / /B B B B Br   r   c                       e Zd Zd ZdS )TestArchOptionc                    |                      t          j        dd          d           |                      t          j        dd          d           |                      t          j        dd          d           t          j                    }|D ]'}|                      t          j        | d|z             (|                      t          j        dd          d|d	         z             d S )
Nr$   r%   
compute_53   
compute_75zcompute_%d%di  r   )r-   r   get_arch_optionrf   )r   supported_ccr)   s      r   test_get_arch_optionz#TestArchOption.test_get_arch_option   s    -a33\BBB-a33\BBB-a33\BBB-//  	Q 	QDT148.4:OPPPP-dA66',r*::	< 	< 	< 	< 	<r   N)rs   rt   ru   r   rR   r   r   rw   rw      s#        
< 
< 
< 
< 
<r   rw   c                       e Zd Zd ZdS )TestLibDevicec                 h    t                      }|                     |j        d d         d           d S )Nr*   s   BC)r   r-   bc)r   	libdevices     r   test_libdevice_loadz!TestLibDevice.test_libdevice_load   s2    KK	bqb)=99999r   N)rs   rt   ru   r   rR   r   r   r   r      s#        : : : : :r   r   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__)ro   llvmliter   numba.cuda.cudadrvr   r   numba.cuda.testingr   numba.cuda.cudadrv.nvvmr   r   r	   r
   TestCaser   rw   r   r   rs   mainrR   r   r   <module>r      s          , , , , , , , , ' ' ' ' ' ' > > > > > > > > > > . . . . . . ;<<xB xB xB xB xBX& xB xB =<xBv ;<<< < < < <X& < < =<< ;<<: : : : :H% : : =<:&R zHMOOOOO r   