ó
    „ñ:i  ã                   ó:   • S SK Jr  S SKJr  S SKrS SKJr  SS jrg)é    )Úcuda)ÚdriverN)Únumpy_supportc           	      óô  ^^• [        U SS5      nU(       dr  U R                  u  p4U R                  R                  U-  U R                  R                  4n[        R
                  R                  R                  XC4UU R                  US9n[        R                  " U R                  5      m[        R                  " 5       R                  n[        [        R                  " S[        R                   " US5      S-  5      5      n[        Xg-  5      nX‡S-   4m[        R"                  UU4S j5       n	[        UR                  S   U-  S-   5      [        UR                  S   U-  S-   5      4n
X‡4nXšX²4   " X5        U$ )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                 óŠ  >• [         R                  R                  T
T	S9n[         R                  R                  n[         R                  R
                  n[         R                  R                  [         R                  R                  -  n[         R                  R
                  [         R                  R
                  -  nXc-   nXT-   nXd-   U R                  S   :  a$  XS-   U R                  S   :  a  XU-   XS-   4   X$U4'   [         R                  " 5         X�R                  S   :  a  XqR                  S   :  a  X#U4   XU4'   g g g )N)Úshaper   r   r
   )
r   ÚsharedÚarrayÚ	threadIdxÚxÚyÚblockIdxÚblockDimr   Úsyncthreads)ÚinputÚoutputÚtileÚtxÚtyÚbxÚbyr   r   ÚdtÚ
tile_shapes            €€Ú_/srv/projetos/modelo_ml_acdoc/venv/lib/python3.13/site-packages/numba/cuda/kernels/transpose.pyÚkernelÚtranspose.<locals>.kernel)   sú   ø€ ô �{‰{× Ñ  z¸Ð Ð<ˆä�^‰^×ÑˆÜ�^‰^×ÑˆÜ�]‰]�_‰_œtŸ}™}Ÿ™Ñ.ˆÜ�]‰]�_‰_œtŸ}™}Ÿ™Ñ.ˆØ‰GˆØ‰Gˆà‰7�U—[‘[ ‘^Ó#¨©°%·+±+¸a±.Ó(@Ø  b¡¨"©'Ð!1Ñ2ˆD�R�‰LÜ×ÒÔØ�|‰|˜A‰Ó 1§|¡|°A¡Ó#6Ø B ™<ˆF�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   Ú	transposer:      sH  ù€ ô �Q˜ !Ó$€FæØ—W‘W‰
ˆØ—'‘'×"Ñ" TÑ)¨1¯7©7×+;Ñ+;Ð;ˆÜ�L‰L×$Ñ$×2Ñ2ØˆLØØ—'‘'Øð	 3ð ˆô 
�Š˜Ÿ™Ó	 €Bä
×
Ò
Ó
×
3Ñ
3€Cä”T—X’X˜a¤§¢¨#¨qÓ!1°AÑ!5Ó6Ó7€JÜ�cÑ&Ó'€Kà¨A™~Ð.€Jä	‡X�Xõ(ó ð(ô$ �—‘˜‘˜kÑ)¨AÑ-Ó.´°A·G±G¸A±JÀÑ4KÈaÑ4OÓ0PÐP€FàÐ%€GØ
�7Ð"Ò# AÔ)à€Hr!   )N)	Únumbar   Únumba.cuda.cudadrv.driverr   r,   Únumba.npr   r'   r:   © r!   r   Ú<module>r?      s   ðÝ Ý ,Û Ý )õ:r!   