Ë
    èÿæi  ã                   ó8   — d dl mZ d dlmZ d dlZd dlmZ dd„Zy)é    )Úcuda)ÚdriverN)Únumpy_supportc           	      óö  ‡‡— t        | dd«      }|sw| j                  \  }}| j                  j                  |z  | j                  j                  f}t        j
                  j                  j                  ||f|| j                  |¬«      }t        j                  | j                  «      Št        j                  «       j                  }t        t        j                  dt        j                   |d«      dz  «      «      }t        ||z  «      }||dz   fŠt        j"                  ˆˆfd„«       }	t        |j                  d   |z  dz   «      t        |j                  d   |z  dz   «      f}
||f} |	|
||f   | |«       |S )aá  Compute the transpose of 'a' and store it into 'b', if given,
    and return it. If 'b' is not given, allocate a new array
    and return that.

    This implements the algorithm documented in
    http://devblogs.nvidia.com/parallelforall/efficient-matrix-transpose-cuda-cc/

    :param a: an `np.ndarray` or a `DeviceNDArrayBase` subclass. If already on
        the device its stream will be used to perform the transpose (and to copy
        `b` to the device if necessary).
    Ústreamr   )Údtyper   é   é   c                 óœ  •— t         j                  j                  ‰
‰	¬«      }t         j                  j                  }t         j                  j
                  }t         j                  j                  t         j                  j                  z  }t         j                  j
                  t         j                  j
                  z  }||z   }||z   }||z   | j                  d   k  r'||z   | j                  d   k  r| ||z   ||z   f   |||f<   t        j                  «        ||j                  d   k  r ||j                  d   k  r|||f   |||f<   y y y )N)Úshaper   r   r
   )
r   ÚsharedÚarrayÚ	threadIdxÚxÚyÚblockIdxÚblockDimr   Úsyncthreads)ÚinputÚoutputÚtileÚtxÚtyÚbxÚbyr   r   ÚdtÚ
tile_shapes            €€úq/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/numba/cuda/kernels/transpose.pyÚkernelztranspose.<locals>.kernel)   s  ø€ ô �{‰{× Ñ  z¸Ð Ó<ˆä�^‰^×ÑˆÜ�^‰^×ÑˆÜ�]‰]�_‰_œtŸ}™}Ÿ™Ñ.ˆÜ�]‰]�_‰_œtŸ}™}Ÿ™Ñ.ˆØ�‰GˆØ�‰Gˆà�‰7�U—[‘[ ‘^Ò#¨¨R©°%·+±+¸a±.Ò(@Ø   b¡¨"¨r©'Ð!1Ñ2ˆD��R�‰LÜ×ÑÔØˆv�|‰|˜A‰Ò 1 v§|¡|°A¡Ò#6Ø  B ™<ˆF�1�a�4ŠLð $7Ðó    )Úgetattrr   r   Úitemsizer   ÚcudadrvÚdevicearrayÚDeviceNDArrayÚnpsÚ
from_dtyper   Ú
get_deviceÚMAX_THREADS_PER_BLOCKÚintÚmathÚpowÚlogÚjit)ÚaÚbr   ÚcolsÚrowsÚstridesÚtpbÚ
tile_widthÚtile_heightr   ÚblocksÚthreadsr   r   s               @@r   Ú	transposer9      sZ  ù€ ô �Q˜ !Ó$€FáØ—W‘W‰
ˆˆdØ—'‘'×"Ñ" TÑ)¨1¯7©7×+;Ñ+;Ð;ˆÜ�L‰L×$Ñ$×2Ñ2Ø�4ˆLØØ—'‘'Øð	 3ó ˆô 
�‰˜Ÿ™Ó	 €Bä
×
Ñ
Ó
×
3Ñ
3€Cä”T—X‘X˜a¤§¡¨#¨qÓ!1°AÑ!5Ó6Ó7€JÜ�c˜JÑ&Ó'€Kà˜z¨A™~Ð.€Jä	‡X�Xô(ó ð(ô$ �—‘˜‘˜kÑ)¨AÑ-Ó.´°A·G±G¸A±JÀÑ4KÈaÑ4OÓ0PÐP€Fà˜:Ð%€GØ#€Fˆ6�7˜FÐ"Ñ# A qÔ)à€Hr    )N)	Únumbar   Únumba.cuda.cudadrv.driverr   r+   Únumba.npr   r&   r9   © r    r   Ú<module>r>      s   ðÝ Ý ,Û Ý )ô:r    