
    xKg8                         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Zd dlmZ ddlmZmZ  ej*                  d	ej                  fd
ej,                  dfg      Z G d de      Z G d de      Zedk(  r e
j6                          yy)    )cudaint32float64void)TypingError)types)unittestCUDATestCaseskip_on_cudasimN)numpy_support   )test_struct_model_type
TestStructij)      c                   <    e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
y	)
TestSharedMemoryIssuec                     t        j                  d      d        t         j                  fd       } |d           y )NT)devicec                  N    t         j                  j                  dt              } y Nr   dtyper   sharedarrayr   )	inner_arrs    c/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_sm.pyinnerzGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.inner   s    ))!5)9I    c                  ^    t         j                  j                  dt              }          y r   r   )	outer_arrr!   s    r    outerzGTestSharedMemoryIssue.test_issue_953_sm_linkage_conflict.<locals>.outer   s!    ))!5)9IGr"   r   r   )r   jit)selfr%   r!   s     @r    "test_issue_953_sm_linkage_conflictz8TestSharedMemoryIssue.test_issue_953_sm_linkage_conflict   sC    			: 
	: 
	 
	 	dr"   c                     t         j                  fd       }t        j                  dt        j                        } |d   |       | j                  |d   |       y )Nc                 n    t         j                  j                  t              }|j                  | d<   y Nr   r   )r   r   r   r   size)aarrshapes     r    sz9TestSharedMemoryIssue._check_shared_array_size.<locals>.s   s)    ++##E#7C88AaDr"   r   r   r&   r   )r   r'   npzerosr   assertEqual)r(   r0   expectedr1   results    `   r    _check_shared_array_sizez.TestSharedMemoryIssue._check_shared_array_size   sP    		 
	 !288,$H-r"   c                 (    | j                  dd       y Nr   r7   r(   s    r    %test_issue_1051_shared_size_broken_1dz;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_1d&   s    %%a+r"   c                 (    | j                  dd       y )N)r   r      r:   r;   s    r    %test_issue_1051_shared_size_broken_2dz;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_2d)   s    %%fa0r"   c                 (    | j                  dd       y )N)r   r         r:   r;   s    r    %test_issue_1051_shared_size_broken_3dz;TestSharedMemoryIssue.test_issue_1051_shared_size_broken_3d,   s    %%i4r"   c                     t         j                  fd       }t        j                  dt        j                        } |d   |       | j                  |d   |       y )Nc                 f    t         j                  j                        }|j                  | d<   y r,   )r   r   r   r-   )r.   r/   r0   tys     r    r1   z>TestSharedMemoryIssue._check_shared_array_size_fp16.<locals>.s1   s)    ++##E#4C88AaDr"   r   r   r&   r   )r   r'   r2   r3   float16r4   )r(   r0   r5   rF   r1   r6   s    ` `  r    _check_shared_array_size_fp16z3TestSharedMemoryIssue._check_shared_array_size_fp160   sP    		 
	 !2::.$H-r"   c                     | j                  ddt        j                         | j                  ddt        j                         y r9   )rH   r   rG   r2   r;   s    r    test_issue_fp16_supportz-TestSharedMemoryIssue.test_issue_fp16_support:   s.    **1a?**1a<r"   c                     dd}dd}t         j                  fd       }t        j                  |t        j                        }t        j
                  |      } |||f   |       t        j                          y)z
        Test issue of warp misalign address due to nvvm not knowing the
        alignment(? but it should have taken the natural alignment of the type)
        r   0   rA   r   c                    t         j                  j                  ft              }t         j                  j                  dt              }t         j                  j
                  }d}t              D ]  }||||f   z  } |d   |z   | d<   y )N   r   )r   r   r   r   	threadIdxxrange)d_block_costs
s_featuress_initialcostrO   
predictionr   examples_per_blocknum_weightss         r    
costs_funcz9TestSharedMemoryIssue.test_issue_2393.<locals>.costs_funcH   s    **,>+L+24J KK--a9M((IJ;'jA66
 (  -Q/*<M!r"   r   N)r   r'   r2   r3   r   	to_devicesynchronize)r(   
num_blocksthreads_per_blockrX   block_costsrR   rV   rW   s         @@r    test_issue_2393z%TestSharedMemoryIssue.test_issue_2393>   sv    
 
		= 
	= hhz<{31
:001-@r"   N)__name__
__module____qualname__r)   r7   r<   r?   rC   rH   rJ   r^    r"   r    r   r      s*    
.,15.=r"   r   c                       e 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d      d        Z ed      d        Zy)TestSharedMemoryc                 T   t        |      }dt        |z        }t        j                  |j                        t
        j                  fd       }t        j                  |      } ||f   ||       |j                         }t        j                  j                  ||       y )N   c                    t         j                  j                  	      }t         j                  j                  }t         j
                  j                  }t         j                  j                  }||z  |z   }|t        |       k  r| |   ||<   t        j                          |dk(  rt        	      D ]  }||   |||z  |z   <    y y r,   
r   r   r   rO   rP   blockIdxblockDimlensyncthreadsrQ   )
rP   ysmtxbxbdr   r   dtnthreadss
           r    use_sm_chunk_copyz8TestSharedMemory._test_shared.<locals>.use_sm_chunk_copyj   s    ""82"6B!!BBB R"A3q6z12 QwxA%'UAb2gkN ) r"   )rk   intnps
from_dtyper   r   r'   device_array_likecopy_to_hostr2   testingassert_array_equal)	r(   r/   nelemnblocksrt   d_resulthost_resultrr   rs   s	          @@r    _test_sharedzTestSharedMemory._test_shared_   s     Ceh&'^^CII&		+ 
	+& ))#.,'8+,S(;++-


%%c;7r"   c                 &   t        j                  dt              }t        t	        |            D ]N  }|||   _        t        j                  dt         j                        }|j                  dd      |z  ||   _	        P | j                  |       y )N   r   r>   r   r   )r2   recarrayrecordwith2darrayrQ   rk   r   arangefloat32reshaper   r   )r(   r/   rP   r   s       r    test_shared_recarrayz%TestSharedMemory.test_shared_recarray   sn    kk#%67s3xACFH		%rzz2AyyA*CFH !
 	#r"   c                     t         j                  j                  ddt         j                        }| j	                  |       y )Nr   )   )r-   r   )r2   randomrandintbool_r   )r(   r/   s     r    test_shared_boolz!TestSharedMemory.test_shared_bool   s/    iirxx@#r"   c                     |j                   |j                  j                  z  } |ddd|f   |       t        j                  j                  ||       y )Nr   r   )r-   r   itemsizer2   rz   r{   )r(   funcr/   r5   nshareds        r    _test_dynshared_slicez&TestSharedMemory._test_dynshared_slice   sH    
 ((SYY///Q1gs#


%%h4r"   c                     t         j                  d        }t        j                  dt        j                        }t        j
                  ddgt        j                        }| j                  |||       y )Nc                     t         j                  j                  dt              }|dd }|dd }d|d<   d|d<   |d   | d<   |d   | d<   y Nr   r   r   r   r   rP   dynsmemsm1sm2s       r    slice_writez@TestSharedMemory.test_dynshared_slice_write.<locals>.slice_write   s]    kk'''7G!A,C!A,CCFCF1:AaD1:AaDr"   r   r   r   r   r'   r2   r3   r   r   r   )r(   r   r/   r5   s       r    test_dynshared_slice_writez+TestSharedMemory.test_dynshared_slice_write   sX    		 
	 hhq)88QF"((3"";X>r"   c                     t         j                  d        }t        j                  dt        j                        }t        j
                  ddgt        j                        }| j                  |||       y )Nc                     t         j                  j                  dt              }|dd }|dd }d|d<   d|d<   |d   | d<   |d   | d<   y r   r   r   s       r    
slice_readz>TestSharedMemory.test_dynshared_slice_read.<locals>.slice_read   s]    kk'''7G!A,C!A,CGAJGAJq6AaDq6AaDr"   r   r   r   r   )r(   r   r/   r5   s       r    test_dynshared_slice_readz*TestSharedMemory.test_dynshared_slice_read   sX    		 
	 hhq)88QF"((3"":sH=r"   c                     t         j                  d        }t        j                  dt        j                        }t        j
                  g dt        j                        }| j                  |||       y )Nc                     t         j                  j                  dt              }|dd }|dd }d|d<   d|d<   d|d<   |d   | d<   |d   | d<   |d   | d<   y )Nr   r   r   r   r   r   r   s       r    slice_diff_sizeszJTestSharedMemory.test_dynshared_slice_diff_sizes.<locals>.slice_diff_sizes   ss    kk'''7G!A,C!A,CGAJGAJGAJq6AaDq6AaDq6AaDr"   r   r   )r   r   r   r   )r(   r   r/   r5   s       r    test_dynshared_slice_diff_sizesz0TestSharedMemory.test_dynshared_slice_diff_sizes   sW     

	 

	 hhq)88IRXX6""#3S(Cr"   c                     t         j                  d        }t        j                  dt        j                        }t        j
                  g dt        j                        }| j                  |||       y )Nc                     t         j                  j                  dt              }|dd }|dd }d|d<   d|d<   d|d<   d|d<   |d   | d<   |d   | d<   |d   | d<   |d   | d<   |d   | d<   y )Nr   r   r   r   rA   r   r   r   s       r    slice_overlapzDTestSharedMemory.test_dynshared_slice_overlap.<locals>.slice_overlap   s    kk'''7G!A,C!A,CGAJGAJGAJGAJq6AaDq6AaDq6AaDq6AaDq6AaDr"      r   )r   r   r   r   rA   r   )r(   r   r/   r5   s       r    test_dynshared_slice_overlapz-TestSharedMemory.test_dynshared_slice_overlap   sT    		 
	 hhq)88O288<""=#x@r"   c                     t         j                  d        }t        j                  dt        j                        }t        j
                  g dt        j                        }| j                  |||       y )Nc                 @   t         j                  j                  dt              }|dd }|dd }d|d<   d|d<   d|d<   d|d<   d|d<   d|d	<   d|d<   d|d<   d|d<   d|d<   d|d<   |d   | d<   |d   | d<   |d   | d<   |d   | d<   |d   | d<   |d	   | d	<   |d   | d<   y )
Nr   r   r   r   rA   r>   c   r   r   r   r   s       r    
slice_gapsz>TestSharedMemory.test_dynshared_slice_gaps.<locals>.slice_gaps   s    kk'''7G!A,C!A,C GAJGAJGAJGAJGAJGAJGAJCFCFCFCF1:AaD1:AaD1:AaD1:AaD1:AaD1:AaD1:AaDr"   rN   r   )r   r   r   r   r   rA   r   r   )r(   r   r/   r5   s       r    test_dynshared_slice_gapsz*TestSharedMemory.test_dynshared_slice_gaps   sW     
	 
	6 hhq)884BHHE"":sH=r"   c                     t         j                  d        }t        j                  dt        j                        }t        j
                  g dt        j                        }| j                  |||       y )Nc                     t         j                  j                  dt              }|dd d   }|ddd   }d|d<   d|d<   d|d<   d|d<   |d   | d<   |d   | d<   |d   | d<   |d   | d<   y )Nr   r   r   r   r   rA   r   r   s       r    slice_write_backwardszTTestSharedMemory.test_dynshared_slice_write_backwards.<locals>.slice_write_backwards  s    kk'''7G!%R%.C!Ab&/CCFCFCFCF1:AaD1:AaD1:AaD1:AaDr"   rA   r   )r   r   rA   r   r   )r(   r   r/   r5   s       r    $test_dynshared_slice_write_backwardsz5TestSharedMemory.test_dynshared_slice_write_backwards  sW     
	 
	 hhq)88L9""#8#xHr"   c                     t         j                  d        }t        j                  dt        j                        }t        j
                  g dt        j                        }| j                  |||       y )Nc                    t         j                  j                  dt              }|d d d   }d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   d|d<   |d   | d<   |d   | d<   |d   | d<   |d   | d<   |d   | d<   |d   | d<   y )	Nr   r   r   r   r   r   rA   r   r   rP   r   r   s      r    slice_nonunit_stridezRTestSharedMemory.test_dynshared_slice_nonunit_stride.<locals>.slice_nonunit_stride!  s    kk'''7G#A#,C GAJGAJGAJGAJGAJGAJCFCFCF1:AaD1:AaD1:AaD1:AaD1:AaD1:AaDr"   r>   r   )r   r   r   r   r   r   r   )r(   r   r/   r5   s       r    #test_dynshared_slice_nonunit_stridez4TestSharedMemory.test_dynshared_slice_nonunit_stride  sX     
	 
	. hhq)881B""#7hGr"   c                     t         j                  d        }t        j                  dt        j                        }t        j
                  g dt        j                        }| j                  |||       y )Nc                    t         j                  j                  dt              }|dd d   }d|d<   d|d<   d|d<   d|d<   d|d	<   d|d
<   d|d<   d|d<   d|d<   |d   | d<   |d   | d<   |d   | d<   |d   | d<   |d	   | d	<   |d
   | d
<   y )Nr   r   r   r   r   r   r   rA   r   r   r   s      r    slice_nonunit_reverse_stridezbTestSharedMemory.test_dynshared_slice_nonunit_reverse_stride.<locals>.slice_nonunit_reverse_stride@  s    kk'''7G"&b&/C GAJGAJGAJGAJGAJGAJCFCFCF1:AaD1:AaD1:AaD1:AaD1:AaD1:AaDr"   r>   r   )r   r   r   r   r   r   r   )r(   r   r/   r5   s       r    +test_dynshared_slice_nonunit_reverse_stridez<TestSharedMemory.test_dynshared_slice_nonunit_reverse_stride=  sX     
	 
	. hhq)881B""#?hOr"   c                   
 t        j                  d      }t        |      }d}t        ||z        }t	        j
                  |j                        
||j                  j                  z  }t        |dz        }t        j                  
fd       }t        j                  |      } |||d|f   |||       |j                         }	t         j                  j                  ||	       y )Nr   rf   r   c                    t         j                  j                  d      }|d| }|||dz   }t         j                  j                  }t         j
                  j                  }t         j                  j                  }||z  |z   }	|	t        |       k  r||k  r	| |	   ||<   n| |	   |||z
  <   t        j                          |dk(  r0t        |      D ]!  }
||
   |||z  |
z   <   ||
   |||z  |
z   |z   <   # y y )Nr   r   r   rh   )rP   rm   	chunksizer   r   r   ro   rp   rq   r   r   rr   s              r    sm_slice_copyz7TestSharedMemory.test_issue_5073.<locals>.sm_slice_copyk  s    kk'''4G!I&C)IM2C!!BBB R"A3q6z	>dCG*+A$CY' Qwy)A%(VAb2gkN14QAb2gkI-. * r"   r   )r2   r   rk   ru   rv   rw   r   r   r   r'   rx   ry   rz   r{   )r(   r/   r|   rs   r}   r   r   r   r~   r   rr   s             @r    test_issue_5073z TestSharedMemory.test_issue_5073\  s     iioCeh&'^^CII&SYY///1%			8 
	82 ))#.4gxG34S(IN++-


%%c;7r"   zCan't check typing in simulatorc                 P   d}d }| j                  t        |      5   t        j                  t	                     |       d d d        d}d }| j                  t        |      5   t        j                  t	                     |       d d d        y # 1 sw Y   QxY w# 1 sw Y   y xY w)Nz+.*Cannot infer the type of variable 'arr'.*c                  l    t         j                  j                  dt        j                  d            } y )N
   Or   )r   r   r   r2   r   r/   s    r    unsupported_typezBTestSharedMemory.test_invalid_array_type.<locals>.unsupported_type  s#    ++##Bbhhsm#<Cr"   z*.*Invalid NumPy dtype specified: 'int33'.*c                  F    t         j                  j                  dd      } y )Nr   int33r   )r   r   r   r   s    r    invalid_string_typezETestSharedMemory.test_invalid_array_type.<locals>.invalid_string_type  s    ++##Bg#6Cr"   )assertRaisesRegexr   r   r'   r   )r(   rgxr   r   s       r    test_invalid_array_typez(TestSharedMemory.test_invalid_array_type  s    ;	=##K5DHHTV-. 6 ;	7##K5DHHTV01 65 65 65s   $B#$BBB%z+Struct model array unsupported in simulatorc           	         dt        j                  t        t        d d d   t        d d d               fd       }t	        j
                  fd      }t	        j
                  fd      } |df   ||       t        |      D ]  \  }}| j                  ||z
  dz
          t        |      D ]   \  }}| j                  ||z
  dz
  dz         " y )N@   r   c                    t         j                  j                  t              }t        j                  d      }|z
  dz
  }|t        |       k  rn|t        |      k  r_t        t        |      t        |dz              }|||<   t        j                          ||   j                  | |<   ||   j                  ||<   y y y )Nr   r   r   )r   r   r   r   gridrk   r   r   rl   rP   rm   )outxoutyr/   r   riobjrs   s         r    write_then_reverse_read_staticzVTestSharedMemory.test_struct_model_type_static.<locals>.write_then_reverse_read_static  s     ++##H4J#KC		!AA!B3t9}SY q5Q<8A  "b'))Qb'))Q "/}r"   r   r   r   )r   r'   r   r   r2   r3   	enumerater4   )r(   r   arrxarryr   rP   rm   rs   s          @r    test_struct_model_type_staticz.TestSharedMemory.test_struct_model_type_static  s    	$uSqSz51:.	/	$ 
0	$" xx73xx733&q({3D$?dODAqQ1q 01 $dODAqQA!1Q 67 $r"   N)r_   r`   ra   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   rb   r"   r    rd   rd   ^   s    "8H5?">"D(A,!>FI,H>P>,8\ 672 82 BC8 D8r"   rd   __main__)numbar   r   r   r   numba.core.errorsr   
numba.corer   numba.cuda.testingr	   r
   r   numpyr2   numba.npr   rv   extensions_usecasesr   r   r   r   r   r   rd   r_   mainrb   r"   r    <module>r      s    , , )  F F  ) CBHHsBHHo"BJJ79 : LL L^Z8| Z8z
 zHMMO r"   