
    J/Ph                         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     d/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_laplace.pyget_maxz3TestCudaLaplace.test_laplace_small.<locals>.get_max   s    1uu    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   [&&wg&>>F!B!BBB
A
A9Q<<DAqF2r6NAvv!a!e))Q1q1u99!a1q5kAaQhK&?()!a%('467Aqk'B CQT
!%ada1g!5r2v qAa%%66%,WVBF^VBFBJ=O%P%PF2r6Na """	 a%% qAa%%66bAgg%,WVBF^VBQJ=O%P%PF2r6Na """	 a%% Qww277 &tb"f w77r   )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	 
$wqqq!!!t}gaaadmWQQQT]CC	D	D&	- &	- &	- &	- 
E	D&	-P   	FBHHFBHHb"XRZ000xR
333q 	 	AAadGDAJJ:!$bHQK&78Xg&&
^Av&&tV,,nZ88ckkdXooOOJ,
:;;;8gx78UKPPP$$Z$???    F:&&**,,E CBEAID# ckkdXooookkkkoor   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       
 
CC
C
s(d d d d dl d d dN zHMOOOOO r   