
    xKg                          d dl Zd dlmZmZmZmZ d dlmZ d dl	m
Z
 d dlmZmZmZ ddlmZmZ d Zd	 Zd
 Z ed       G d de             Zedk(  r ej.                          yy)    N)cudaint32
complex128void)types)TypingError)unittestCUDATestCaseskip_on_cudasim   )test_struct_model_type
TestStructc                     t         j                  j                  dt              }t	        |j
                  d         D ]
  }| |   ||<    t	        |j
                  d         D ]
  }||   ||<    y )N  dtyper   r   localarrayr   rangeshapeABCis       i/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_localmem.pyculocalr   
   c    

U+A1771:t! 1771:t!     c                     t         j                  j                  dt              }t	        |j
                  d         D ]
  }| |   ||<    t	        |j
                  d         D ]
  }||   ||<    y )Nd   r   r   )r   r   r   r   r   r   r   s       r   culocalcomplexr#      sc    

J/A1771:t! 1771:t! r    c                     t         j                  j                  dt              }t	        |j
                  d         D ]
  }| |   ||<    t	        |j
                  d         D ]
  }||   ||<    y )N)   r   r   r   r   s       r   culocal1tupler&      r   r    z'PTX inspection not available in cudasimc                       e Zd Zd Zd Zd Zd Z ed      d        Z ed      d        Z	 ed      d        Z
 ed      d	        Zd
 Zd Zd Zd Zy)TestCudaLocalMemc                 h   t         d d  t         d d  f} t        j                  |      t              }| j	                  d|j                  |      v        t        j                  dd      }t        j                  |      } |d   ||       | j	                  t        j                  ||k(               y )Nz.localr   r   r   r   r   )
r   r   jitr   
assertTrueinspect_asmnparange
zeros_likeall)selfsigjculocalr   r   s        r   test_local_arrayz!TestCudaLocalMem.test_local_array$   s    Qxq" 488C=)H$8$8$==>IId'*MM!q!qAv'r    c                      t        j                  d      t              }t        j                  dd      }t        j
                  |      } |d   ||       | j                  t        j                  ||k(               y)zGEnsure that local arrays can be constructed with 1-tuple shape
        zvoid(int32[:], int32[:])r%   r   r   r*   N)r   r+   r&   r.   r/   r0   r,   r1   )r2   r4   r   r   s       r   test_local_array_1_tuplez)TestCudaLocalMem.test_local_array_1_tuple-   sb     848867F IIaw'MM!q!qAv'r    c                    d} t        j                  |      t              }t        j                  dd      dz
  dz  }t        j
                  |      } |d   ||       | j                  t        j                  ||k(               y )Nz"void(complex128[:], complex128[:])r"   r   r   r   y               @r*   )r   r+   r#   r.   r/   r0   r,   r1   )r2   r3   jculocalcomplexr   r   s        r   test_local_array_complexz)TestCudaLocalMem.test_local_array_complex8   sl    2'$((3-7YYs,/!3r9MM!a#qAv'r    c                     t        t        |j                  j                                     j                  }|j
                  d   j                  }| j                  ||       y )Nl)nextiter	overloadsvalues_type_annotationtypemapr   assertEqual)r2   fr   
annotationl_dtypes        r   check_dtypezTestCudaLocalMem.check_dtype@   sJ    $q{{11345FF
$$S)//%(r    zCan't check typing in simulatorc                     t        j                  t        t        d d d               d        }| j	                  |t               y )Nr   c                 n    t         j                  j                  dt              }| d   |d<   |d   | d<   y N
   r   r   )r   r   r   r   xr<   s     r   rD   z,TestCudaLocalMem.test_numba_dtype.<locals>.fJ   s5    

  5 1AQ4AaDQ4AaDr    r   r+   r   r   rG   r2   rD   s     r   test_numba_dtypez!TestCudaLocalMem.test_numba_dtypeG   >     
$uSqSz"	#	 
$	
 	E"r    c                     t        j                  t        t        d d d               d        }| j	                  |t               y )Nr   c                     t         j                  j                  dt        j                        }| d   |d<   |d   | d<   y rJ   )r   r   r   r.   r   rL   s     r   rD   z,TestCudaLocalMem.test_numpy_dtype.<locals>.fU   s9    

  288 4AQ4AaDQ4AaDr    rN   rO   s     r   test_numpy_dtypez!TestCudaLocalMem.test_numpy_dtypeR   rQ   r    c                     t        j                  t        t        d d d               d        }| j	                  |t               y )Nr   c                 f    t         j                  j                  dd      }| d   |d<   |d   | d<   y )NrK   r   r   r   r   r   r   rL   s     r   rD   z-TestCudaLocalMem.test_string_dtype.<locals>.f`   s5    

  7 3AQ4AaDQ4AaDr    rN   rO   s     r   test_string_dtypez"TestCudaLocalMem.test_string_dtype]   rQ   r    c           	          d}| j                  t        |      5  t        j                  t	        t
        d d d               d        }d d d        y # 1 sw Y   y xY w)Nz*.*Invalid NumPy dtype specified: 'int33'.*r   c                 f    t         j                  j                  dd      }| d   |d<   |d   | d<   y )NrK   int33r   r   rW   rL   s     r   rD   z5TestCudaLocalMem.test_invalid_string_dtype.<locals>.fm   s5    JJ$$Rw$7t!t!r    )assertRaisesRegexr   r   r+   r   r   )r2   rerD   s      r   test_invalid_string_dtypez*TestCudaLocalMem.test_invalid_string_dtypeh   sM     :##K4XXd51:&' ( 544s   /AAc                     t        j                  t        t        d d d               d        }| j	                  |t               y )Nr   c                 n    t         j                  j                  dt              }| d   |d<   |d   | d<   y rJ   )r   r   r   r   rL   s     r   rD   z<TestCudaLocalMem.test_type_with_struct_data_model.<locals>.ft   s6    

  +A BAQ4AaDQ4AaDr    )r   r+   r   r   rG   rO   s     r    test_type_with_struct_data_modelz1TestCudaLocalMem.test_type_with_struct_data_models   s?    	$-cc23	4	 
5	
 	23r    c           	         t        j                  t        t        d d d   t        d d d               d        }t	        j
                  dd      }t	        j
                  dd      } |d   ||       t        |      D ]  \  }}| j                  ||        t        |      D ]  \  }}| j                  ||dz          y )Nr   c                 F   t         j                  j                  dt              }t	        t        |            D ](  }t        t        |      t        |dz              }|||<   * t	        t        |            D ]&  }||   j                  | |<   ||   j                  ||<   ( y )NrK   r      )
r   r   r   r   r   lenr   r   rM   y)outxoutyarrr   objs        r   rD   z6TestCudaLocalMem.test_struct_model_type_arr.<locals>.f}   s     **""2-C"DC3s8_ q5Q<8A % 3s8_a&((Qa&((Q %r    )rK   r   r   r*   rd   )r   r+   r   r   r.   r   	enumeraterC   )r2   rD   arrxarryr   rM   rf   s          r   test_struct_model_type_arrz+TestCudaLocalMem.test_struct_model_type_arr|   s    	$uSqSz51:.	/
	# 
0
	# xxW-xxW-$ddODAqQ" $dODAqQA& $r    c                     t         j                  fd       }t        j                  dt        j                        } |d   |       | j                  |d   |       y )Nc                 f    t         j                  j                        }|j                  | d<   y )Nr   r   )r   r   r   size)ari   r   tys     r   sz8TestCudaLocalMem._check_local_array_size_fp16.<locals>.s   s)    **""5"3C88AaDr    r   r   r*   r   )r   r+   r.   zerosfloat16rC   )r2   r   expectedrs   rt   results    ` `  r   _check_local_array_size_fp16z-TestCudaLocalMem._check_local_array_size_fp16   sP    		 
	 !2::.$H-r    c                     | j                  ddt        j                         | j                  ddt        j                         y )Nrd   )ry   r   rv   r.   )r2   s    r   test_issue_fp16_supportz(TestCudaLocalMem.test_issue_fp16_support   s.    ))!Q>))!Q

;r    N)__name__
__module____qualname__r5   r7   r:   rG   r   rP   rT   rX   r^   ra   rn   ry   r{    r    r   r(   r(   "   s    (	(() 67# 8# 67# 8# 67# 8# 67 84'0.<r    r(   __main__)numpyr.   numbar   r   r   r   
numba.corer   numba.core.errorsr   numba.cuda.testingr	   r
   r   extensions_usecasesr   r   r   r#   r&   r(   r|   mainr   r    r   <module>r      sk     / /  ) F F C :;}<| }< <}<@ zHMMO r    