
    shT                     H   S SK r S SKJr  S SKJrJr  S SKJr  S SKJ	r	J
r
Jr  S SKJr  \" S5       " S S	\R                  5      5       r\" S5       " S
 S\R                  5      5       r\" S5       " S S\R                  5      5       rSr\S:X  a  \R&                  " 5         gg)    N)ir)nvvmruntime)unittest)	LibDevice	NvvmErrorNVVM)skip_on_cudasimz(NVVM Driver unsupported in the simulatorc                   P    \ rS rSrS rS rS rS rS rS r	S r
S	 rS
 rS rSrg)TestNvvmDriver
   c                     [        5       R                  5       n[        5       R                  n[        R	                  X!S9$ )N)data_layoutv)r	   get_ir_versionr   nvvmir_genericformat)selfversionsr   s      }/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudadrv/test_nvvm_driver.py
get_nvvmirTestNvvmDriver.get_nvvmir   s3    6((*f(($$$II    c                     U R                  5       n[        R                  " U5      R                  S5      nU R	                  SU;   5        U R	                  SU;   5        g )Nutf8simpleave)r   r   
compile_irdecode
assertTrue)r   nvvmirptxs      r   test_nvvm_compile_simple'TestNvvmDriver.test_nvvm_compile_simple   sG    "oof%,,V4C(%r   c                     [         R                  " 5       S:  a  U R                  S5        U R                  5       n[        R
                  " USS SS9nU R                  US S S5        g )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_option/TestNvvmDriver.test_nvvm_compile_nullary_option   s[      7*MMHI"At,O 	r$78r   c                     SnU R                  [        U5         [        R                  " SSS9  S S S 5        g ! , (       d  f       g = f)Nz*-made-up-option=2 is an unsupported option    )made_up_option)assertRaisesRegexr   r   r   )r   msgs     r   test_nvvm_bad_option#TestNvvmDriver.test_nvvm_bad_option'   s3     ;##Is3OOBq1 433s	   8
Ac                    [         R                  " S5      nSUl        [        R                  " U5        [         R
                  " [         R                  " 5       [         R                  " S5      /5      n[         R                  " XSS9n[         R                  " UR                  S5      5      nUR                  5         [        R                  " U5        [        5       R                  Ul        [        R                  " [!        U5      5      R#                  S5      nU R%                  SU;   5        U R%                  SU;   5        g )	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   r=   "TestNvvmDriver.test_nvvm_from_llvm.   s    II+,(Aoobkkmbjjn-=>Q.9||F55g>?V$**ooc!f%,,V4#-.*c12r   c                    [         R                  " S5      nSUl        [        5       R                  Ul        [
        R                  " U5        [         R                  " [         R                  " 5       [         R                  " S5      /5      n[         R                  " XSS9n[         R                  " UR                  S5      5      nUR                  5         [
        R                  " U5        [        U5      R!                  5        Vs/ s H  nSU;   d  M  UPM     nnSnU R#                  [%        U5      S	U5        US
   nU R'                  SU5        U R'                  SU5        U R'                  SU5        g s  snf )Ntest_used_listr>   r?   r@   rA   rC   z	llvm.usedz'Expected exactly one @"llvm.used" array   r   zappending globalzsection "llvm.metadata")r   rD   rE   r	   r   r   rF   rG   rH   rI   rJ   rK   rL   rM   rN   rO   
splitlinesr0   lenassertIn)	r   rP   rQ   rR   rS   line
used_linesr9   	used_lines	            r   rV   TestNvvmDriver.test_used_list=   s!   II&'(**A oobkkmbjjn-=>Q.9||F55g>?V$ (+1v'8'8': .':t$, ':
 .7Z!S1qM	ni0()4/;.s   
E7E7c                 :   [         R                  " S5      nSUl        [        5       R                  Ul        [
        R                  " U5        U R                  [        S5         [
        R                  " [        U5      5        S S S 5        g ! , (       d  f       g = f)Ntest_bad_irzunknown-unknown-unknownzInvalid target triple)r   rD   rE   r	   r   r   rF   r8   r   r   rO   )r   rP   s     r   test_nvvm_ir_verify_fail'TestNvvmDriver.test_nvvm_ir_verify_failY   sb    IIm$,**A##I/FGOOCF# HGGs   # B
Bc                    SR                   " U6 nU R                  5       n[        R                  " X2SSSS9R	                  S5      nU R                  SR                   " U6 U5        U R                  SU5        U R                  SU5        g )	Nzcompute_{0}{1}rW   r   )r,   ftz	prec_sqrtprec_divr   z.target sm_{0}{1}r   r   )r   r   r   r   r   rZ   )r   r,   
compute_xxr!   r"   s        r   _test_nvvm_support!TestNvvmDriver._test_nvvm_supporta   sw    %,,d3
"oof1'(**0&. 	)00$7=h$eS!r   c                 Z    [         R                  " 5        H  nU R                  US9  M     g)z"Test supported CC by NVVM
        )r,   N)r   get_supported_ccsrh   )r   r,   s     r   test_nvvm_support TestNvvmDriver.test_nvvm_supportj   s(     **,D###. -r   c                    [         R                  " S5      nSUl        [        5       R                  Ul        [
        R                  " U5        [         R                  " [         R                  " 5       / 5      n[         R                  " XSS9n[         R                  " UR                  S5      5      nUR                  5         [
        R                  " U5        UR                  R                  S5        [         R"                  " SS9 n[
        R$                  " ['        U5      5        S S S 5        U R)                  [+        W5      S	5        U R-                  S
['        US   5      5        g ! , (       d  f       NH= f)Ntest_nvvm_warningr>   inlinekernelrA   rC   noinlineT)recordrW   zoverriding noinline attributer   )r   rD   rE   r	   r   r   rF   rG   rH   rJ   rK   rL   rM   rN   
attributesaddwarningscatch_warningsr   rO   r0   rY   rZ   )r   rP   rQ   rR   builderws         r   ro    TestNvvmDriver.test_nvvm_warningp   s    II)*(**AoobkkmR0Q.9,,v88ABV$ 	j)$$D1QOOCF# 2 	Q#5s1Q4yA	 21s    E((
E6 N)__name__
__module____qualname____firstlineno__r   r#   r2   r:   r=   rV   ra   rh   rl   ro   __static_attributes__rz   r   r   r   r   
   s6    J
&9 23<8$"/Br   r   c                       \ rS rSrS rSrg)TestArchOption   c                    U R                  [        R                  " SS5      S5        U R                  [        R                  " SS5      S5        U R                  [        R                  " SS5      S5        [        R                  " 5       nU H)  nU R                  [        R                  " U6 SU-  5        M+     U R                  [        R                  " SS5      SUS	   -  5        g )
Nr'   r(   
compute_53   
compute_75zcompute_%d%di  r   )r0   r   get_arch_optionrk   )r   supported_ccr,   s      r   test_get_arch_option#TestArchOption.test_get_arch_option   s    --a3\B--a3\B--a3\B--/ DT1148.4:OP !--dA6',r*::	<r   rz   N)r{   r|   r}   r~   r   r   rz   r   r   r   r      s    
<r   r   c                       \ rS rSrS rSrg)TestLibDevice   c                 V    [        5       nU R                  UR                  S S S5        g )Nr-   s   BC)r   r0   bc)r   	libdevices     r   test_libdevice_load!TestLibDevice.test_libdevice_load   s$    K	bq)=9r   rz   N)r{   r|   r}   r~   r   r   rz   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__)ru   llvmliter   numba.cuda.cudadrvr   r   numba.cuda.testingr   numba.cuda.cudadrv.nvvmr   r   r	   r
   TestCaser   r   r   r   r{   mainrz   r   r   <module>r      s      , ' > > . ;<xBX&& xB =xBv ;<<X&& < =< ;<:H%% : =:&R zMMO r   