
    ΅i                      |   S SK r S SKJrJrJr  S SKJr  \ R                  " \S5      r\ R                  " \=(       d    \S5      r	\ R                  " \S5      r
\" 5       (       Ga  S SKrS SKJr  S SKrS\\R                      4S	 jr\R$                    SKS j5       r\R$                    SKS j5       r\R$                      SLS j5       r\R$                      SLS j5       r\R.                  " \R                   " S
S0SSS9\R                   " S
S0SSS9\R                   " S
S0SSS9\R                   " S
S0SSS9// S9\R$                    SKS j5       5       r\R.                  " \R                   " S
S0SSS9\R                   " S
S0SSS9\R                   " S
S0SSS9\R                   " S
S0SSS9// S9\R$                    SKS j5       5       r\R.                  " \R                   " S
S0SSS9// S9\R$                    SKS j5       5       r\R.                  " \R                   " SSS.SSS9\R                   " SSS.SSS9\R                   " SSS.SSS9\R                   " SSS.SSS9// S9\R$                      SMS j5       5       rS r\R.                  " \R                   " S
S0SSS9\R                   " S
S0SSS9// SSS \0S!9\R$                    SKS" j5       5       r\R$                    SKS# j5       r\R$                    SKS$ j5       r\R$                      SMS% j5       r \R$                    SKS& j5       r!\R$                      SMS' j5       r"\R$                    SKS( j5       r#\R$                    SKS) j5       r$\R$                    SKS* j5       r%\R$                    SKS+ j5       r&\R$                  S, 5       r'\R$                      SNS- j5       r(\R$                      SOS. j5       r)\R$                      SPS/ j5       r*\R$                      SPS0 j5       r+\R$                  S
\RX                  4S1 j5       r-\R$                  S
\RX                  4S2 j5       r.S S3K/J0r0J1r1  \R$                    SKS4 j5       r2\R$                    SKS5 j5       r3\R$                    SKS6 j5       r4\R$                    SKS7 j5       r5\R$                    SKS8 j5       r6\R.                  " \R                   " SSSSS9.SSS9\R                   " SSS:SS9.SSS9// S;QS9\R$                  S<\RX                  S=\RX                  S>\RX                  S?\RX                  4S@ j5       5       r7\R$                  S
\RX                  4SA j5       r8\R$                  S
\RX                  4SB j5       r9\R$                  S
\RX                  4SC j5       r:\R$                  S
\RX                  4SD j5       r;\R$                  S
\RX                  4SE j5       r<\R$                  S
\RX                  4SF j5       r=\R$                    SKSG j5       r> SQSH\\?   SI\@4SJ jjrAgg)R    N)HAS_CUDA_AND_TRITONHAS_GPUHAS_XPU_AND_TRITON)
has_tritonzrequires cuda and tritonzrequires gpu and tritonzrequires gpu)languagereturnc                  ,   [         R                  R                  (       a>  [        R                  " SSSSSSSS.SSS9[        R                  " SSSSSSSS.SSS9/n U $ [        R                  " SSSSS	.SSS9[        R                  " SSS
SS	.SSS9/n U $ )N            )BLOCK_SIZE_MBLOCK_SIZE_NBLOCK_SIZE_KGROUP_SIZE_Mmatrix_instr_nonkdimwaves_per_eukpack
num_stages	num_warps   @   r   r   r   r          )torchversionhiptritonConfig)configss    ^/home/james-whalen/.local/lib/python3.13/site-packages/torch/testing/_internal/triton_utils.py_get_strange_configsr$      s    ==(*(*(*()02()!"  ! (+(*(*()02()!"  !Gf + (*(*(*()	  !	 (+(*(*()	  !	G,     
BLOCK_SIZEc                    [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
X-   n[         R                  " X'-   XS9  g Nr   axismasktl
program_idarangeloadstorein_ptr0in_ptr1out_ptr
n_elementsr&   pidblock_startoffsetsr,   xyoutputs               r#   
add_kernelr>   S   u     mm#&		!Z 88#GGG%D1GGG%D1
"F6r%   c                    [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
X-
  n[         R                  " X'-   XS9  g r(   r-   r3   s               r#   
sub_kernelrA   d   r?   r%   c                    [         R                  " SS9nXe-  nU[         R                  " SU5      -   nX:  n	[         R                  " X-   U	S9n
US:X  a  [         R                  " X-   U	S9nX-   nOU
n[         R                  " X(-   XS9  g Nr   r)   r+   twor-   )r4   r5   r6   r7   ARGS_PASSEDr&   r8   r9   r:   r,   r;   r<   r=   s                r#   add_kernel_with_optional_paramrF   u   s     mm#&		!Z 88#GGG%D1%)5AUFF
"F6r%   c                     [         R                  " SS9nXv-  nU[         R                  " SU5      -   n	X:  n
[         R                  " X	U-  -   U
S9nUS:X  a  [         R                  " X-   U
S9nX-   nOUn[         R                  " X)U-  -   XS9  g rC   r-   )r4   r5   r6   r7   striderE   r&   r8   r9   r:   r,   r;   r<   r=   s                 r#   -add_kernel_with_none_param_and_equal_to_1_argrI      s     mm#&		!Z 88#GGG..T:%)5AUFF
V++V?r%   r   r   r   r   r   r   )r"   keyc                    [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
X-   n[         R                  " X'-   XS9  g r(   r-   r3   s               r#   add_kernel_autotunedrL      u    " mm#&		!Z 88#GGG%D1GGG%D1
"F6r%   c                    [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
X-
  n[         R                  " X'-   XS9  g r(   r-   r3   s               r#   sub_kernel_autotunedrO      rM   r%   r
   r   c                    [         R                  " SS9nXS-  nU[         R                  " SU5      -   nXr:  n[         R                  " X-   US9n	[         R                  " X-   US9n
X-   n[         R                  " XG-   XS9  g r(   r-   )r4   r5   r7   r&   r6   r8   r9   r:   r,   r;   r<   r=   s               r#   &add_kernel_autotuned_weird_param_orderrQ      su      mm#&		!Z 88#GGG%D1GGG%D1
"F6r%   )BLOCK_SIZE_XBLOCK_SIZE_Yc                    [         R                  " S5      U-  nU[         R                  " SU5      S S 2S 4   -   nX:  n	[         R                  " S5      U-  n
U
[         R                  " SU5      S S S 24   -   nX:  nUnUn[         R                  " XX>-  -   -   X-  5      n[         R                  " XXM-  -   -   X-  5      nUU-   n[         R                  " X-X>-  -   -   UX-  5        g )Nr      r-   )r4   r5   r6   
x_elements
y_elementsrR   rS   xoffsetxindexxmaskyoffsetyindexymaskx1y0tmp0tmp1tmp2s                     r#   add_kernel_2d_autotunedrc      s    6 --"\1299Q5ag>>#--"\1299Q5dAg>>#www
"895=Iwww
"895=Id{
*/23T5=Ir%   c                     U $ )N )r"   ___s      r#   _dummy_early_config_prunerh     s    r%   
      early_config_prune)r"   rJ   warmuprepprune_configs_byc                    [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
X-   n[         R                  " X'-   XS9  g r(   r-   r3   s               r#   *add_kernel_autotuned_with_unsupported_argsrp     su    $ mm#&		!Z 88#GGG%D1GGG%D1
"F6r%   c                    [         R                  " SS9nXe-  nU[         R                  " SU5      -   nX:  n	[         R                  " X-   U	S9n
[         R                  " X-   U	S9nX-   U-  n[         R                  " X(-   XS9  g r(   r-   )r4   r5   r6   r7   scaling_factorr&   r8   r9   r:   r,   r;   r<   r=   s                r#   add_kernel_with_scalingrs   3  sz     mm#&		!Z 88#GGG%D1GGG%D1%>)
"F6r%   c                    [         R                  " SS9nXC-  n[         R                  " U U/U/[         R                  5      n[         R                  " UU/U/[         R                  5      nXg-   n[         R                  " UUU/5        g Nr   r)   r.   r/   _experimental_descriptor_loadfloat32_experimental_descriptor_store	in_desc_ptr0in_desc_ptr1out_desc_ptrr&   r8   offsetabr=   s	            r#   add_kernel_with_tma_1d_old_apir   E  s     mm#!,,HLJJ	
 ,,HLJJ	
 
))H	
r%   c                 B   [         R                  " SS9n[         R                  " SS9nXS-  nXd-  n[         R                  " U Xx/X4/[         R                  5      n	[         R                  " UXx/X4/[         R                  5      n
X-   n[         R                  " UUXx/5        g Nr   r)   rU   rv   r{   r|   r}   rR   rS   pid_xpid_yoffset_xoffset_yr;   r<   r=   s               r#   add_kernel_with_tma_2d_old_apir   d  s     1%1%'',, (JJ	
 ,, (JJ	
 
)) 	
r%   c                     [         R                  " SS9nXC-  n[         R                  " U U/5      n[         R                  " UU/5      nXg-   n[         R                  " UU/U5        g ru   r.   r/   load_tensor_descriptorstore_tensor_descriptorrz   s	            r#   add_kernel_with_tma_1d_new_apir     sn     mm#!%%H
 %%H

 
""H	
r%   c                     [         R                  " SS9n[         R                  " SS9nXS-  nXd-  n[         R                  " U Xx/5      n	[         R                  " UXx/5      n
X-   n[         R                  " UXx/U5        g r   r   r   s               r#   add_kernel_with_tma_2d_new_apir     s     1%1%''%% 
 %% 

 
"" 	
r%   c                    UnUS-   nUS-   n	[         R                  R                  R                  UU Xf/X4/U R                  R
                  S9  [         R                  R                  R                  UUXf/X4/UR                  R
                  S9  [         R                  R                  R                  U	UXf/X4/UR                  R
                  S9  [         R                  R                  R                  U5        [         R                  R                  R                  U5        [         R                  R                  R                  U	5        [         R                  " SS9n
[         R                  " SS9nX-  nX-  n[         R                  " UX/Xf/[         R                  5      n[         R                  " UX/Xf/[         R                  5      nX-   n[         R                  " U	UX/5        g )Nr      )desc_ptrglobal_address	load_sizeglobal_size
element_tyr   r)   rU   )r.   extracuda&experimental_device_tensormap_create2ddtyper   )experimental_tensormap_fenceproxy_acquirer/   rw   rx   ry   )a_ptrb_ptrc_ptrmn	workspacer&   
a_desc_ptr
b_desc_ptr
c_desc_ptrr   r   r   r   r   r   r=   s                    r#    add_kernel_on_device_tma_old_apir     s    
_
_

<< !.{{-- 	= 	
 	<< !.{{-- 	= 	
 	<< !.{{-- 	= 	
 	??
K
??
K
??
K1%1%%% ,, $JJ	
 ,, $JJ	
  	)) 	
r%   c                    [         R                  " U X4/US/Xf/S9n[         R                  " UX4/US/Xf/S9n[         R                  " UX4/US/Xf/S9n	[         R                  " SS9n
[         R                  " SS9nX-  nX-  n[         R                  " UX/5      n[         R                  " UX/5      nX-   n[         R                  " U	X/U5        g )NrU   )baseshapestridesblock_shaper   r)   )r.   make_tensor_descriptorr/   r   r   )r   r   r   r   r   r   r&   a_descb_descc_descr   r   r   r   r   r   r=   s                    r#    add_kernel_on_device_tma_new_apir     s     **&F#0	
 **&F#0	
 **&F#0	
 1%1%%% %% 
 %% 
  	"" 	
r%   c                     [         R                  " SS9nXC-  nU[         R                  " SU5      -   nXb:  n[         R                  " X-   US9nSU-  n	[         R                  " X-   XS9  g Nr   r)   r+   r   r-   )
r4   r6   r7   r&   r8   r9   r:   r,   r;   r=   s
             r#   mul2_kernelr   :  sd     mm#&		!Z 88#GGG%D1Q
"F6r%   c                     [         R                  " SS9nX2-  nU[         R                  " SU5      -   nXQ:  n[         R                  " X-   US9nSU-  n[         R                  " X-   XS9  g r   r-   )	ptrr7   r&   r8   r9   r:   r,   r;   r=   s	            r#   mul2_inplace_kernelr   I  sb     mm#&		!Z 88#GGCM-Q
2r%   c                 8    [         R                  " U S:  U S5      $ )Nr   )r.   where)r;   s    r#   	zero_negsr   W  s    xxQ1%%r%   c                    [         R                  " SS9nXS-  nU[         R                  " SU5      -   nXr:  nUS:X  a  [        XUS9  OUS:X  a  [	        X XUS9  [         R
                  " X-   US9n	[         R                  " X-   XS9  g )Nr   r)   r   )r&   r>   r+   )r.   r/   r0   r   r>   r1   r2   )
r4   r6   r7   r&   
ACTIVATIONr8   r9   r:   r,   r;   s
             r#   indirection_kernelr   [  s     mm#&		!Z 88#..
K<'wTGGG%D1
"A1r%   c                    [         R                  " SS9n[         R                  " SS9nXd-  nXu-  n	U[         R                  " SU5      -   n
U	[         R                  " SU5      -   nUS S 2S 4   U-  U
S S S 24   -   nUS S 2S 4   U-  U
S S S 24   -   n[         R                  " X-   5      n[         R                  " X-   US-  5        g )Nr   r)   rU   g       @r-   )in_ptrr6   in_y_strideout_y_strideX_BLOCK_SIZEY_BLOCK_SIZExidyidx_starty_start	x_offsets	y_offsetssrc_offsetsdst_offsetssrcs                  r#   double_strided_kernelr   n  s     mm#mm#$$bii<88	bii<88	4(;6479KK4(<7)D!G:LLggf*+
&c	2r%   c           	         [         R                  " U [         R                  " SU5      -   5      n[         R                  " U[         R                  " SU5      -   5      n[         R                  " U/U[         R                  5      n[         R
                  " SSXVU/[         R                  SSS9n[         R                  " U[         R                  " SU5      -   U5        g )Nr   shf.l.wrap.b32 $0, $1, $2, $3;
=r,r, r, rTrU   r   is_purepackr.   r1   r0   fullint32inline_asm_elementwiser2   	XYZr   BLOCKr;   r<   szs	            r#   inline_asm_kernel_is_pure_truer     s     GGA		!U++,GGA		!U++,GGUGQ)%%,1I((
 	RYYq%((!,r%   c           	         [         R                  " U [         R                  " SU5      -   5      n[         R                  " U[         R                  " SU5      -   5      n[         R                  " U/U[         R                  5      n[         R
                  " SSXVU/[         R                  SSS9n[         R                  " U[         R                  " SU5      -   U5        g )Nr   r   r   FrU   r   r   r   s	            r#   inline_asm_kernel_is_pure_falser     s     GGA		!U++,GGA		!U++,GGUGQ)%%,1I((
 	RYYq%((!,r%   c                 j   [         R                  " SS9nXT-  n[         R                  " [         R                  " U U/S/U/U/S/S9S/S9n[         R                  " [         R                  " UU/S/U/U/S/S9S/S9nXx-   n	[         R                  " [         R                  " UU/S/U/U/S/S9U	S/S9  g Nr   r)   rU   )r   r   r   r:   r   order)boundary_checkr.   r/   r1   make_block_ptrr2   )
x_ptry_ptr
output_ptrr7   r&   r8   r9   r;   r<   r=   s
             r#   add_kernel_with_block_ptrr     s     mm#&GG!l$'Lc 3

 GG!l$'Lc 3

 
!l$'Lc 3	
r%   c                    [         R                  " SS9nXC-  n[         R                  " [         R                  " U US/SS/US/US/SS/S9S/S9nUn[         R                  " [         R                  " UUS/SS/US/US/SS/S9US/S9  g r   r   )r   r   r7   r&   r8   r9   r;   r=   s           r#   kernel_with_block_ptr_2dr     s     mm#&GG!1oA$a('O!f 3

 
!1oA$a('O!f 3	
r%   )r1   r2   c                     [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[        X-   US9n	[        X-   US9n
X-   n[	        X'-   XS9  g r(   r-   r3   s               r#   add_kernel_with_importr     si     mm#&		!Z 88#".".g3r%   c                 @   [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
[         R                  " S5      S:X  a  X-   nOX-  n[         R                  " X'-   XS9  g r(   r-   r3   s               r#   cond_op_kernelr     s     mm#&		!Z 88#GGG%D1GGG%D1==q UFUF
"F6r%   c                    [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
X-   n[         R                  " X'-   XS9  g r(   )r.   r/   r0   r1   
atomic_addr3   s               r#   atomic_add_kernelr     su     mm#&		!Z 88#GGG%D1GGG%D1
g';r%   c                    [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXs:  n[         R                  " X-   US9n	[         R                  " X-   US9n
[	        S5       H  nX-   n[         R
                  " X'-   XS9  M      SnUS:  a)  US-  nX-   n[         R
                  " X'-   XS9  US:  a  M(  g g )Nr   r)   r+   r   rU   )r.   r/   r0   r1   ranger2   )r4   r5   r6   r7   r&   r8   r9   r:   r,   r;   r<   rf   r=   is                 r#   add_4_times_kernelr   ,  s     mm#&		!Z 88#GGG%D1GGG%D1qAUFHHW&:  !eFAUFHHW&: !er%   c                    [         R                  " SS9nXT-  nU[         R                  " SU5      -   nXr:  n[         R                  " X-   US9n	[         R                  " X-   US9n
X-   n[         R                  " X7-   XS9  g r(   r-   )r4   r5   r7   r6   r&   r8   r9   r:   r,   r;   r<   r=   s               r#   add_kernel_out_of_order_fn2r   C  r?   r%   r   r   )M_ptrNKr   r   r   r   c
                    [         R                  " SS9n
[         R                  " U5      nUS:X  a	  US:  a  SnOUS:X  a  g [         R                  " X5      n[         R                  " XG5      nX-  nX-  nX-  n[	        UU-
  U	5      nUX-  U-  -   nX-  U-  nUU-  [         R
                  " SU5      -   U-  nUU-  [         R
                  " SU5      -   U-  n[         R
                  " SU5      nU US S 2S 4   US S S 24   -   -   nUUS S 2S 4   US S S 24   -   -   n[         R                  " Xg4[         R                  S9n[        [         R                  " XX5      5       Hq  n[         R                  " UUS S S 24   UUU-  -
  :  SS9n[         R                  " UUS S 2S 4   UUU-  -
  :  SS9n[         R                  " UUU5      nUU-  nUU-  nMs     UR                  [         R                  5      nUU-  [         R
                  " SU5      -   nUU-  [         R
                  " SU5      -   nUUS S 2S 4   -   US S S 24   -   n US S 2S 4   U:  US S S 24   U:  -  n![         R                  " U UU!S9  g )	Nr   r)   r   i   r   g        )r,   otherr+   )r.   r/   r1   cdivminr0   zerosrx   r   dottofloat16r2   )"r   r   r   r   r   r   r   r   r   r   r8   M	num_pid_m	num_pid_nnum_pid_in_groupgroup_idfirst_pid_mgroup_size_mpid_mpid_noffs_amoffs_bnoffs_ka_ptrsb_ptrsaccumulatorkr   r   coffs_cmoffs_cnc_ptrsc_masks"                                     r#   strange_config_matmul_kernelr  T  sx   N mm#GGEN6lR'A!VGGA,	GGA,	'3*-9{2LA 6,FG'L8<'"))A|*DDI<'"))A|*DDI1l+'!T'*VD!G_<=&D/GD!G,<<=hh;2::Nrwwq/0AVD!G_q1|;K7K%KSVWAVAtG_q1|;K7K%KSVWA&&A{3Kl"Fl"F 1 NN2::&,&1l)CC,&1l)CCD))GD!G,<<!T'"Q&747+;a+?@
(r%   c                     [         R                  " SS9n[         R                  " SU5      X2-  -   n[         R                  " U/S[         R                  S9n[         R
                  " X-   XTU:  S9  g)zq
This kernel contains a triple-quote docstring w/ double quotes.
Make sure that codegen sanitizes the docstring.
r   r)         ?r   r+   Nr.   r/   r0   r   rx   r2   r6   numelr&   r8   r:   oness         r#   #kernel_with_docstring_double_quotesr"    sW     mm#))Az*S-==ww
|S

;
"D?r%   c                     [         R                  " SS9n[         R                  " SU5      X2-  -   n[         R                  " U/S[         R                  S9n[         R
                  " X-   XTU:  S9  g)z
This kernel contains a triple-quote docstring w/ single quotes
Make sure that codegen sanitizes the docstring.
To prevent it from being linted to double quotes: """!!!"""
r   r)   r  r   r+   Nr  r  s         r#   #kernel_with_docstring_single_quotesr$    sW     mm#))Az*S-==ww
|S

;
"D?r%   c           	         [         R                  " SS9n[         R                  " SU5      XC-  -   n[         R                  " X-   XR:  S9n[         R                  " SSU/[         R
                  SSS9n[         R                  " X-   XuU:  S9  g )	Nr   r)   r+   z{
            {
                cos.approx.f32 $0, $1;
                ex2.approx.f32 $0, $0;
            }
                =r, rTrU   asmconstraintsargsr   r   r   r.   r/   r0   r1   r   rx   r2   r   r6   r   r&   r8   r:   datacos_pows           r#   kernel_inline_asm_double_quotesr/    s     mm#))Az*S-==wwv'go>++ !**
 	"GE/Br%   c           	         [         R                  " SS9n[         R                  " SU5      XC-  -   n[         R                  " X-   XR:  S9n[         R                  " SSU/[         R
                  SSS9n[         R                  " X-   XuU:  S9  g )	Nr   r)   r+   z
            {
                // double quotes to pacify the linter """!!!"""
                cos.approx.f32 $0, $1;
                ex2.approx.f32 $0, $0;
            }
                r&  TrU   r'  r+  r,  s           r#   kernel_inline_asm_single_quotesr1    s     mm#))Az*S-==wwv'go>++ !**
 	"GE/Br%   c           	         [         R                  " SS9n[         R                  " SU5      XC-  -   n[         R                  " X-   XR:  S9n[         R                  " SSU/[         R
                  SSS9n[         R                  " X-   XuU:  S9  g 	Nr   r)   r+   zK
            v_sin_f32 $0, $1
            v_exp_f32 $0, $0
                z=v, vTrU   r'  r+  r,  s           r#   $kernel_inline_asm_rocm_double_quotesr4         mm#))Az*S-==wwv'go>++ !**

 	"GE/Br%   c           	         [         R                  " SS9n[         R                  " SU5      XC-  -   n[         R                  " X-   XR:  S9n[         R                  " SSU/[         R
                  SSS9n[         R                  " X-   XuU:  S9  g r3  r+  r,  s           r#   $kernel_inline_asm_rocm_single_quotesr7    r5  r%   c                    [         R                  " SS9nXe-  nU[         R                  " SU5      -   nX:  n	[         R                  " X-   U	S9n
U(       a  [         R                  " X-   U	S9nX-   nOU
n[         R                  " X(-   XS9  g r(   r-   )r4   r5   r6   r7   add_xyr&   r8   r9   r:   r,   r;   r<   r=   s                r#   add_kernel_with_boolean_paramr:    s}     mm#&		!Z 88#GGG%D1)5AUFF
"F6r%   block_sizesnew_apic           	      F   U(       a3  [         R                  R                  R                  R	                  X5      $ [        U5      S:X  aZ  [         R                  R                  R                  U R                  5       U R                  S5      US   U R                  5       5      $ [        U5      S:X  d   e[         R                  R                  R                  U R                  5       U R                  S5      U R                  S5      US   US   U R                  5       5      $ )NrU   r   r   )r    toolstensor_descriptorTensorDescriptorfrom_tensorlenexperimental_descriptorcreate_1d_tma_descriptordata_ptrsizeelement_sizecreate_2d_tma_descriptor)tensorr;  r<  s      r#   create_tensor_descriptor_shimrJ  #  s     <<11BBNN  ;1$||;;TTOO%KKNN'')	  ;'1,,,||;;TTOO%KKNKKNNN'') r%   )r&   tl.constexpr)rE   rK  r&   rK  )rR   rK  rS   rK  )r&   rK  r   rK  )r   rK  r   rK  )r   rK  r   rK  )T)Bunittest&torch.testing._internal.inductor_utilsr   r   r   torch.utils._tritonr   
skipUnlessrequires_cuda_and_tritonrequires_gpu_and_tritonrequires_gpur    r   r.   r   listr!   r$   jitr>   rA   rF   rI   autotunerL   rO   rQ   rc   rh   rp   rs   r   r   r   r   r   r   r   r   r   r   r   r   r   	constexprr   r   triton.languager1   r2   r   r   r   r   r   r  r"  r$  r/  r1  r4  r7  r:  intboolrJ  re   r%   r#   <module>rZ     s    
 + $..3  #----/H  ""7N;<<%5$v}}"5 5p ZZ7
 #7 7  ZZ7
 #7 7  ZZ7
 $7 #7 7( ZZ@ $@ #@ @* __MM<-!qIMM<-!qIMM<,aHMM<,aH	
  ZZ7
 #7 7  __MM<-!qIMM<-!qIMM<,aHMM<,aH	
  ZZ7
 #7 7  __MM<,aH
 	 ZZ7 #	7 7$ __MM!$c:qTU MM!$c:qTU MM!#R8QRS MM!#R8QRS
 " ZZJ %J %J #$J, __MM<-!qIMM<,aH
 .0IJ	 ZZ7
 #7 	7  ZZ7 #7 7" ZZ
 #	
 
< ZZ
 %	

 %
 
B ZZ
 #	
 
4 ZZ
 %	

 %
 
: ZZA
 #A
 A
F ZZ4
 #4
 4
l ZZ7 #	7 7 ZZ3 #3 3 ZZ& & ZZ2 #	2
 #2 2$ ZZ3
 %3 %3 3& ZZ-"-+9- -  ZZ-"-+9- -  ZZ+

 LL+
 +
Z ZZ
 LL	
 
B ,ZZ4
 #4 4  ZZ7
 #7 7& ZZ<
 #< <  ZZ;
 #; ;, ZZ7
 #7 7  __MM$&$&$&$%	 	 MM$'$&$&$%	 	
,  /2 ZZ1) ll1) ll1) ll1) ll1) 341)f ZZ@ @ @ ZZ	@ 	@ 	@ ZZC,.LLC C* ZZC,.LLC C, ZZC,.LLC C& ZZC,.LLC C& ZZ7 #7 7, 9=!#Y15]  r%   