U
    hâËdf;  ã                   @   s  d dl Z d dlmZ d dlmZ d dlmZ d dlm	Z	m
Z
mZmZmZ d dlmZ d dlmZ d dlmZ d dlmZ d d	lmZ d d
lmZ ddlmZ d dlmZmZmZ d dlm Z  d dl!m"Z" G dd„ de	jƒZ#e  $de j%¡Z&G dd„ deƒZ'G dd„ deƒZ(dS )é    N)Úcached_property)Úir)ÚtypingÚtypesÚ	debuginfoÚitanium_manglerÚcgutils)Ú
Dispatcher)ÚNumbaInvalidConfigWarning)ÚBaseContext)ÚMinimalCallConv)Ú	cmathdecl)Ú	datamodelé   )Únvvm)ÚcodegenÚ	nvvmutilsÚufuncs)Úcuda_data_manager)Úwarnc                       s$   e Zd Zdd„ Z‡ fdd„Z‡  ZS )ÚCUDATypingContextc                 C   s€   ddl m}m}m}m} ddlm}m} |  |j	¡ |  |j	¡ |  |j	¡ |  t
j	¡ |  |j	¡ |  |j	¡ |  |j¡ d S )Nr   )ÚcudadeclÚcudamathÚlibdevicedeclÚvector_typesr   )ÚenumdeclÚ
cffi_utils)Ú r   r   r   r   Únumba.core.typingr   r   Úinstall_registryÚregistryr   Ztyping_registry)Úselfr   r   r   r   r   r   © r"   úJ/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/target.pyÚload_additional_registries   s    z,CUDATypingContext.load_additional_registriesc                    s¦   ddl m} t|tƒr–t||ƒs–z
|j}W nj tk
r”   |jsHtdƒ‚|j 	¡ }d|d< | 
dd¡|d< | 
dd¡|d< ||j|ƒ}||_|}Y nX tt| ƒ |¡S )	Nr   )ÚCUDADispatcherz<using cpu function on device but its compilation is disabledTZdeviceÚdebugFÚopt)Znumba.cuda.dispatcherr%   Ú
isinstancer	   Z_CUDATypingContext__dispatcherÚAttributeErrorZ_can_compileÚ
ValueErrorÚtargetoptionsÚcopyÚgetZpy_funcÚsuperr   Úresolve_value_type)r!   Úvalr%   r+   Zdisp©Ú	__class__r"   r#   r/   %   s"    
ÿ


z$CUDATypingContext.resolve_value_type)Ú__name__Ú
__module__Ú__qualname__r$   r/   Ú__classcell__r"   r"   r1   r#   r      s   r   z	[^a-z0-9]c                       sÌ   e Zd ZdZdZd*‡ fdd„	Zedd„ ƒZedd„ ƒZd	d
„ Z	dd„ Z
dd„ Zdd„ Zedd„ ƒZedd„ ƒZedd„ ƒZdddœdd„Zd+dd„Zdd„ Zd d!„ Zd"d#„ Zd$d%„ Zd&d'„ Zd(d)„ Z‡  ZS ),ÚCUDATargetContextTÚcudac                    s    t ƒ  ||¡ t tj¡| _d S ©N)r.   Ú__init__r   Úchainr   Zdefault_managerZdata_model_manager)r!   Z	typingctxÚtargetr1   r"   r#   r:   I   s    ÿzCUDATargetContext.__init__c                 C   s*   t  ¡ jrtjS d}tt|ƒƒ tjS d S )Nz3debuginfo is not generated for CUDA toolkits < 11.2)r   ÚNVVMÚ	is_nvvm70r   Ú	DIBuilderr   r
   ZDummyDIBuilder)r!   Úmsgr"   r"   r#   r?   O   s
    
zCUDATargetContext.DIBuilderc                 C   s   dS )NFr"   ©r!   r"   r"   r#   Úenable_boundscheckX   s    z$CUDATargetContext.enable_boundscheckc                 C   s   | j  |¡S r9   )Ú_internal_codegenZ_create_empty_module)r!   Únamer"   r"   r#   Úcreate_module^   s    zCUDATargetContext.create_modulec                 C   s   t  d¡| _d | _d S )Nznumba.cuda.jit)r   ZJITCUDACodegenrC   Ú_target_datarA   r"   r"   r#   Úinita   s    zCUDATargetContext.initc                 C   sè   ddl m}m}m} ddl m}m}m} ddl m}m} ddl m	}	 ddl
m}
 ddlm} ddlm} d	d
lm}m}m}m}m} ddlm} |  |j¡ |  |
j¡ |  |j¡ |  |j¡ |  |	j¡ |  |j¡ |  |j¡ d S )Nr   )ÚnumbersÚtupleobjÚslicing)ÚrangeobjÚ	iteratorsÚenumimpl)ÚunicodeÚcharseq)Ú	cmathimpl)Úcffiimpl)Úarrayobj)Ú
npdatetimer   )ÚcudaimplÚ	printimplÚlibdeviceimplÚmathimplr   )Úndarray)Znumba.cpythonrH   rI   rJ   rK   rL   rM   rN   rO   rP   Z
numba.miscrQ   Znumba.nprR   rS   r   rT   rU   rV   rW   r   Znumba.np.unsaferX   r   r    Zimpl_registry)r!   rH   rI   rJ   rK   rL   rM   rN   rO   rP   rQ   rR   rS   rT   rU   rV   rW   r   rX   r"   r"   r#   r$   e   s     z,CUDATargetContext.load_additional_registriesc                 C   s   | j S r9   )rC   rA   r"   r"   r#   r   }   s    zCUDATargetContext.codegenc                 C   s"   | j d krt t ¡ j¡| _ | j S r9   )rF   ÚllZcreate_target_datar   r=   Zdata_layoutrA   r"   r"   r#   Útarget_data€   s    
zCUDATargetContext.target_datac                    s*   ddl m‰  d}t‡ fdd„|D ƒƒ}|S )z×
        Some CUDA intrinsics are at the module level, but cannot be treated as
        constants, because they are loaded from a special register in the PTX.
        These include threadIdx, blockDim, etc.
        r   ©r8   )Z	threadIdxZblockDimZblockIdxZgridDimZlaneidZwarpsizec                    s   g | ]}t  ˆ ¡|f‘qS r"   )r   ÚModule)Ú.0Úncr[   r"   r#   Ú
<listcomp>�   s   ÿz;CUDATargetContext.nonconst_module_attrs.<locals>.<listcomp>)Znumbar8   Útuple)r!   Z	nonconstsZnonconsts_with_modr"   r[   r#   Únonconst_module_attrs†   s    ÿz'CUDATargetContext.nonconst_module_attrsc                 C   s   t | ƒS r9   )ÚCUDACallConvrA   r"   r"   r#   Ú	call_conv”   s    zCUDATargetContext.call_convr"   N©Úabi_tagsÚuidc                C   s   t j||||d�S )Nrd   )r   Zmangle)r!   rD   Úargtypesre   rf   r"   r"   r#   Úmangler˜   s    
ÿzCUDATargetContext.manglerc	              	   C   sV   t j|jdd�}	|  ¡ j|j› d�|	||d�}
|
 |¡ |  |
||	||||¡}|
|fS )a»  
        Adapt a code library ``codelib`` with the numba compiled CUDA kernel
        with name ``fname`` and arguments ``argtypes`` for NVVM.
        A new library is created with a wrapper function that can be used as
        the kernel entry point for the given kernel.

        Returns the new code library and the wrapper function.

        Parameters:

        codelib:       The CodeLibrary containing the device function to wrap
                       in a kernel call.
        fndesc:        The FunctionDescriptor of the source function.
        debug:         Whether to compile with debug.
        lineinfo:      Whether to emit line info.
        nvvm_options:  Dict of NVVM options used when compiling the new library.
        filename:      The source filename that the function is contained in.
        linenum:       The source line that the function is on.
        max_registers: The max_registers argument for the code library.
        Úcudapy©ÚnsZ_kernel_)Z
entry_nameÚnvvm_optionsÚmax_registers)r   Úprepend_namespaceÚllvm_func_namer   Zcreate_libraryrD   Zadd_linking_libraryÚgenerate_kernel_wrapper)r!   ZcodelibÚfndescr&   Úlineinforl   ÚfilenameÚlinenumrm   Úkernel_nameÚlibraryÚwrapperr"   r"   r#   Úprepare_cuda_kernelœ   s"     ÿý

  þz%CUDATargetContext.prepare_cuda_kernelc           $   
      sâ  |j }|  |¡}	t|	jƒ}
t t ¡ |
¡}|  d¡‰t t d¡| j	 
tj¡g|
 ¡}t ˆ||j¡}tj|jdd�}t ˆ||¡‰ t ˆ  d¡¡}|sœ|rØ|o¤| }| jˆ|| |d�}| ˆ ||j||¡ | ||¡ ‡ ‡fdd„}|d	ƒ}g }g }d
D ](}| |d| ƒ¡ | |d| ƒ¡ qú|	 |ˆ j¡}| j	 ||tj||¡\}}|�ršt ||j¡� | ¡  W 5 Q R X |  | !|j"¡¡�� t #|j$j%d¡}t& '¡ j(�rÈ| )|||j*dd¡}| +|d¡}nPt |j$|j$|j$|j$g¡}d}tjˆ||d�}| ,||||j*g¡}| -d||¡}t. /|¡} |  |¡�^ t0d
|ƒD ] \}!}"|  1|!¡}#| 2|#|"¡ �q8t0d
|ƒD ] \}!}"|  3|!¡}#| 2|#|"¡ �qdW 5 Q R X W 5 Q R X | ¡  t& 4ˆ ¡ | 5ˆ¡ |�sÂ|�rÊ| 6¡  | 6¡  | 7ˆ j¡‰ ˆ S )z°
        Generate the kernel wrapper in the given ``library``.
        The function being wrapped is described by ``fndesc``.
        The wrapper function is returned.
        zcuda.kernel.wrapperé    ri   rj   r   )ÚmoduleÚfilepathZcgctxÚdirectives_onlyc                    s4   ˆ j |  }t ˆt d¡|¡}t |jjd ¡|_|S )Nry   )	rD   r   Úadd_global_variabler   ÚIntTypeÚConstantÚtypeÚpointeeÚinitializer)ZpostfixrD   Úgv©ZwrapfnZwrapper_moduler"   r#   Údefine_error_gvâ   s    
ÿzBCUDATargetContext.generate_kernel_wrapper.<locals>.define_error_gvZ__errcode__Zxyzz	__tid%s__z__ctaid%s__NÚ	monotonicr   Z___numba_atomic_i32_cas_hack)rD   z==)8rg   Zget_arg_packerÚlistZargument_typesr   ÚFunctionTypeZVoidTyperE   r~   rc   Zget_return_typer   ZpyobjectÚFunctionro   r   rn   rD   Z	IRBuilderZappend_basic_blockr?   Zmark_subprogramÚargsZmark_locationÚappendZfrom_argumentsZcall_functionÚvoidr   Z	if_likelyZis_okZret_voidZif_thenÚnot_Zis_python_excr   r€   r�   r   r=   r>   ZcmpxchgÚcodeZextract_valueÚcallZicmp_unsignedr   ZSRegBuilderÚzipÚtidÚstoreZctaidZset_cuda_kernelZadd_ir_moduleÚfinalizeZget_function)$r!   rv   rq   ru   r&   rr   rs   rt   rg   ZarginfoZargtysZwrapfntyZfntyÚfuncÚprefixedÚbuilderr|   r   r…   Zgv_excZgv_tidZgv_ctaidÚiZcallargsÚstatusÚ_ÚoldZxchgÚchangedZcasfntyZcas_hackZcasfnZsregÚdimZptrr0   r"   r„   r#   rp   À   sœ    


ÿÿ
ý    ÿ    ÿ ÿÿ


$

z)CUDATargetContext.generate_kernel_wrapperc              	      s   |j }‡ fdd„t|jdd�ƒD ƒ}t t d¡t|ƒ¡}t ||¡}tj	}t
j||jd|d�}	d|	_d	|	_||	_ˆ  |j¡}
ˆ  |
¡}d
|d  ¡  |	_t t d¡¡}| |	|d¡}ˆ  |¡ˆ |ƒ}‡ fdd„|jD ƒ}‡ fdd„|jD ƒ}ˆ j|| ||jj¡|||j|jdd� | ¡ S )úi
        Unlike the parent version.  This returns a a pointer in the constant
        addrspace.
        c                    s   g | ]}ˆ   tj|¡‘qS r"   )Úget_constantr   Úbyte)r]   r—   rA   r"   r#   r_   *  s   ÿz9CUDATargetContext.make_constant_array.<locals>.<listcomp>ÚA)Úorderé   Z_cudapy_cmem©Ú	addrspaceÚinternalTé   r   Úgenericc                    s   g | ]}ˆ   tj|¡‘qS r"   ©rž   r   Zintp©r]   ÚsrA   r"   r#   r_   C  s     c                    s   g | ]}ˆ   tj|¡‘qS r"   r¨   r©   rA   r"   r#   r_   D  s     N)ÚdataÚshapeÚstridesÚitemsizeÚparentZmeminfo) rz   ÚiterÚtobytesr   Ú	ArrayTyper~   Úlenr   r   ÚADDRSPACE_CONSTANTr   r}   r€   ÚlinkageÚglobal_constantr‚   Zget_data_typeZdtypeZget_abi_sizeofÚ
bit_lengthÚalignÚPointerTypeÚaddrspacecastZ
make_arrayr¬   r­   Zpopulate_arrayÚbitcastr«   r®   r¯   Z	_getvalue)r!   r–   ZarytyZarrÚlmodZ	constvalsZ
constarytyZconstaryr¤   rƒ   Zlldtyper¸   ZptrtyZgenptrZaryZkshapeZkstridesr"   rA   r#   Úmake_constant_array"  s:    
þÿ
 üz%CUDATargetContext.make_constant_arrayc                 C   s€   t  | d¡d ¡}d dt |¡g¡}|j |¡}|dkrdt j||j	|t
jd�}d|_d|_||_|j	jj}| | t
j¡¡S )	r�   zutf-8ó    ú$Z__conststring__Nr£   r¥   T)r   Zmake_bytearrayÚencodeÚjoinr   Zmangle_identifierÚglobalsr-   r}   r€   r   r´   rµ   r¶   r‚   r�   Úelementr»   Z
as_pointer)r!   ÚmodÚstringÚtextrD   rƒ   Zchartyr"   r"   r#   Úinsert_const_stringM  s    ÿÿ
z%CUDATargetContext.insert_const_stringc                 C   s0   |j }|  ||¡}t t d¡¡}| ||d¡S )z«
        Insert a constant string in the constant addresspace and return a
        generic i8 pointer to the data.

        This function attempts to deduplicate.
        r¢   r§   )rz   rÇ   r   r¹   r~   rº   )r!   r–   rÅ   r¼   rƒ   Z	charptrtyr"   r"   r#   Úinsert_string_const_addrspacec  s    z/CUDATargetContext.insert_string_const_addrspacec                 C   s   dS )zRun O1 function passes
        Nr"   )r!   r”   r"   r"   r#   Úoptimize_functiono  s    z#CUDATargetContext.optimize_functionc                 C   s
   t  |¡S r9   )r   Úget_ufunc_info)r!   Z	ufunc_keyr"   r"   r#   rÊ   |  s    z CUDATargetContext.get_ufunc_info)r8   )N)r3   r4   r5   Zimplement_powi_as_math_callZstrict_alignmentr:   Úpropertyr?   rB   rE   rG   r$   r   rZ   r   ra   rc   rh   rx   rp   r½   rÇ   rÈ   rÉ   rÊ   r6   r"   r"   r1   r#   r7   E   s4   




 þ
$b+r7   c                   @   s   e Zd ZdS )rb   N)r3   r4   r5   r"   r"   r"   r#   rb   €  s   rb   ))ÚreÚ	functoolsr   Zllvmlite.bindingZbindingrY   Zllvmliter   Z
numba.corer   r   r   r   r   Znumba.core.dispatcherr	   Znumba.core.errorsr
   Znumba.core.baser   Znumba.core.callconvr   r   r   r   Zcudadrvr   Z
numba.cudar   r   r   Znumba.cuda.modelsr   Úwarningsr   r   ÚcompileÚIZVALID_CHARSr7   rb   r"   r"   r"   r#   Ú<module>   s(   *  =