o
    8ήc                     @   s   d dl mZ d dlmZmZmZ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rse  dS dS )    )ir)llvm_to_ptxset_cuda_kernelfix_data_layoutget_arch_optionget_supported_ccs)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e	dddd Z
dS )TestNvvmDriverc                 C   s    t  jr
t}t| S t}t| S )N)r   	is_nvvm70metadata_nvvm70metadata_nvvm34nvvmir_generic)selfmetadata r   P/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudadrv/test_nvvm_driver.py
get_nvvmir   s
   zTestNvvmDriver.get_nvvmirc                 C   s6   |   }t|d}| d|v  | d|v  d S )Nutf8simpleave)r   r   decode
assertTrue)r   nvvmirptxr   r   r   test_nvvm_compile_simple   s   z'TestNvvmDriver.test_nvvm_compile_simplec                 C   s   t d}t t  t dg}t j||dd}t |d}|  t	| t
| tt|d}| d|v  trH| d|v  d S | d|v  d S )	Ntest_nvvm_from_llvm    mycudakernel)nameentryr   z.address_size 64z.address_size 32)r   ModuleFunctionTypeVoidTypeIntTypeFunction	IRBuilderappend_basic_blockret_voidr   r   r   strr   r   is64bit)r   mftykernelbldrr    r   r   r   r"      s   
z"TestNvvmDriver.test_nvvm_from_llvmc                 C   sX   dj | }|  }t||ddddd}| dj | | | d| | d| d S )	Nzcompute_{0}{1}   r   )archftz	prec_sqrtprec_divr   z.target sm_{0}{1}r   r   )formatr   r   r   assertIn)r   r6   
compute_xxr   r    r   r   r   _test_nvvm_support-   s   

z!TestNvvmDriver._test_nvvm_supportc                 C   s   t  D ]}| j|d qdS )z"Test supported CC by NVVM
        r6   N)r   r=   )r   r6   r   r   r   test_nvvm_support6   s   
z TestNvvmDriver.test_nvvm_supportTzNo new CC unknown to NVVM yetc              	   C   s`   g }|D ])}dj | }| t}| j|d W d   n1 s!w   Y  | ||j qdS )z>Test unsupported CC to help track the feature support
        z-arch=compute_{0}{1}r>   N)r:   assertRaisesr   r=   r;   msg)r   future_archsr6   patraisesr   r   r   test_nvvm_future_support<   s   
z'TestNvvmDriver.test_nvvm_future_supportN)__name__
__module____qualname__r   r!   r"   r=   r?   r   skipIfrE   r   r   r   r   r      s    	
r   c                   @      e Zd Zdd ZdS )TestArchOptionc                 C   sx   |  tddd |  tddd |  tddd t }|D ]}|  t| d|  q |  tddd|d	   d S )
N      
compute_53   
compute_75zcompute_%d%di  r   )assertEqualr   r   )r   supported_ccr6   r   r   r   test_get_arch_optionM   s   
z#TestArchOption.test_get_arch_optionN)rF   rG   rH   rT   r   r   r   r   rK   K       rK   c                   @   rJ   )TestLibDevicec                 C   s    t  }| |jd d d d S )N   s   BC)r   rR   bc)r   	libdevicer   r   r   test_libdevice_load\   s   z!TestLibDevice.test_libdevice_loadN)rF   rG   rH   rZ   r   r   r   r   rV   Z   rU   rV   a  target triple="nvptx64-"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"

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
z
!nvvmir.version = !{!1}
!1 = !{i32 1, i32 6, i32 3, i32 0}

!nvvm.annotations = !{!2}
!2 = !{void (i32*)* @simple, !"kernel", i32 1}
z\
!nvvm.annotations = !{!1}
!1 = metadata !{void (i32*)* @simple, metadata !"kernel", i32 1}
__main__N)llvmliter   numba.cuda.cudadrv.nvvmr   r   r   r   r   ctypesr   r	   r
   numba.cuda.testingr   r   r   r   r   r0   TestCaser   rK   rV   r   r   r   rF   mainr   r   r   r   <module>   s&    =!	