§
    ƒÛg‹  ã                   ó¨   — d dl Zd dlmZmZmZ d dlmZmZ d dl	m
Z
 e
j        rdZndZeefZ G d„ de¦  «        Zed	k    r ej        ¦   «          dS dS )
é    N)ÚcudaÚfloat64Úvoid)ÚunittestÚCUDATestCase)Úconfigé   é   c                   ó   — e Zd Zd„ ZdS )ÚTestCudaLaplacec           
      óÌ  ‡— t          j        t          t          t          ¦  «        dd¬¦  «        d„ ¦   «         Št          j        t          t          d d …d d …f         t          d d …d d …f         t          d d …d d …f         ¦  «        ¦  «        ˆfd„¦   «         }t          j        rd\  }}d}nd\  }}d}t          j        ||ft          j        ¬	¦  «        }t          j        ||ft          j        ¬	¦  «        }|}d
}d}	t          |¦  «        D ]}
d||
df<   d||
df<   Œd}t          t          f}||d         z  ||d         z  f}t          j        |¦  «        }t          j
        ¦   «         }t          j        ||¦  «        }t          j        ||¦  «        }t          j        ||¦  «        }|	|k    r®||k     r¬|                      |j        t          j        k    ¦  «          ||||f         |||¦  «         |                     ||¬¦  «         |                     ¦   «          t          j        |¦  «                             ¦   «         }	|}|}|}|dz  }|	|k    r
||k     °¦d S d S d S d S )NT)ÚdeviceÚinlinec                 ó   — | |k    r| S |S )N© )ÚaÚbs     úZ/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_laplace.pyÚget_maxz3TestCudaLaplace.test_laplace_small.<locals>.get_max   s   € à�1ŠuˆuØ�à�ó    c                 ó  •— t           j                             t          t          ¬¦  «        }t           j        j        }t           j        j        }t           j        j        }t           j        j        }| j	        d         }| j	        d         }	t          j
        d¦  «        \  }
}d|||f<   |dk    rq||dz
  k     rh|
dk    rb|
|	dz
  k     rYd| ||
dz   f         | ||
dz
  f         z   | |dz
  |
f         z   | |dz   |
f         z   z  |||
f<   |||
f         | ||
f         z
  |||f<   t          j        ¦   «          t          dz  }|dk    rH||k     r$ ‰|||f         |||z   |f         ¦  «        |||f<   |dz  }t          j        ¦   «          |dk    °Ht          dz  }|dk    rN||k     r*|dk    r$ ‰|||f         ||||z   f         ¦  «        |||f<   |dz  }t          j        ¦   «          |dk    °N|dk    r|dk    r|d         |||f<   d S d S d S )N©Údtyper   é   é   g      Ð?)r   r   )r   ÚsharedÚarrayÚSM_SIZEr   Ú	threadIdxÚxÚyÚblockIdxÚshapeÚgridÚsyncthreadsÚtpb)ÚAÚAnewÚerrorÚerr_smÚtyÚtxÚbxÚbyÚnÚmÚiÚjÚtr   s                €r   Újocabi_relax_corez=TestCudaLaplace.test_laplace_small.<locals>.jocabi_relax_core   sC  ø€ å”[×&Ò&¥wµgÐ&Ñ>Ô>ˆFå”Ô!ˆBÝ”Ô!ˆBÝ””ˆBÝ””ˆBà”˜”
ˆAØ”˜”
ˆAå”9˜Q‘<”<‰DˆAˆqàˆF�2�r�6‰NØ�AŠvˆv˜!˜a !™eš)˜)¨¨Qª¨°1°q¸1±u²9°9Ø! a¨¨1¨q©5¨¤k°A°a¸¸Q¹°h´KÑ&?Ø()¨!¨a©%°¨(¬ñ'4Ø67¸¸A¹¸q¸´kñ'Bñ C��Q˜�T‘
à!% a¨ d¤¨a°°1°¬gÑ!5��r˜2�v‘åÔÑÔÐõ �q‘ˆAØ�a’%�%Ø˜’6�6Ø%, W¨V°B¸°F¬^¸VÀBÈÁFÈBÀJÔ=OÑ%PÔ%P�F˜2˜r˜6‘NØ�a‘�ÝÔ Ñ"Ô"Ð"ð	 �a’%�%õ �q‘ˆAØ�a’%�%Ø˜’6�6˜b Ašg˜gØ%, W¨V°B¸°F¬^¸VÀBÈÈQÉÀJÔ=OÑ%PÔ%P�F˜2˜r˜6‘NØ�a‘�ÝÔ Ñ"Ô"Ð"ð	 �a’%�%ð �QŠwˆw˜2 š7˜7Ø & t¤��b˜"�f‘��ð ˆw˜7˜7r   )r	   r	   é   )é   r6   iè  r   g�íµ ÷Æ°>g      ð?r   r   )Ústream)r   Újitr   r   r   ÚENABLE_CUDASIMÚnpÚzerosÚranger&   r7   Ú	to_deviceÚ
assertTruer   Úcopy_to_hostÚsynchronizeÚabsÚmax)Úselfr4   ÚNNÚNMÚiter_maxr'   r(   r/   Útolr)   r2   ÚiterÚblockdimÚgriddimÚ
error_gridr7   ÚdAÚdAnewÚderror_gridÚtmpr   s                       @r   Útest_laplace_smallz"TestCudaLaplace.test_laplace_small   s�  ø€ å	Œ•'�'¥7Ñ+Ô+°DÀÐ	FÑ	FÔ	Fð	ð 	ñ 
GÔ	Fð	õ 
Œ•$•w˜q˜q˜q ! ! !˜t”}¥g¨a¨a¨a°°°¨d¤mµW¸Q¸Q¸QÀÀÀ¸T´]ÑCÔCÑ	DÔ	Dð&	-ð &	-ð &	-ð &	-ñ 
EÔ	Dð&	-õP Ô ð 	Ø‰FˆB�ØˆHˆHà‰FˆB�ØˆHåŒH�b˜"�X¥R¤ZÐ0Ñ0Ô0ˆÝŒx˜˜R˜­¬
Ð3Ñ3Ô3ˆàˆàˆØˆå�q‘”ð 	ð 	ˆAØˆAˆa�ˆd‰GØˆD��A�‰JˆJàˆå��:ˆØ˜ !œÑ$ b¨H°Q¬KÑ&7Ð8ˆå”X˜gÑ&Ô&ˆ
å”‘”ˆåŒ^˜A˜vÑ&Ô&ˆÝ”˜t VÑ,Ô,ˆÝ”n Z°Ñ8Ô8ˆà�cŠkˆk˜d Xšo˜oØ�OŠO˜JÔ,µ´
Ò:Ñ;Ô;Ð;à8Ð˜g x°Ð7Ô8¸¸UÀKÑPÔPÐPà×$Ò$ Z¸Ð$Ñ?Ô?Ð?ð ×ÒÑ Ô Ð å”F˜:Ñ&Ô&×*Ò*Ñ,Ô,ˆEð ˆCØˆBØˆEà�A‰IˆDð# �cŠkˆk˜d Xšo˜o˜o˜oˆkˆkˆkˆk˜o˜or   N)Ú__name__Ú
__module__Ú__qualname__rP   r   r   r   r   r      s(   € € € € € ðcð cð cð cð cr   r   Ú__main__)Únumpyr:   Únumbar   r   r   Únumba.cuda.testingr   r   Ú
numba.corer   r9   r&   r   r   rQ   Úmainr   r   r   ú<module>rZ      sÔ   ðØ Ð Ð Ð Ø %Ð %Ð %Ð %Ð %Ð %Ð %Ð %Ð %Ð %Ø 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ð 5Ø Ð Ð Ð Ð Ð ð 
Ôð Ø
€C€Cà
€CØ
ˆsˆ(€ðdð dð dð dð d�lñ dô dð dðN ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   