§
    1@Äjî(  ã                   ó`  — d dl Zd dlZd dlmZ d dlmZmZ d dlmZmZ d dl	m
Z
mZmZ d dlmZ d dlmZ d dlmZ d d	lmZmZmZmZmZmZmZ  ej        d
ej        ¬¦  «        Zd„ Zd„ Zd„ Zd„ Z d„ Z!dZ"d„ Z# ed¦  «         G d„ de¦  «        ¦   «         Z$e%dk    r ej&        ¦   «          dS dS )é    N)Úunittest)Úskip_on_cudasimÚskip_if_cuda_includes_missing)ÚCUDATestCaseÚtest_data_dir)ÚCudaAPIErrorÚLinkerÚLinkerError)Ú
NvrtcError)Úrequire_context)Úignore_internal_warnings)ÚcudaÚvoidÚfloat64Úint64Úint32ÚtypeofÚfloat32é
   ©Údtypec                 ó’   — t           j                             t          ¦  «        }t          j        d¦  «        }||         dz   | |<   d S )Né   ç      ð?)r   ÚconstÚ
array_likeÚCONST1DÚgrid)ÚAÚCÚis      úUC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/tests/cudadrv/test_linker.pyÚsimple_const_memr#      s9   € ÝŒ
×Ò�gÑ&Ô&€AÝŒ	�!‰Œ€AàˆQŒ4�#‰:€A€a�D€D€Dó    c                 óœ  — d}d}d}	d}
d}d}d}d}d}d}d}d}d}d}d}d}d}d}d}d}t          |¦  «        D ]f}||z  }||z  }|	|z  }	|
|z  }
||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }||z  }Œg||z   |	z   |
z   |z   | t          j        d¦  «        <   | t          j        d¦  «        xx         ||z   |z   |z   |z   z  cc<   | t          j        d¦  «        xx         ||z   |z   |z   |z   z  cc<   | t          j        d¦  «        xx         ||z   |z   |z   |z   z  cc<   d S )Nr   r   r   )Úranger   r   )ÚxÚaÚbÚcÚdÚeÚfÚa1Úa2Úa3Úa4Úa5Úb1Úb2Úb3Úb4Úb5Úc1Úc2Úc3Úc4Úc5Úd1Úd2Úd3Úd4Úd5r!   s                               r"   Úfunc_with_lots_of_registersrB      sè  € Ø	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BØ	€BÝ�1‰XŒXð ð ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆa‰ˆØ
ˆq‰ˆØ
ˆq‰ˆØ
ˆq‰ˆØ
ˆq‰ˆØ
ˆq‰ˆˆØ˜2‘g ‘l RÑ'¨"Ñ,€A…d„i��l„l�OØ…d„i��l„l€O€O„O�r˜B‘w ‘| bÑ(¨2Ñ-Ñ-€O€O�OØ…d„i��l„l€O€O„O�r˜B‘w ‘| bÑ(¨2Ñ-Ñ-€O€O�OØ…d„i��l„l€O€O„O�r˜B‘w ‘| bÑ(¨2Ñ-Ñ-€O€O�O€O€Or$   c                 óä   — t           j                             d|¦  «        }t          j        d¦  «        }|dk    rt	          d¦  «        D ]}|||<   Œt          j        ¦   «          ||         | |<   d S )Néd   r   r   )r   ÚsharedÚarrayr   r&   Úsyncthreads)ÚaryÚdtyÚsmr!   Újs        r"   Úsimple_smemrL   H   so   € Ý	Œ×	Ò	˜3 Ñ	$Ô	$€BÝŒ	�!‰Œ€AØˆA‚v€vÝ�s‘”ð 	ð 	ˆAØˆBˆq‰EˆEÝÔÑÔÐØ�ŒU€Cˆ�F€F€Fr$   c                 óâ   — t          j        d¦  «        \  }}t           j                             dt          ¦  «        }|dz   |dz   z  |||f<   t          j        ¦   «          |||f         | ||f<   d S )Né   )r   é   r   )r   r   rE   rF   r   rG   )rH   r!   rK   rJ   s       r"   Úcoop_smem2drP   R   sl   € ÝŒ9�Q‰<Œ<�D€A€qÝ	Œ×	Ò	˜8¥WÑ	-Ô	-€BØ�A‘˜!˜a™%Ñ €B€qˆ!€t�HÝÔÑÔÐØ�1�a�4”€Cˆˆ1ˆ�I€I€Ir$   c                 ó8   — t          j        d¦  «        }|| |<   d S ©Nr   )r   r   )rH   r!   s     r"   Úsimple_maxthreadsrS   Z   s   € ÝŒ	�!‰Œ€AØ€Cˆ�F€F€Fr$   iè  c                 óð   — t           j                             t          |¦  «        }t	          |j        d         ¦  «        D ]}| |         ||<   Œt	          |j        d         ¦  «        D ]}||         ||<   Œd S ©Nr   )r   ÚlocalrF   Ú	LMEM_SIZEr&   Úshape)r   ÚBrI   r    r!   s        r"   Úsimple_lmemrZ   b   s{   € ÝŒ
×Ò� CÑ(Ô(€AÝ�1”7˜1”:ÑÔð ð ˆØ�Œtˆˆ!‰ˆÝ�1”7˜1”:ÑÔð ð ˆØ�Œtˆˆ!‰ˆðð r$   z$Linking unsupported in the simulatorc                   óÀ   — e Zd ZddiZed„ ¦   «         Zd„ Zd„ Zd„ Zd„ Z	d„ Z
d	„ Zd
„ Zd„ Ze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S )Ú
TestLinkerÚNUMBA_CUDA_USE_NVIDIA_BINDINGÚ0c                 ó2   — t          j        d¬¦  «        }~dS )z9Simply go through the constructor and destructor
        )é   é   )ÚccN)r	   Únew)ÚselfÚlinkers     r"   Útest_linker_basiczTestLinker.test_linker_basicn   s   € õ ”˜vÐ&Ñ&Ô&ˆØˆFˆFr$   c                 óŠ  — t          j        dd¦  «        at          t          dz  ¦  «        }|rdg}ng }t          j        |d|giŽd„ ¦   «         }t          j        dgt          j        ¬¦  «        }t          j        d	gt          j        ¬¦  «        } |d
         ||¦  «         |  	                    |d         dk    ¦  «         d S )NÚbarúint32(int32)zjitlink.ptxzvoid(int32[:], int32[:])Úlinkc                 ót   — t          j        d¦  «        }| |xx         t          ||         ¦  «        z  cc<   d S rR   )r   r   rh   )r'   Úyr!   s      r"   Úfooz%TestLinker._test_linking.<locals>.foo€   s3   € å”	˜!‘”ˆAØˆaˆDˆDŒD•C˜˜!œ‘I”IÑˆDˆD‰DˆDˆDr$   é{   r   iA  )r   r   r   iý  )
r   Údeclare_devicerh   Ústrr   ÚjitÚnprF   r   Ú
assertTrue)rd   Úeagerrj   Úargsrm   r   rY   s          r"   Ú_test_linkingzTestLinker._test_linkingu   sÎ   € åÔ! %¨Ñ8Ô8ˆå•= =Ñ0Ñ1Ô1ˆàð 	Ø.Ð/ˆDˆDàˆDå	Œ�4Ð	%˜t˜fÐ	%Ð	%ð	ð 	ñ 
&Ô	%ð	õ ŒH�c�U¥"¤(Ð+Ñ+Ô+ˆÝŒH�c�U¥"¤(Ð+Ñ+Ô+ˆàˆˆDŒ	�!�Q‰Œˆà�Š˜˜!œ Ò-Ñ.Ô.Ð.Ð.Ð.r$   c                 ó2   — |                       d¬¦  «         d S )NF©rt   ©rv   ©rd   s    r"   Útest_linking_lazy_compilez$TestLinker.test_linking_lazy_compileŒ   s   € Ø×Ò ÐÑ'Ô'Ð'Ð'Ð'r$   c                 ó2   — |                       d¬¦  «         d S )NTrx   ry   rz   s    r"   Útest_linking_eager_compilez%TestLinker.test_linking_eager_compile�   s   € Ø×Ò ÐÑ&Ô&Ð&Ð&Ð&r$   c                 óx  ‡— t          j        dd¦  «        Št          t          dz  ¦  «        }t          j        |g¬¦  «        ˆfd„¦   «         }t          j        dt
          j        ¬¦  «        }t          j        |¦  «        } |d         ||¦  «         |d	z  }t
          j	         
                    ||¦  «         d S )
Nrh   ri   z
jitlink.cu©rj   c                 ó‚   •— t          j        d¦  «        }|t          | ¦  «        k     r ‰||         ¦  «        | |<   d S d S rR   )r   r   Úlen)Úrr'   r!   rh   s      €r"   Úkernelz*TestLinker.test_linking_cu.<locals>.kernel—   s?   ø€ å”	˜!‘”ˆAà•3�q‘6”6ŠzˆzØ�s˜1˜Qœ4‘y”y��!‘��ð ˆzr$   r   r   )r   é    rN   )r   ro   rp   r   rq   rr   Úaranger   Ú
zeros_likeÚtestingÚassert_array_equal)rd   rj   rƒ   r'   r‚   Úexpectedrh   s         @r"   Útest_linking_cuzTestLinker.test_linking_cu’   sÀ   ø€ ÝÔ! %¨Ñ8Ô8ˆå•= <Ñ/Ñ0Ô0ˆå	Œ˜�vÐ	Ñ	Ô	ð	!ð 	!ð 	!ð 	!ñ 
Ô	ð	!õ ŒI�b¥¤Ð)Ñ)Ô)ˆÝŒM˜!ÑÔˆàˆˆuŒ�a˜ÑÔÐð �q‘5ˆÝ
Œ
×%Ò% a¨Ñ2Ô2Ð2Ð2Ð2r$   c                 ó  ‡— t          j        dd¦  «        Št          t          dz  ¦  «        }t	          j        d¬¦  «        5 }t          ¦   «          t          j        d|g¬¦  «        ˆfd„¦   «         }d d d ¦  «         n# 1 swxY w Y   |                      t          |¦  «        d	d
¦  «         |  
                    dt          |d         j        ¦  «        ¦  «         |  
                    dt          |d         j        ¦  «        ¦  «         d S )Nrh   ri   zwarn.cuT)Úrecordúvoid(int32)r   c                 ó   •—  ‰| ¦  «         d S ©N© ©r'   rh   s    €r"   rƒ   z6TestLinker.test_linking_cu_log_warning.<locals>.kernel¯   ó   ø€ à��A‘”���r$   r   zExpected warnings from NVRTCzNVRTC log messagesr   zdeclared but never referenced)r   ro   rp   r   ÚwarningsÚcatch_warningsr   rq   ÚassertEqualr�   ÚassertInÚmessage)rd   rj   Úwrƒ   rh   s       @r"   Útest_linking_cu_log_warningz&TestLinker.test_linking_cu_log_warning§   s;  ø€ ÝÔ! %¨Ñ8Ô8ˆå•= 9Ñ,Ñ-Ô-ˆåÔ$¨DÐ1Ñ1Ô1ð 	°QÝ$Ñ&Ô&Ð&åŒX�m¨4¨&Ð1Ñ1Ô1ðð ð ð ñ 2Ô1ðð		ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	×Ò�˜Q™œ Ð$BÑCÔCÐCà�ŠÐ*­C°°!´´Ñ,=Ô,=Ñ>Ô>Ð>à�ŠÐ5µs¸1¸Q¼4¼<Ñ7HÔ7HÑIÔIÐIÐIÐIs   Á1B Â BÂBc                 ó°  ‡— t          j        dd¦  «        Št          t          dz  ¦  «        }|                      t
          ¦  «        5 }t          j        d|g¬¦  «        ˆfd„¦   «         }d d d ¦  «         n# 1 swxY w Y   |j        j        d         }|  	                    d|¦  «         |  	                    d	|¦  «         |  	                    d
|¦  «         d S )Nrh   ri   zerror.cur�   r   c                 ó   •—  ‰| ¦  «         d S r�   r�   r‘   s    €r"   rƒ   z0TestLinker.test_linking_cu_error.<locals>.kernel¿   r’   r$   r   zNVRTC Compilation failurez identifier "SYNTAX" is undefinedz in the compilation of "error.cu")
r   ro   rp   r   ÚassertRaisesr   rq   Ú	exceptionru   r–   )rd   rj   r,   rƒ   Úmsgrh   s        @r"   Útest_linking_cu_errorz TestLinker.test_linking_cu_error¹   s  ø€ ÝÔ! %¨Ñ8Ô8ˆå•= :Ñ-Ñ.Ô.ˆà×Ò�zÑ*Ô*ð 	¨aÝŒX�m¨4¨&Ð1Ñ1Ô1ðð ð ð ñ 2Ô1ðð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð
 ŒkÔ˜qÔ!ˆà�ŠÐ1°3Ñ7Ô7Ð7à�ŠÐ8¸#Ñ>Ô>Ð>à�ŠÐ8¸#Ñ>Ô>Ð>Ð>Ð>s   Á#A7Á7A;Á>A;c                 ó²   — d}|                       t          |¦  «        5  t          j        ddg¬¦  «        d„ ¦   «         }d d d ¦  «         d S # 1 swxY w Y   d S )Nz/Don't know how to link file with extension .cuhúvoid()z
header.cuhr   c                  ó   — d S r�   r�   r�   r$   r"   rƒ   z>TestLinker.test_linking_unknown_filetype_error.<locals>.kernelÎ   ó   € à�r$   ©ÚassertRaisesRegexÚRuntimeErrorr   rq   ©rd   Úexpected_errrƒ   s      r"   Ú#test_linking_unknown_filetype_errorz.TestLinker.test_linking_unknown_filetype_errorË   s¬   € ØHˆØ×#Ò#¥L°,Ñ?Ô?ð 	ð 	ÝŒX�h l ^Ð4Ñ4Ô4ðð ñ 5Ô4ðð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	ð 	ó   ž!AÁAÁAc                 ó²   — d}|                       t          |¦  «        5  t          j        ddg¬¦  «        d„ ¦   «         }d d d ¦  «         d S # 1 swxY w Y   d S )Nz-Don't know how to link file with no extensionr¡   Údatar   c                  ó   — d S r�   r�   r�   r$   r"   rƒ   zDTestLinker.test_linking_file_with_no_extension_error.<locals>.kernelÕ   r£   r$   r¤   r§   s      r"   Ú)test_linking_file_with_no_extension_errorz4TestLinker.test_linking_file_with_no_extension_errorÒ   s¬   € ØFˆØ×#Ò#¥L°,Ñ?Ô?ð 	ð 	ÝŒX�h f XÐ.Ñ.Ô.ðð ñ /Ô.ðð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	ð 	rª   c                 ót   — t          t          dz  ¦  «        }t          j        d|g¬¦  «        d„ ¦   «         }d S )Nzcuda_include.cur¡   r   c                  ó   — d S r�   r�   r�   r$   r"   rƒ   z7TestLinker.test_linking_cu_cuda_include.<locals>.kernelß   s   € àˆDr$   )rp   r   r   rq   )rd   rj   rƒ   s      r"   Útest_linking_cu_cuda_includez'TestLinker.test_linking_cu_cuda_includeÙ   sL   € å•=Ð#4Ñ4Ñ5Ô5ˆõ 
Œ�( $ Ð	(Ñ	(Ô	(ð	ð 	ñ 
)Ô	(ð	ð 	ð 	r$   c                 óê   — |                       t          ¦  «        5 }t          j        ddg¬¦  «        d„ ¦   «         }d d d ¦  «         n# 1 swxY w Y   |                      d|j        j        ¦  «         d S )Núvoid(int32[::1])znonexistent.ar   c                 ó   — d| d<   d S rU   r�   )r'   s    r"   r-   z2TestLinker.test_try_to_link_nonexistent.<locals>.få   s   € à��!‘��r$   znonexistent.a not found)rœ   r
   r   rq   r–   r�   ru   )rd   r,   r-   s      r"   Útest_try_to_link_nonexistentz'TestLinker.test_try_to_link_nonexistentã   s¸   € Ø×Ò�{Ñ+Ô+ð 	¨qÝŒXÐ(°Ð/@ÐAÑAÔAðð ñ BÔAðð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	�ŠÐ/°´Ô1AÑBÔBÐBÐBÐBs   ›!AÁAÁAc                 óâ   — t          j        t          ¦  «        } |j        t	          j        d¦  «        gt          d¦  «        ¢R Ž }|                      |                     ¦   «         d¦  «         dS )a  Ensure that the jitted kernel used in the test_set_registers_* tests
        uses more than 57 registers - this ensures that test_set_registers_*
        are really checking that they reduced the number of registers used from
        something greater than the maximum.r„   é   é9   N)	r   rq   rB   Ú
specializerr   Úemptyr&   ÚassertGreaterÚget_regs_per_thread©rd   Úcompileds     r"   Útest_set_registers_no_maxz$TestLinker.test_set_registers_no_maxê   sd   € õ
 ”8Õ7Ñ8Ô8ˆØ&�8Ô&¥r¤x°¡|¤|Ð?µe¸A±h´hÐ?Ð?Ð?ˆØ×Ò˜8×7Ò7Ñ9Ô9¸2Ñ>Ô>Ð>Ð>Ð>r$   c                 óö   —  t          j        d¬¦  «        t          ¦  «        } |j        t	          j        d¦  «        gt          d¦  «        ¢R Ž }|                      |                     ¦   «         d¦  «         d S )Nr¸   ©Úmax_registersr„   r·   ©	r   rq   rB   r¹   rr   rº   r&   ÚassertLessEqualr¼   r½   s     r"   Útest_set_registers_57z TestLinker.test_set_registers_57ó   óo   € Ø-•4”8¨"Ð-Ñ-Ô-Õ.IÑJÔJˆØ&�8Ô&¥r¤x°¡|¤|Ð?µe¸A±h´hÐ?Ð?Ð?ˆØ×Ò˜X×9Ò9Ñ;Ô;¸RÑ@Ô@Ð@Ð@Ð@r$   c                 óö   —  t          j        d¬¦  «        t          ¦  «        } |j        t	          j        d¦  «        gt          d¦  «        ¢R Ž }|                      |                     ¦   «         d¦  «         d S )Né&   rÁ   r„   r·   rÃ   r½   s     r"   Útest_set_registers_38z TestLinker.test_set_registers_38ø   rÆ   r$   c           	      ó   — t          t          d d d…         t          t          t          t          t          t          ¦  «        } t          j        |d¬¦  «        t
          ¦  «        }|                      |                     ¦   «         d¦  «         d S )Nr   rÈ   rÁ   )r   r   r   r   rq   rB   rÄ   r¼   )rd   Úsigr¾   s      r"   Útest_set_registers_eagerz#TestLinker.test_set_registers_eagerý   sl   € Ý•7˜3˜3˜Q˜3”<¥­­uµe½UÅEÑJÔJˆØ2•4”8˜C¨rÐ2Ñ2Ô2Õ3NÑOÔOˆØ×Ò˜X×9Ò9Ñ;Ô;¸RÑ@Ô@Ð@Ð@Ð@r$   c                 óì   — t          t          d d d…         ¦  «        } t          j        |¦  «        t          ¦  «        }|                     ¦   «         }|                      |t          j        ¦  «         d S rR   )	r   r   r   rq   r#   Úget_const_mem_sizeÚassertGreaterEqualr   Únbytes)rd   rË   r¾   Úconst_mem_sizes       r"   Útest_get_const_mem_sizez"TestLinker.test_get_const_mem_size  s`   € Ý•7˜3˜3˜Q˜3”<Ñ Ô ˆØ •4”8˜C‘=”=Õ!1Ñ2Ô2ˆØ!×4Ò4Ñ6Ô6ˆØ×Ò µ´Ñ?Ô?Ð?Ð?Ð?r$   c                 óæ   — t          j        t          ¦  «        } |j        t	          j        d¦  «        gt          d¦  «        ¢R Ž }|                     ¦   «         }|                      |d¦  «         d S )Nr„   r·   r   )	r   rq   rB   r¹   rr   rº   r&   Úget_shared_mem_per_blockr•   )rd   r¾   Úshared_mem_sizes      r"   Útest_get_no_shared_memoryz$TestLinker.test_get_no_shared_memory  sg   € Ý”8Õ7Ñ8Ô8ˆØ&�8Ô&¥r¤x°¡|¤|Ð?µe¸A±h´hÐ?Ð?Ð?ˆØ"×;Ò;Ñ=Ô=ˆØ×Ò˜¨!Ñ,Ô,Ð,Ð,Ð,r$   c                 ó  — t          t          d d d…         t          t          j        ¦  «        ¦  «        } t	          j        |¦  «        t          ¦  «        }|                     ¦   «         }|                      |d¦  «         d S )Nr   i�  )	r   r   r   rr   r   rq   rL   rÔ   r•   )rd   rË   r¾   rÕ   s       r"   Útest_get_shared_mem_per_blockz(TestLinker.test_get_shared_mem_per_block  si   € Ý•5˜˜˜1˜”:�v¥b¤hÑ/Ô/Ñ0Ô0ˆØ •4”8˜C‘=”=¥Ñ-Ô-ˆØ"×;Ò;Ñ=Ô=ˆØ×Ò˜¨#Ñ.Ô.Ð.Ð.Ð.r$   c                 ó  — t          j        t          ¦  «        }|                     t	          j        dt          j        ¬¦  «        t          j        ¦  «        }|                     ¦   «         }|  	                    |d¦  «         d S )NrD   r   i   )
r   rq   rL   r¹   rr   Úzerosr   r   rÔ   r•   )rd   r¾   Úcompiled_specializedrÕ   s       r"   Ú#test_get_shared_mem_per_specializedz.TestLinker.test_get_shared_mem_per_specialized  sl   € Ý”8�KÑ(Ô(ˆØ'×2Ò2ÝŒH�S¥¤Ð)Ñ)Ô)­2¬:ñ 7ô  7Ðà.×GÒGÑIÔIˆØ×Ò˜¨#Ñ.Ô.Ð.Ð.Ð.r$   c                 óž   —  t          j        d¦  «        t          ¦  «        }|                     ¦   «         }|                      |d¦  «         d S )Nzvoid(float32[:,::1])r   )r   rq   rP   Úget_max_threads_per_blockr»   )rd   r¾   Úmax_threadss      r"   Útest_get_max_threads_per_blockz)TestLinker.test_get_max_threads_per_block  sI   € Ø3•4”8Ð2Ñ3Ô3µKÑ@Ô@ˆØ×8Ò8Ñ:Ô:ˆØ×Ò˜;¨Ñ*Ô*Ð*Ð*Ð*r$   c                 óJ  —  t          j        d¦  «        t          ¦  «        }|                     ¦   «         }|dz   }t	          j        |t          j        ¬¦  «        }	  |d|f         |¦  «         d S # t          $ r&}|                      d|j	        ¦  «         Y d }~d S d }~ww xY w)Nr³   r   r   ÚcuLaunchKernel)
r   rq   rS   rÞ   rr   rº   r   r   r–   rž   )rd   r¾   rß   ÚnelemrH   r,   s         r"   Útest_max_threads_exceededz$TestLinker.test_max_threads_exceeded   s¹   € Ø/•4”8Ð.Ñ/Ô/Õ0AÑBÔBˆØ×8Ò8Ñ:Ô:ˆØ˜a‘ˆÝŒh�u¥B¤HÐ-Ñ-Ô-ˆð	3ØˆH�Q˜�XÔ˜sÑ#Ô#Ð#Ð#Ð#øÝð 	3ð 	3ð 	3Ø�MŠMÐ*¨A¬EÑ2Ô2Ð2Ð2Ð2Ð2Ð2Ð2Ð2øøøøð	3øøøs   ÁA2 Á2
B"Á<BÂB"c                 ó|  — t          t          d d d…         t          d d d…         t          t          j        ¦  «        ¦  «        } t	          j        |¦  «        t          ¦  «        }|                     ¦   «         }t          j        t          j        ¦  «        j	        t          z  }|                      ||¦  «         d S rR   )r   r   r   rr   r   rq   rZ   Úget_local_mem_per_threadr   ÚitemsizerW   rÏ   )rd   rË   r¾   Úlocal_mem_sizeÚ	calc_sizes        r"   Útest_get_local_mem_per_threadz(TestLinker.test_get_local_mem_per_thread*  sŽ   € Ý•5˜˜˜1˜”:�u S S q Sœz­6µ"´(Ñ+;Ô+;Ñ<Ô<ˆØ •4”8˜C‘=”=¥Ñ-Ô-ˆØ!×:Ò:Ñ<Ô<ˆÝ”H�RœXÑ&Ô&Ô/µ)Ñ;ˆ	Ø×Ò °	Ñ:Ô:Ð:Ð:Ð:r$   c                 ó°  — t          j        t          ¦  «        }|                     t	          j        t          t          j        ¬¦  «        t	          j        t          t          j        ¬¦  «        t          j        ¦  «        }| 	                    ¦   «         }t	          j
        t          j        ¦  «        j        t          z  }|                      ||¦  «         d S )Nr   )r   rq   rZ   r¹   rr   rÚ   rW   r   r   ræ   r   rç   rÏ   )rd   r¾   rÛ   rè   ré   s        r"   Ú"test_get_local_mem_per_specializedz-TestLinker.test_get_local_mem_per_specialized1  sœ   € Ý”8�KÑ(Ô(ˆØ'×2Ò2ÝŒH•Y¥b¤hÐ/Ñ/Ô/ÝŒH•Y¥b¤hÐ/Ñ/Ô/ÝŒJñ ô  Ðð .×FÒFÑHÔHˆÝ”H�RœZÑ(Ô(Ô1µIÑ=ˆ	Ø×Ò °	Ñ:Ô:Ð:Ð:Ð:r$   N)Ú__name__Ú
__module__Ú__qualname__Ú_NUMBA_NVIDIA_BINDING_0_ENVr   rf   rv   r{   r}   rŠ   r™   rŸ   r©   r®   r   r±   rµ   r¿   rÅ   rÉ   rÌ   rÒ   rÖ   rØ   rÜ   rà   rä   rê   rì   r�   r$   r"   r\   r\   j   s¡  € € € € € à#BÀCÐ"HÐàðð ñ „_ðð/ð /ð /ð.(ð (ð (ð'ð 'ð 'ð3ð 3ð 3ð*Jð Jð Jð$?ð ?ð ?ð$ð ð ðð ð ð #ðð ñ #Ô"ððCð Cð Cð?ð ?ð ?ðAð Að Að
Að Að Að
Að Að Að
@ð @ð @ð-ð -ð -ð/ð /ð /ð/ð /ð /ð+ð +ð +ð
3ð 3ð 3ð;ð ;ð ;ð;ð ;ð ;ð ;ð ;r$   r\   Ú__main__)'Únumpyrr   r“   Únumba.cuda.testingr   r   r   r   r   Únumba.cuda.cudadrv.driverr   r	   r
   Únumba.cuda.cudadrv.errorr   Ú
numba.cudar   Únumba.tests.supportr   Únumbar   r   r   r   r   r   r   r…   r   r#   rB   rL   rP   rS   rW   rZ   r\   rí   Úmainr�   r$   r"   ú<module>rú      sç  ðØ Ð Ð Ð Ø €€€Ø 'Ð 'Ð 'Ð 'Ð 'Ð 'Ø OÐ OÐ OÐ OÐ OÐ OÐ OÐ OØ :Ð :Ð :Ð :Ð :Ð :Ð :Ð :ð4ð 4ð 4ð 4ð 4ð 4ð 4ð 4ð 4ð 4à /Ð /Ð /Ð /Ð /Ð /Ø &Ð &Ð &Ð &Ð &Ð &Ø 8Ð 8Ð 8Ð 8Ð 8Ð 8Ø DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ Dð ˆ"Œ)�B˜bœjÐ
)Ñ
)Ô
)€ðð ð ð-.ð -.ð -.ð`ð ð ðð ð ðð ð ð
 €	ðð ð ð €Ð7Ñ8Ô8ðN;ð N;ð N;ð N;ð N;�ñ N;ô N;ñ 9Ô8ðN;ðb ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr$   