ó
    „ñ:iŽ/  ã                   óÄ   • S SK Jr  S SKJrJr  S SKJrJr  SSKJ	r	J
r
JrJr  S SKJr  S SKrS SKrS SKrSrS	 rS
 rS r " S S\R,                  \5      r " S S\5      rg)é    )Úir)ÚconfigÚ	serialize)ÚCodegenÚCodeLibraryé   )ÚdevicesÚdriverÚnvvmÚruntime)Úget_cudalibNznvptx64-nvidia-cudac                 óf  • S nS n [         R                  " 5       u  p#[        US5       nUR                  U 5        S S S 5         [        R
                  " S/UQUPS[        R                  [        R                  S9nUR                  R                  S5      Ub  [        R                  " U5        Ub  [        R                  " U5        $ $ ! , (       d  f       N•= f! [         a  nSn[        U5      UeS nAff = f! Ub  [        R                  " U5        Ub  [        R                  " U5        f f = f)NÚwbÚnvdisasmT)ÚcheckÚstdoutÚstderrztnvdisasm has not been found. You may need to install the CUDA toolkit and ensure that it is available on your PATH.
zutf-8)ÚtempfileÚmkstempÚopenÚwriteÚ
subprocessÚrunÚPIPEÚFileNotFoundErrorÚRuntimeErrorr   ÚdecodeÚosÚcloseÚunlink)ÚcubinÚflagsÚfdÚfnameÚfÚcpÚeÚmsgs           ÚU/srv/projetos/modelo_ml_acdoc/venv/lib/python3.13/site-packages/numba/cuda/codegen.pyÚrun_nvdisasmr*      s  € ð 
€BØ€EðÜ×$Ò$Ó&‰	ˆÜ�%˜Ô !Ø�G‰G�EŒN÷ ð	+Ü—’ Ð ;¨eÐ ;°UÐ ;À4Ü'1§¡Ü'1§¡ñ8ˆBð �y‰y×Ñ Ó(à‰>Ü�HŠH�RŒLØÑÜ�IŠI�eÕð ÷! Õûô !ó 	+ð5ˆCô ˜sÓ#¨Ð*ûð		+ûð ‰>Ü�HŠH�RŒLØÑÜ�IŠI�eÕð úsE   †#C; ©C
»C; Á8C Á<C; Ã

CÃC; Ã
C8Ã%C3Ã3C8Ã8C; Ã;5D0c                 ó   • S/n[        X5      $ )Nz-gi©r*   ©r!   r"   s     r)   Údisassemble_cubinr.   +   s   € àˆG€EÜ˜Ó%Ð%ó    c                 ó   • S/n[        X5      $ )Nz-cfgr,   r-   s     r)   Údisassemble_cubin_for_cfgr1   1   s   € àˆH€EÜ˜Ó%Ð%r/   c                   óâ   ^ • \ rS rSrSr  SU 4S jjr\S 5       rS rS r	SS jr
SS jrSS	 jrS
 rS rSS jrSS jrS rS rS rS r\S 5       r\S 5       rS rS r\S 5       rSrU =r$ )ÚCUDACodeLibraryé7   zÉ
The CUDACodeLibrary generates PTX, SASS, cubins for multiple different
compute capabilities. It also loads cubins to multiple devices (via
get_cufunc), which may be of different compute capabilities.
c                 óþ   >• [         TU ]  X5        SU l        [        5       U l        [        5       U l        SU l        SU l        0 U l        0 U l	        0 U l
        0 U l        0 U l        X@l        Uc  0 nXPl        X0l        g)a?  
codegen:
    Codegen object.
name:
    Name of the function in the source.
entry_name:
    Name of the kernel function in the binary, if this is a global
    kernel and not a device function.
max_registers:
    The maximum register usage to aim for when linking.
nvvm_options:
        Dict of options to pass to NVVM.
NF)ÚsuperÚ__init__Ú_moduleÚsetÚ_linking_librariesÚ_linking_filesÚneeds_cudadevrtÚ
_llvm_strsÚ
_ptx_cacheÚ_ltoir_cacheÚ_cubin_cacheÚ_linkerinfo_cacheÚ_cufunc_cacheÚ_max_registersÚ_nvvm_optionsÚ_entry_name)ÚselfÚcodegenÚnameÚ
entry_nameÚmax_registersÚnvvm_optionsÚ	__class__s         €r)   r7   ÚCUDACodeLibrary.__init__>   sˆ   ø€ ô 	‰Ñ˜Ô'ð ˆŒô
 #&£%ˆÔô "›eˆÔà$ˆÔð ˆŒàˆŒàˆÔàˆÔà!#ˆÔàˆÔà+ÔØÑØˆLØ)ÔØ%Õr/   c                 ó”   • U R                   c+  U R                   Vs/ s H  n[        U5      PM     snU l         U R                   $ s  snf ©N)r=   ÚmodulesÚstr©rF   Úmods     r)   Ú	llvm_strsÚCUDACodeLibrary.llvm_strso   s:   € à�?‰?Ñ"Ø37·<²<Ó@²<¨Cœs 3žx±<Ñ@ˆDŒOØ�‰Ðùò As   œAc                 ó8   • SR                  U R                  5      $ )Nz

)ÚjoinrT   )rF   s    r)   Úget_llvm_strÚCUDACodeLibrary.get_llvm_stru   s   € Ø�{‰{˜4Ÿ>™>Ó*Ð*r/   c                 ób   • Ub  U$ [         R                  " 5       R                  nUR                  $ rO   )r	   Úget_contextÚdeviceÚcompute_capability)rF   Úccr\   s      r)   Ú
_ensure_ccÚCUDACodeLibrary._ensure_ccx   s,   € Ø‰>ØˆIä×$Ò$Ó&×-Ñ-ˆØ×(Ñ(Ð(r/   c                 ó6  • U R                  U5      nU R                  R                  US 5      nU(       a  U$ [        R                  " U6 nU R
                  R                  5       nX4S'   U R                  n[        R                  " U40 UD6nUR                  5       R                  S5      R                  5       n[        R                  (       a>  [        SU R                  -  R                  SS5      5        [        U5        [        S5        X`R                  U'   U$ )NÚarchÚ zASSEMBLY %séP   Ú-zP================================================================================)r_   r>   Úgetr   Úget_arch_optionrD   ÚcopyrT   Ú
compile_irr   Ústripr   ÚDUMP_ASSEMBLYÚprintÚ_nameÚcenter)rF   r^   Úptxesrb   ÚoptionsÚirsÚptxs          r)   Úget_asm_strÚCUDACodeLibrary.get_asm_str   sÜ   € Ø�_‰_˜RÓ ˆà—‘×#Ñ# B¨Ó-ˆÞØˆLä×#Ò# RÐ(ˆØ×$Ñ$×)Ñ)Ó+ˆØ�‰à�n‰nˆä�oŠo˜cÑ- WÑ-ˆð
 �j‰j‹l× Ñ  Ó(×.Ñ.Ó0ˆä××Ü�= 4§:¡:Ñ-×5Ñ5°b¸#Ó>Ô?Ü�#ŒJÜ�(ŒOà!�‰˜Ñàˆ
r/   c                 ó8  • U R                  U5      nU R                  R                  US 5      nUb  U$ [        R                  " U6 nU R
                  R                  5       nX4S'   S US'   U R                  n[        R                  " U40 UD6nX R                  U'   U$ )Nrb   zgen-lto)	r_   r?   rf   r   rg   rD   rh   rT   ri   )rF   r^   Últoirrb   rp   rq   s         r)   Ú	get_ltoirÚCUDACodeLibrary.get_ltoirœ   s—   € Ø�_‰_˜RÓ ˆà×!Ñ!×%Ñ% b¨$Ó/ˆØÑØˆLä×#Ò# RÐ(ˆØ×$Ñ$×)Ñ)Ó+ˆØ�‰Ø!ˆ�	Ñà�n‰nˆÜ—’ Ñ/ wÑ/ˆØ %×Ñ˜"Ñàˆr/   c                 óŒ  • U R                  U5      nU R                  R                  US 5      nU(       a  U$ [        R                  R                  U R                  US9nUR                  (       a!  U R                  US9nUR                  U5        O.U R                  US9nUR                  UR                  5       5        U R                   H  nUR                  U5        M     U R                  (       a  UR                  [!        SSS95        UR#                  5       nX R                  U'   UR$                  U R&                  U'   U$ )N)rJ   r^   ©r^   Ú	cudadevrtT)Ústatic)r_   r@   rf   r
   ÚLinkerÚnewrC   Últorw   Ú	add_ltoirrs   Úadd_ptxÚencoder;   Úadd_file_guess_extr<   r   ÚcompleteÚinfo_logrA   )rF   r^   r!   Úlinkerrv   rr   Úpaths          r)   Ú	get_cubinÚCUDACodeLibrary.get_cubin®   s  € Ø�_‰_˜RÓ ˆà×!Ñ!×%Ñ% b¨$Ó/ˆÞØˆLä—‘×"Ñ"°×1DÑ1DÈÐ"ÐLˆà�:�:Ø—N‘N b�NÐ)ˆEØ×Ñ˜UÕ#à×"Ñ" bÐ"Ð)ˆCØ�N‰N˜3Ÿ:™:›<Ô(à×'Ô'ˆDØ×%Ñ% dÖ+ñ (à××Ø×%Ñ%¤k°+ÀdÑ&KÔLà—‘Ó!ˆØ %×Ñ˜"ÑØ%+§_¡_ˆ×Ñ˜rÑ"àˆr/   c                 ó”  • U R                   c  Sn[        U5      e[        R                  " 5       nUR                  nU R
                  R                  UR                  S 5      nU(       a  U$ U R                  UR                  S9nUR                  U5      nUR                  U R                   5      nX@R
                  UR                  '   U$ )NzLMissing entry_name - are you trying to get the cufunc for a device function?rz   )rE   r   r	   r[   r\   rB   rf   Úidrˆ   r]   Úcreate_module_imageÚget_function)rF   r(   Úctxr\   Úcufuncr!   Úmodules          r)   Ú
get_cufuncÚCUDACodeLibrary.get_cufuncÉ   s²   € Ø×ÑÑ#ð+ˆCä˜sÓ#Ð#ä×!Ò!Ó#ˆØ—‘ˆà×#Ñ#×'Ñ'¨¯	©	°4Ó8ˆÞØˆMà—‘ &×";Ñ";�Ð<ˆØ×(Ñ(¨Ó/ˆð ×$Ñ$ T×%5Ñ%5Ó6ˆð )/×Ñ˜6Ÿ9™9Ñ%àˆr/   c                 óZ   •  U R                   U   $ ! [         a    [        SU 35      ef = f)NzNo linkerinfo for CC )rA   ÚKeyError©rF   r^   s     r)   Úget_linkerinfoÚCUDACodeLibrary.get_linkerinfoá   s:   € ð	9Ø×)Ñ)¨"Ñ-Ð-øÜó 	9ÜÐ2°2°$Ð7Ó8Ð8ð	9ús   ‚ ‘*c                 ó2   • [        U R                  US95      $ ©Nrz   )r.   rˆ   r•   s     r)   Úget_sassÚCUDACodeLibrary.get_sassç   s   € Ü  §¡°2 Ð!6Ó7Ð7r/   c                 ó2   • [        U R                  US95      $ r™   )r1   rˆ   r•   s     r)   Úget_sass_cfgÚCUDACodeLibrary.get_sass_cfgê   s   € Ü(¨¯©¸2¨Ð)>Ó?Ð?r/   c                 ó`   • U R                  5         U R                  b  [        S5      eXl        g )Nz(CUDACodeLibrary only supports one module)Ú_raise_if_finalizedr8   r   rR   s     r)   Úadd_ir_moduleÚCUDACodeLibrary.add_ir_moduleí   s)   € Ø× Ñ Ô"Ø�<‰<Ñ#ÜÐIÓJÐJØ�r/   c                 óz   • UR                  5         U R                  5         U R                  R                  U5        g rO   )Ú_ensure_finalizedr    r:   Úadd)rF   Úlibrarys     r)   Úadd_linking_libraryÚ#CUDACodeLibrary.add_linking_libraryó   s0   € Ø×!Ñ!Ô#ð
 	× Ñ Ô"à×Ñ×#Ñ# GÕ,r/   c                 ó:   • U R                   R                  U5        g rO   )r;   r¥   )rF   Úfilepaths     r)   Úadd_linking_fileÚ CUDACodeLibrary.add_linking_fileý   s   € Ø×Ñ×Ñ Õ)r/   c                 ó‚   • U R                   R                   H  nUR                  U:X  d  M  Us  $    [        SU S35      e)Nz	Function z
 not found)r8   Ú	functionsrH   r”   )rF   rH   Úfns      r)   r�   ÚCUDACodeLibrary.get_function   s<   € Ø—,‘,×(Ô(ˆBØ�w‰w˜$�Ø’	ñ )ô ˜ 4 &¨
Ð3Ó4Ð4r/   c                 ó�   • U R                   /U R                   VVs/ s H  nUR                    H  nUPM     M     snn-   $ s  snnf rO   )r8   r:   rP   )rF   ÚlibrS   s      r)   rP   ÚCUDACodeLibrary.modules  sK   € à—‘ˆ~°×0GÒ0Gô !9Ò0G¨Ø,/¯K­K Só "%Ù,7ñ "%Ñ0Gò !9ñ 9ð 	9ùó !9s   œ Ac                 óˆ   • / nU R                    H/  nUR                  UR                  5        UR                  U5        M1     U$ rO   )r:   ÚextendÚlinking_librariesÚappend)rF   Úlibsr²   s      r)   r¶   Ú!CUDACodeLibrary.linking_libraries  s>   € ð
 ˆØ×*Ô*ˆCØ�K‰K˜×-Ñ-Ô.Ø�K‰K˜Öñ +ð ˆr/   c                 óØ   • U R                  5         U R                   HC  nUR                   H0  nUR                   H  nUR                  (       a  M  SUl        M     M2     ME     SU l        g )NÚlinkonce_odrT)r    r:   rP   r®   Úis_declarationÚlinkageÚ
_finalized)rF   r¦   rS   r¯   s       r)   ÚfinalizeÚCUDACodeLibrary.finalize  sZ   € ð 	× Ñ Ô"ð ×.Ô.ˆGØ—”�ØŸ-œ-�BØ×,×,Ñ,Ø%3˜ž
ó (ó 'ñ /ð ˆ�r/   c                 óN  • U R                   (       a  Sn[        U5      eU R                  (       d  [        S5      e[        SU R                  U R
                  U R                  U R                  U R                  U R                  U R                  U R                  U R                  S9
$ )z 
Reduce the instance for serialization. We retain the PTX and cubins,
but loaded functions are discarded. They are recreated when needed
after deserialization.
z0Cannot pickle CUDACodeLibrary with linking filesz)Cannot pickle unfinalized CUDACodeLibraryN)
rG   rH   rI   rT   Ú	ptx_cacheÚcubin_cacheÚlinkerinfo_cacherJ   rK   r<   )r;   r   r¾   ÚdictrH   rE   rT   r>   r@   rA   rC   rD   r<   )rF   r(   s     r)   Ú_reduce_statesÚCUDACodeLibrary._reduce_states0  s‹   € ð ××ØDˆCÜ˜sÓ#Ð#Ø��ÜÐJÓKÐKÜØØ—‘Ø×'Ñ'Ø—n‘nØ—o‘oØ×)Ñ)Ø!×3Ñ3Ø×-Ñ-Ø×+Ñ+Ø ×0Ñ0ñ
ð 	
r/   c                 óv   • U " XUS9nXKl         X[l        Xkl        X{l        X‹l        X›l        X«l        SUl        U$ )z
Rebuild an instance.
)rI   T)r=   r>   r@   rA   rC   rD   r<   r¾   )ÚclsrG   rH   rI   rT   rÂ   rÃ   rÄ   rJ   rK   r<   Úinstances               r)   Ú_rebuildÚCUDACodeLibrary._rebuildH  sG   € ñ �w°Ñ<ˆà'ÔØ'ÔØ +ÔØ%5Ô"à"/ÔØ!-ÔØ#2Ô à"ˆÔàˆr/   )r@   rB   rE   r¾   rA   r;   r:   r=   r?   rC   r8   rD   r>   r<   )NNNrO   )Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__r7   ÚpropertyrT   rX   r_   rs   rw   rˆ   r‘   r–   rš   r�   r¡   r§   r«   r�   rP   r¶   r¿   rÆ   ÚclassmethodrË   Ú__static_attributes__Ú__classcell__)rL   s   @r)   r3   r3   7   s´   ø† ñð FJØ"÷/&ðb ñó ðò
+ò)ôô:ô$ò6ò09ô8ô@òò-ò*ò5ð ñ9ó ð9ð ñó ðòò4
ð0 ñó ör/   r3   c                   ó4   • \ rS rSrSr\rS rS rS r	S r
Srg)	ÚJITCUDACodegeni_  z�
This codegen implementation for CUDA only generates optimized LLVM IR.
Generation of PTX code is done separately (see numba.cuda.compiler).
c                 ó   • g rO   © )rF   Úmodule_names     r)   r7   ÚJITCUDACodegen.__init__g  ó   € Ør/   c                 ó¼   • [         R                  " U5      n[        Ul        [        R
                  " 5       R                  Ul        [        R                  " U5        U$ rO   )r   ÚModuleÚCUDA_TRIPLEÚtripler   ÚNVVMÚdata_layoutÚadd_ir_version)rF   rH   Ú	ir_modules      r)   Ú_create_empty_moduleÚ#JITCUDACodegen._create_empty_modulej  s?   € Ü—I’I˜d“Oˆ	Ü&ˆ	ÔÜ $§	¢	£× 7Ñ 7ˆ	ÔÜ×Ò˜IÔ&ØÐr/   c                 ó   • g rO   rÙ   )rF   r�   s     r)   Ú_add_moduleÚJITCUDACodegen._add_moduleq  rÜ   r/   c                 ó˜   • [         R                  " 5       nUR                  R                  n[        R                  R                  5       U4$ )z@
Return a tuple unambiguously describing the codegen behaviour.
)r	   r[   r\   r]   r   Úget_version)rF   rŽ   r^   s      r)   Úmagic_tupleÚJITCUDACodegen.magic_tuplet  s9   € ô ×!Ò!Ó#ˆØ�Z‰Z×*Ñ*ˆÜ—‘×+Ñ+Ó-¨rÐ2Ð2r/   rÙ   N)rÍ   rÎ   rÏ   rÐ   rÑ   r3   Ú_library_classr7   rå   rè   rì   rÔ   rÙ   r/   r)   r×   r×   _  s"   † ñð
 %€Nòòòõ3r/   r×   )Úllvmliter   Ú
numba.corer   r   Únumba.core.codegenr   r   Úcudadrvr	   r
   r   r   Únumba.cuda.cudadrv.libsr   r   r   r   rß   r*   r.   r1   ÚReduceMixinr3   r×   rÙ   r/   r)   Ú<module>rõ      sY   ðÝ ç (ß 3ß 3Ó 3Ý /ã 	Û Û ð $€òò6&ò&ôe�i×+Ñ+¨[ô eôP	3�Wõ 3r/   