o
    8ήc                     @   s  d dl mZmZ d dlmZ d dlmZmZmZ d dl	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d Zdd ZdZdZdZd Zd!Zd"Zd#Zd$Zd%Zd&Z d'Z!d(Z"d)Z#d*Z$eed+d,d-d. Z%eed/d,d0d1 Z&eed+d,d2d3 Z'eed/d,d4d5 Z(eed+d,d6d7 Z)eed+d,d8d9 Z*eed/d,d:d; Z+eed/d,d<d= Z,eed+d,d>d? Z-eed/d,d@dA Z.eed+d,dBdC Z/eed/d,dDdE Z0eed+d,dFdG Z1eed/d,dHdI Z2edJG dKdL dLeZ3e4dMkre5  dS dS )N    )cudanjit)overload)CUDATestCaseskip_on_cudasimunittestNc                   C      d S N r
   r
   r
   L/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudapy/test_overload.pygeneric_func_1
      r   c                   C   r   r	   r
   r
   r
   r
   r   cuda_func_1   r   r   c                   C   r   r	   r
   r
   r
   r
   r   generic_func_2   r   r   c                   C   r   r	   r
   r
   r
   r
   r   cuda_func_2   r   r   c                   C   r   r	   r
   r
   r
   r
   r   generic_calls_generic   r   r   c                   C   r   r	   r
   r
   r
   r
   r   generic_calls_cuda   r   r   c                   C   r   r	   r
   r
   r
   r
   r   cuda_calls_generic"   r   r   c                   C   r   r	   r
   r
   r
   r
   r   cuda_calls_cuda&   r   r   c                   C   r   r	   r
   r
   r
   r
   r   target_overloaded*   r   r   c                   C   r   r	   r
   r
   r
   r
   r   generic_calls_target_overloaded.   r   r   c                   C   r   r	   r
   r
   r
   r
   r   cuda_calls_target_overloaded2   r   r   c                   C   r   r	   r
   r
   r
   r
   r   )target_overloaded_calls_target_overloaded6   r   r                                    %   )   +   generic)targetc                 C      dd }|S )Nc                 S      | d  t 9  < d S Nr   )GENERIC_FUNCTION_1xr
   r
   r   implS      zol_generic_func_1.<locals>.implr
   r.   r/   r
   r
   r   ol_generic_func_1Q      r2   r   c                 C   r)   )Nc                 S   r*   r+   )CUDA_FUNCTION_1r-   r
   r
   r   r/   Z   r0   zol_cuda_func_1.<locals>.implr
   r1   r
   r
   r   ol_cuda_func_1X   r3   r5   c                 C   r)   )Nc                 S   r*   r+   )GENERIC_FUNCTION_2r-   r
   r
   r   r/   a   r0   zol_generic_func_2.<locals>.implr
   r1   r
   r
   r   ol_generic_func_2_   r3   r7   c                 C   r)   )Nc                 S   r*   r+   )CUDA_FUNCTION_2r-   r
   r
   r   r/   h   r0   zol_cuda_func.<locals>.implr
   r1   r
   r
   r   ol_cuda_funcf   r3   r9   c                 C   r)   )Nc                 S      | d  t 9  < t|  d S r+   )GENERIC_CALLS_GENERICr   r-   r
   r
   r   r/   o      z&ol_generic_calls_generic.<locals>.implr
   r1   r
   r
   r   ol_generic_calls_genericm      r=   c                 C   r)   )Nc                 S   r:   r+   )GENERIC_CALLS_CUDAr   r-   r
   r
   r   r/   w   r<   z#ol_generic_calls_cuda.<locals>.implr
   r1   r
   r
   r   ol_generic_calls_cudau   r>   r@   c                 C   r)   )Nc                 S   r:   r+   )CUDA_CALLS_GENERICr   r-   r
   r
   r   r/      r<   z#ol_cuda_calls_generic.<locals>.implr
   r1   r
   r
   r   ol_cuda_calls_generic}   r>   rB   c                 C   r)   )Nc                 S   r:   r+   )CUDA_CALLS_CUDAr   r-   r
   r
   r   r/      r<   z ol_cuda_calls_cuda.<locals>.implr
   r1   r
   r
   r   ol_cuda_calls_cuda   r>   rD   c                 C   r)   )Nc                 S   r*   r+   )GENERIC_TARGET_OLr-   r
   r
   r   r/      r0   z*ol_target_overloaded_generic.<locals>.implr
   r1   r
   r
   r   ol_target_overloaded_generic   r3   rF   c                 C   r)   )Nc                 S   r*   r+   )CUDA_TARGET_OLr-   r
   r
   r   r/      r0   z'ol_target_overloaded_cuda.<locals>.implr
   r1   r
   r
   r   ol_target_overloaded_cuda   r3   rH   c                 C   r)   )Nc                 S   r:   r+   )GENERIC_CALLS_TARGET_OLr   r-   r
   r
   r   r/      r<   z0ol_generic_calls_target_overloaded.<locals>.implr
   r1   r
   r
   r   "ol_generic_calls_target_overloaded   r>   rJ   c                 C   r)   )Nc                 S   r:   r+   )CUDA_CALLS_TARGET_OLr   r-   r
   r
   r   r/      r<   z-ol_cuda_calls_target_overloaded.<locals>.implr
   r1   r
   r
   r   ol_cuda_calls_target_overloaded   r>   rL   c                 C   r)   )Nc                 S   r:   r+   )!GENERIC_TARGET_OL_CALLS_TARGET_OLr   r-   r
   r
   r   r/      r<   z8ol_generic_calls_target_overloaded_generic.<locals>.implr
   r1   r
   r
   r   *ol_generic_calls_target_overloaded_generic   r>   rN   c                 C   r)   )Nc                 S   r:   r+   )CUDA_TARGET_OL_CALLS_TARGET_OLr   r-   r
   r
   r   r/      r<   z5ol_generic_calls_target_overloaded_cuda.<locals>.implr
   r1   r
   r
   r   'ol_generic_calls_target_overloaded_cuda   r>   rP   z$Overloading not supported in cudasimc                   @   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d Zdd Zdd Zdd Zdd ZdS ) TestOverloadc                 C   s6   t jdt jd}t|d | | |d | d S )N   dtype)rR   rR   r   )nponesint32r   jitassertEqualselfkernelexpectedr.   r
   r
   r   check_overload   s   zTestOverload.check_overloadc                 C   s0   t jdt jd}t|| | |d | d S )NrR   rS   r   )rU   rV   rW   r   rY   rZ   r
   r
   r   check_overload_cpu   s   zTestOverload.check_overload_cpuc                 C      dd }t }| || d S )Nc                 S      t |  d S r	   )r   r-   r
   r
   r   r\         z)TestOverload.test_generic.<locals>.kernel)r,   r^   r[   r\   r]   r
   r
   r   test_generic      zTestOverload.test_genericc                 C   r`   )Nc                 S   ra   r	   )r   r-   r
   r
   r   r\      rb   z&TestOverload.test_cuda.<locals>.kernel)r4   r^   rc   r
   r
   r   	test_cuda   re   zTestOverload.test_cudac                 C       dd }t t }| || d S )Nc                 S      t |  t|  d S r	   )r   r   r-   r
   r
   r   r\         z2TestOverload.test_generic_and_cuda.<locals>.kernel)r,   r4   r^   rc   r
   r
   r   test_generic_and_cuda      z"TestOverload.test_generic_and_cudac                 C   rg   )Nc                 S   rh   r	   )r   r   r-   r
   r
   r   r\      ri   z8TestOverload.test_call_two_generic_calls.<locals>.kernel)r,   r6   r^   rc   r
   r
   r   test_call_two_generic_calls   rk   z(TestOverload.test_call_two_generic_callsc                 C   rg   )Nc                 S   rh   r	   )r   r   r-   r
   r
   r   r\      ri   z5TestOverload.test_call_two_cuda_calls.<locals>.kernel)r4   r8   r^   rc   r
   r
   r   test_call_two_cuda_calls   rk   z%TestOverload.test_call_two_cuda_callsc                 C   rg   )Nc                 S   ra   r	   )r   r-   r
   r
   r   r\      rb   z7TestOverload.test_generic_calls_generic.<locals>.kernel)r;   r,   r^   rc   r
   r
   r   test_generic_calls_generic      z'TestOverload.test_generic_calls_genericc                 C   rg   )Nc                 S   ra   r	   )r   r-   r
   r
   r   r\      rb   z4TestOverload.test_generic_calls_cuda.<locals>.kernel)r?   r4   r^   rc   r
   r
   r   test_generic_calls_cuda   ro   z$TestOverload.test_generic_calls_cudac                 C   rg   )Nc                 S   ra   r	   )r   r-   r
   r
   r   r\      rb   z4TestOverload.test_cuda_calls_generic.<locals>.kernel)rA   r,   r^   rc   r
   r
   r   test_cuda_calls_generic   ro   z$TestOverload.test_cuda_calls_genericc                 C   rg   )Nc                 S   ra   r	   )r   r-   r
   r
   r   r\     rb   z1TestOverload.test_cuda_calls_cuda.<locals>.kernel)rC   r4   r^   rc   r
   r
   r   test_cuda_calls_cuda  ro   z!TestOverload.test_cuda_calls_cudac                 C   r`   )Nc                 S   ra   r	   )r   r-   r
   r
   r   r\   
  rb   z8TestOverload.test_call_target_overloaded.<locals>.kernel)rG   r^   rc   r
   r
   r   test_call_target_overloaded	  re   z(TestOverload.test_call_target_overloadedc                 C   rg   )Nc                 S   ra   r	   )r   r-   r
   r
   r   r\     rb   zATestOverload.test_generic_calls_target_overloaded.<locals>.kernel)rI   rG   r^   rc   r
   r
   r   $test_generic_calls_target_overloaded  ro   z1TestOverload.test_generic_calls_target_overloadedc                 C   rg   )Nc                 S   ra   r	   )r   r-   r
   r
   r   r\     rb   z>TestOverload.test_cuda_calls_target_overloaded.<locals>.kernel)rK   rG   r^   rc   r
   r
   r   !test_cuda_calls_target_overloaded  ro   z.TestOverload.test_cuda_calls_target_overloadedc                 C   s4   dd }t t }| || tt }| || d S )Nc                 S   ra   r	   )r   r-   r
   r
   r   r\     rb   zKTestOverload.test_target_overloaded_calls_target_overloaded.<locals>.kernel)rO   rG   r^   rM   rE   r_   rc   r
   r
   r   .test_target_overloaded_calls_target_overloaded  s
   z;TestOverload.test_target_overloaded_calls_target_overloadedN)__name__
__module____qualname__r^   r_   rd   rf   rj   rl   rm   rn   rp   rq   rr   rs   rt   ru   rv   r
   r
   r
   r   rQ      s     rQ   __main__)6numbar   r   numba.core.extendingr   numba.cuda.testingr   r   r   numpyrU   r   r   r   r   r   r   r   r   r   r   r   r   r,   r4   r6   r8   r;   r?   rA   rC   rE   rG   rI   rK   rM   rO   r2   r5   r7   r9   r=   r@   rB   rD   rF   rH   rJ   rL   rN   rP   rQ   rw   mainr
   r
   r
   r   <module>   s~    	




























o