
    PhZ                        d dl Z d dlZd dlmZ  e j                  ej
                  e d      Zer>d dlZd dlmZ	 ej                  	 	 dd       Zej                  	 	 	 	 dd       Z ej                   ej                  ddid	d
       ej                  ddid	d
      gg       ej                  	 	 dd              Z ej                   ej                  dddd	d
       ej                  dddd	d
      gg       ej                  	 	 	 	 dd              Zej                  	 	 dd       Zej                  	 	 dd       Zej                  d        Zej                  	 	 	 	 dd       Zyy)    N)HAS_CUDAzrequires cuda)language
BLOCK_SIZEc                    t        j                  d      }||z  }|t        j                  d|      z   }||k  }t        j                  | |z   |      }	t        j                  ||z   |      }
|	|
z   }t        j                  ||z   ||       y Nr   axismasktl
program_idarangeloadstorein_ptr0in_ptr1out_ptr
n_elementsr   pidblock_startoffsetsr   xyoutputs               oC:\Users\daisl\Desktop\realtime-object-detection\venv\Lib\site-packages\torch/testing/_internal/triton_utils.py
add_kernelr      s     mm#J&		!Z 88#GGGg%D1GGGg%D1Q
7"F6    c                 .   t        j                  d      }||z  }|t        j                  d|      z   }||k  }	t        j                  | |z   |	      }
|dk(  r t        j                  ||z   |	      }|
|z   }n|
}t        j                  ||z   ||	       y )Nr   r   r
   twor   )r   r   r   r   ARGS_PASSEDr   r   r   r   r   r   r   r   s                r   add_kernel_with_optional_paramr#      s     mm#J&		!Z 88#GGGg%D1%')5AUFF
7"F6r            )
num_stages	num_warps@   )configskeyc                    t        j                  d      }||z  }|t        j                  d|      z   }||k  }t        j                  | |z   |      }	t        j                  ||z   |      }
|	|
z   }t        j                  ||z   ||       y r   r   r   s               r   add_kernel_autotunedr-   3   s     mm#J&		!Z 88#GGGg%D1GGGg%D1Q
7"F6r   )BLOCK_SIZE_XBLOCK_SIZE_Yc                    t        j                  d      |z  }|t        j                  d|      d d d f   z   }||k  }	t        j                  d      |z  }
|
t        j                  d|      d d d f   z   }||k  }|}|}t        j                  | |||z  z   z   |	|z        }t        j                  | |||z  z   z   |	|z        }||z   }t        j                  ||||z  z   z   ||	|z         y )Nr      r   )r   r   r   
x_elements
y_elementsr.   r/   xoffsetxindexxmaskyoffsetyindexymaskx1y0tmp0tmp1tmp2s                     r   add_kernel_2d_autotunedr?   K   s    * --"\1299Q5ag>>#--"\1299Q5dAg>>#www"
R"8955=Iwww"
R"8955=Id{
B*r/23T55=Ir   c                     t        j                  d      }||z  }|t        j                  d|      z   }||k  }t        j                  | |z   |      }d|z  }	t        j                  ||z   |	|       y Nr   r   r
      r   )
r   r   r   r   r   r   r   r   r   r   s
             r   mul2_kernelrC   m   sn     mm#J&		!Z 88#GGGg%D1Q
7"F6r   c                     t        j                  d      }||z  }|t        j                  d|      z   }||k  }t        j                  | |z   |      }d|z  }t        j                  | |z   ||       y rA   r   )	ptrr   r   r   r   r   r   r   r   s	            r   mul2_inplace_kernelrF   |   sl     mm#J&		!Z 88#GGC'M-Q
wT2r   c                 6    t        j                  | dk\  | d      S )Nr   )r   where)r   s    r   	zero_negsrI      s    xxQ1%%r   c                    t        j                  d      }||z  }|t        j                  d|      z   }||k  }|dk(  rt        | ||       t        j                  | |z   |      }	t        j
                  ||z   |	|       y )Nr   r   rF   )r   r
   )r   r   r   rF   r   r   )
r   r   r   r   
ACTIVATIONr   r   r   r   r   s
             r   indirection_kernelrL      sz     mm#J&		!Z 88#..
KGGGg%D1
7"AD1r   )r   tl.constexpr)r"   rM   r   rM   )r.   rM   r/   rM   )r   rM   rK   rM   )	functoolsunittest&torch.testing._internal.inductor_utilsr   partialskipIfrequires_cudatritonr   r   jitr   r#   autotuneConfigr-   r?   rC   rF   rI   rL    r   r   <module>rY      s     ;!	!!(//x<Q% ZZ7
 #7 7  ZZ7
 $7 #7 7( V__FMM<-!qIFMM<,aH
  ZZ7
 #7 7  V__FMM!$c:qTU FMM!#R8QRS	
 
 ZZJ %J %J 
J, ZZ7 #	7 7 ZZ3 #3 3 ZZ& & ZZ2 #	2
 #2 2O r   