
    wgT                     F   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&                          yy)    N)ir)nvvmruntime)unittest)	LibDevice	NvvmErrorNVVM)skip_on_cudasimz(NVVM Driver unsupported in the simulatorc                   H    e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
d	 Zd
 Zy)TestNvvmDriverc                     t               j                         }t               j                  }t        j	                  ||      S )N)data_layoutv)r	   get_ir_versionr   nvvmir_genericformat)selfversionsr   s      n/home/mcse/projects/flask/flask-venv/lib/python3.12/site-packages/numba/cuda/tests/cudadrv/test_nvvm_driver.py
get_nvvmirzTestNvvmDriver.get_nvvmir   s5    6((*f(($$$II    c                     | j                         }t        j                  |      j                  d      }| j	                  d|v        | j	                  d|v        y )Nutf8simpleave)r   r   
compile_irdecode
assertTrue)r   nvvmirptxs      r   test_nvvm_compile_simplez'TestNvvmDriver.test_nvvm_compile_simple   sG    "oof%,,V4C(%r   c                     t        j                         dk  r| j                  d       | j                         }t	        j
                  |dd d      }| j                  |d d d       y )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   s[      7*MMHI"At,O 	r$78r   c                     d}| j                  t        |      5  t        j                  dd       d d d        y # 1 sw Y   y xY w)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 	2OOBq1	2 	2 	2s	   :Ac                 x   t        j                  d      }d|_        t        j                  |       t        j
                  t        j                         t        j                  d      g      }t        j                  ||d      }t        j                  |j                  d            }|j                          t        j                  |       t               j                  |_        t        j                  t!        |            j#                  d      }| j%                  d|v        | j%                  d|v        y )	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    II+,(Aoobkkmbjjn-=>Q.9||F55g>?V$**ooc!f%,,V4#-.*c12r   c                    t        j                  d      }d|_        t               j                  |_        t        j                  |       t        j                  t        j                         t        j                  d      g      }t        j                  ||d      }t        j                  |j                  d            }|j                          t        j                  |       t        |      j!                         D cg c]  }d|v r|
 }}d}| j#                  t%        |      d	|       |d
   }| j'                  d|       | j'                  d|       | j'                  d|       y c c}w )Ntest_used_listr9   r:   r;   r<   r>   z	llvm.usedz'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   line
used_linesr5   	used_lines	            r   rP   z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   E,c                 *   t        j                  d      }d|_        t               j                  |_        t        j                  |       | j                  t        d      5  t        j                  t        |             d d d        y # 1 sw Y   y xY w)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   sj    IIm$,**A##I/FG 	$OOCF#	$ 	$ 	$s   !B		Bc                     dj                   | }| j                         }t        j                  ||ddd      j	                  d      }| j                   dj                   | |       | j                  d|       | j                  d|       y )	Nzcompute_{0}{1}rQ   r   )r)   ftz	prec_sqrtprec_divr   z.target sm_{0}{1}r   r   )r   r   r   r   r   rT   )r   r)   
compute_xxr   r    s        r   _test_nvvm_supportz!TestNvvmDriver._test_nvvm_supporta   s    ,%,,d3
"oof:1'(**0&. 	0)00$7=h$eS!r   c                 Z    t        j                         D ]  }| j                  |        y)z"Test supported CC by NVVM
        )r)   N)r   get_supported_ccsr`   )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                  |j                  d            }|j                          t        j                  |       |j                  j                  d       t!        j"                  d      5 }t        j$                  t'        |             d d d        | j)                  t+              d	       | j-                  d
t'        |d                y # 1 sw Y   CxY w)Ntest_nvvm_warningr9   inlinekernelr<   r>   noinlineT)recordrQ   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-   rS   rT   )r   rK   rL   rM   builderws         r   re   z TestNvvmDriver.test_nvvm_warningp   s   II)*(**AoobkkmR0Q.9,,v88ABV$ 	j)$$D1 	$QOOCF#	$ 	Q#5s1Q4yA		$ 	$s   E&&E/N)__name__
__module____qualname__r   r!   r/   r6   r8   rP   rZ   r`   rc   re    r   r   r   r   
   s6    J
&9 23<8$"/Br   r   c                       e Zd Zd Zy)TestArchOptionc                    | j                  t        j                  dd      d       | j                  t        j                  dd      d       | j                  t        j                  dd      d       t        j                         }|D ]'  }| j                  t        j                  | d|z         ) | j                  t        j                  dd      d|d	   z         y )
Nr$   r%   
compute_53   
compute_75zcompute_%d%di  r   )r-   r   get_arch_optionrb   )r   supported_ccr)   s      r   test_get_arch_optionz#TestArchOption.test_get_arch_option   s    --a3\B--a3\B--a3\B--/  	QDT1148.4:OP	Q--dA6',r*::	<r   N)ro   rp   rq   r|   rr   r   r   rt   rt      s    
<r   rt   c                       e Zd Zd Zy)TestLibDevicec                 V    t               }| j                  |j                  d d d       y )Nr*   s   BC)r   r-   bc)r   	libdevices     r   test_libdevice_loadz!TestLibDevice.test_libdevice_load   s$    K	bq)=9r   N)ro   rp   rq   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__)rk   llvmliter   numba.cuda.cudadrvr   r   numba.cuda.testingr   numba.cuda.cudadrv.nvvmr   r   r	   r
   TestCaser   rt   r~   r   ro   mainrr   r   r   <module>r      s      , ' > > . ;<xBX&& xB =xBv ;<<X&& < =< ;<:H%% : =:&R zHMMO r   