o
    װi                      @   s  d dl Z d dlZd dlZd dlZd dlZd dlZd dlZd dlmZm	Z	m
Z
mZmZmZmZmZ d dlZd dlZd dlZd dlmZ d dlmZ d dl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& d dl m'Z' d dl m(Z( dZ)e*ej+ej,fZ-dej.krd dlm/Z/ ee/d  Z0ne1Z0e2dg dZ3G dd de4Z5dd Z6dd Z7G dd dZ8dd Z9d d! Z:d\d#e
e1 d$e1d%e
e1 fd&d'Z;d(e#j<d%e*fd)d*Z=G d+d, d,Z>d-d. Z?d/d0 Z@d1e jAd2eee&e%f  d3e>d%ee&e%f fd4d5ZBd6ejCd2eee&e%f  d7eejD d3e>d%e%f
d8d9ZEd:e
e jF d;e*d3e>d%efd<d=ZGe6d>e jFd;e*d3e>d%efd?d@ZHe6dAe jId3e>d%e#j<fdBdCZJdAee jI d3e>d%e#j<fdDdEZKdFee&e%f dGe%d3e>d%efdHdIZL	 d]dJe jId3e>dKe%d;e*dLeMd%efdMdNZNdOe#j<dPe#j<d3e>d%ee%e&f fdQdRZOedSe&e%ee&e%f ZPdTePdUe!jQdVe0d3e>d%ePf
dWdXZRdTee&e%f dYee&e%f d3e>d%e%fdZd[ZSdS )^    N)AnyDictListOptionalSequenceTupleTypeVarUnion)ComplexWarning)runtime)	CodeBlock	_CodeType)_kernel)_raise_if_invalid_cast)jit)_cuda_types)_cuda_typerules)_internal_typesDataConstant)_builtin_funcs)
_interfaceF)      )Literal)noequivsafe	same_kindunsafeResult)	func_namecodereturn_typeenable_cooperative_groupsbackendoptionsjitifyc                   @   s   e Zd Zdd Zdd ZdS )_JitCompileErrorc                 C   s   t || _t|| _|| _d S N)type
error_typestrmesnode)selfer0    r3   F/home/ubuntu/.local/lib/python3.10/site-packages/cupyx/jit/_compile.py__init__7   s   


z_JitCompileError.__init__c                    sN   | j jt| j d d fddt|dD }| | jd | )N
end_lineno
c                    sD   g | ]\}}|d    kr krn nd| nd|   qS )   z>   )rstrip).0ilineendstartr3   r4   
<listcomp>?   s    2z,_JitCompileError.reraise.<locals>.<listcomp>z

)r0   linenogetattrjoin	enumeratesplitr-   r/   )r1   pycoder3   r>   r4   reraise<   s   z_JitCompileError.reraiseN)__name__
__module____qualname__r5   rH   r3   r3   r3   r4   r*   5   s    r*   c                    s   dt jf fdd}|S )Nr0   c              
      sJ   z | g|R i |W S  t y     ty$ } zt || d }~ww r+   )r*   	Exception)r0   argskwargsr2   funcr3   r4   new_funcF   s   
z,transpile_function_wrapper.<locals>.new_func)astAST)rP   rQ   r3   rO   r4   transpile_function_wrapperE   s   rT   c           	         s  t | stdzt| }W n ty   d}Y nw |dkr'td|  d| jdkrjt| \}}t|d t|d 	  d
fd	d
|D }t|}t|tjsZJ t|jdkscJ |jd |fS |du rvtd|  dd
t|}t| \}t|  d
|}t|} fdd
t|D }t|dkrtd|  d|d }tjd|jt|jgg ddd|fS )a/  Returns the tuple of ``ast.FunctionDef`` object and the source string
    for the given callable ``func``.

    ``func`` can be a ``def`` function or a ``lambda`` expression.

    The source is returned only for informational purposes (i.e., rendering
    an exception message in case of an error).
    z!`func` must be a callable object.Nz<stdin>z/JIT needs access to the Python source code for zf but it cannot be retrieved within the Python interactive interpreter. Consider using IPython instead.z<lambda>r    c                    s   g | ]}| d   ddqS ) rU   r8   )replacer;   r=   )
num_indentr3   r4   rA   p   s    z*_parse_function_object.<locals>.<listcomp>r8   z+JIT needs access to Python source code for zX but could not be located.
(hint: it is likely you passed a built-in function or method)c                    s6   g | ]}t |tjr|j  kr k rn n|qS r3   )
isinstancerR   LambdarB   )r;   r0   )end_line
start_liner3   r4   rA      s    
z4Multiple callables are found near the definition of z4, and JIT could not identify the source code for it._lambda_kernel)namerM   bodydecorator_listreturnstype_comment)callable
ValueErrorinspectgetsourcefile	TypeErrorRuntimeErrorrI   getsourcelineslenlstriprD   rR   parserZ   Moduler`   	linecachegetlineswalkFunctionDefrM   Return)	rP   filenamelines_sourcetreefull_sourcenodesr0   r3   )r\   rY   r]   r4   _parse_function_objectQ   sV   	






r{   c                   @   s(   e Zd ZdddZdeddfddZdS )		GeneratedreturnNc                 C   sD   g | _ i | _d| _d| _d| _d| _d| _tjrdnd| _	d| _
d S )NF)z-DCUPY_JIT_MODEz--std=c++17z-DCUB_DISABLE_BF16_SUPPORTnvccnvrtc)codesdevice_function	enable_cg
include_cginclude_cg_memcpy_asyncinclude_cuda_barrierr(   r   is_hipr'   r)   )r1   r3   r3   r4   r5      s   
zGenerated.__init__r$   c                 C   s6   || j vr| j | t| j tjkrtdd S d S )Nz(Number of functions exceeds upper limit.)r   appendrk   r   _n_functions_upperlimitre   )r1   r$   r3   r3   r4   add_code   s   
zGenerated.add_code)r}   N)rI   rJ   rK   r5   r.   r   r3   r3   r3   r4   r|      s    
r|   c              	   C   s|   t  }t|}t| |||||\}}|j| |f \}}	d|j}
|j}|j}|j}|j	}t
r3t|
 t||
|||||dS )as  Transpiles the target function.

    Args:
        func (function): Target function.
        attributes (list of str): Attributes of the generated CUDA function.
        mode ('numpy' or 'cuda'): The rule for typecast.
        in_types (list of _cuda_types.TypeBase): Types of the arguments.
        ret_type (_cuda_types.TypeBase or None): Type of the return value.
    r7   )r#   r$   r%   r)   r&   r'   r(   )r|   tuple_transpile_func_objr   rD   r   r'   r(   r)   r   _is_debug_modeprintr"   )rP   
attributesmodein_typesret_type	generatedr_   r%   r#   rv   r$   r'   r(   r)   r   r3   r3   r4   	transpile   s"   
r   c                 C   s   | |f|j v r|j | |f }|d u rtd|S t| \}}t| }	tdi |	j|	j|	j}
d	|}|j
}t|j dkrK|dtt|j  7 }d |j | |f< t|||||
||||d	\}}||jf|j | |f< || ||jfS )Nz$Recursive function is not supported.rV   r   rv   )rw   r3   )r   re   r{   rf   getclosurevarsdictglobals	nonlocalsbuiltinsrD   r_   rk   r.   _transpile_functionr   r   )rP   r   r   r   r   r   resultrx   rw   cvarsconstsr_   	cuda_codeenvr3   r3   r4   r      s(   





r   r9   ru   spacesr}   c                    s    fdd| D S )Nc                    s   g | ]} | qS r3   r3   rX   r   r3   r4   rA      s    z_indent.<locals>.<listcomp>r3   )ru   r   r3   r   r4   _indent   s   r   valuesc                  G   s(   t dd | D sJ t dd | D S )Nc                 s   s    | ]	}t |tjV  qd S r+   )rZ   r   Exprr;   xr3   r3   r4   	<genexpr>   s    zis_constants.<locals>.<genexpr>c                 s   s    | ]}t |tV  qd S r+   rZ   r   r   r3   r3   r4   r      s    )all)r   r3   r3   r4   is_constants   s   r   c                
   @   sx   e Zd ZdZdedeeef deeef dej	de
f
ddZd	ed
eeeef  fddZ	ddeded
efddZdS )Environmenta:  Environment of the scope

    Attributes:
        mode ('numpy' or 'cuda'): The rule for typecast.
        consts (dict): The dictionary with keys as the variable names and
            the values as the data that is determined at compile-time.
        params (dict): The dictionary of function arguments with keys as
            the variable names and the values as the Data.
        locals (dict): The dictionary with keys as the variable names and the
            values as the Data stored at the local scope of the function.
        ret_type (_cuda_types.TypeBase):
            The type of return value of the function.
            If it is initialized to be ``None``, the return type must be
            inferred until the end of transpilation of the function.
        generated (Generated): Generated CUDA functions.
    r   r   paramsr   r   c                 C   s4   || _ || _|| _i | _i | _|| _|| _d| _d S )Nr   )r   r   r   localsdeclsr   r   count)r1   r   r   r   r   r   r3   r3   r4   r5     s   
zEnvironment.__init__keyr}   c                 C   s@   || j v r
| j | S || jv r| j| S || jv r| j| S d S r+   )r   r   r   )r1   r   r3   r3   r4   __getitem__$  s   





zEnvironment.__getitem__rU   prefixsuffixc                 C   s<   |  j d7  _ | | j  | }| | d u r|S | ||S )Nr8   )r   get_fresh_variable_name)r1   r   r   r_   r3   r3   r4   r   -  s
   z#Environment.get_fresh_variable_nameN)rU   rU   )rI   rJ   rK   __doc__r.   r   r   r   r   TypeBaser|   r5   r   r	   r   r   r3   r3   r3   r4   r     s,    



r   c             
   C   s^   zt | |||||||W S  ty' }	 z|	}
tr|
| W Y d}	~	nd}	~	ww |
| J )a  Transpile the function
    Args:
        func (ast.FunctionDef): Target function.
        name (str): Function name.
        attributes (str): The attributes of target function.
        mode ('numpy' or 'cuda'): The rule for typecast.
        consts (dict): The dictionary with keys as variable names and
            values as concrete data object.
        in_types (list of _cuda_types.TypeBase): The types of arguments.
        ret_type (_cuda_types.TypeBase): The type of return value.

    Returns:
        code (str): The generated CUDA code.
        env (Environment): More details of analysis result of the function,
            which includes preambles, estimated return type and more.
    N)_transpile_function_internalr*   r   rH   )rP   r_   r   r   r   r   r   r   rw   r2   excr3   r3   r4   r   7  s   


r   c                    s  t dd | D }t| tjstdt| t| j	dkrFt
jdkrF| j	D ]}t| t fdddD sEtd	  d
t q)| j}	|	jd urRtdt|	jdkr]td|	jd urftdt|	jdkrqtddd |	jD }
t|
t|krt| dt|
 dt| dt dd t|
|D }t|||||}t| jd|}ddd t|
|D }dd |j D }|jd u rtj|_| d|j d| d| d}t ||| }t!||fS )Nc                 S   s   g | ]
\}}|t |fqS r3   r   r;   kvr3   r3   r4   rA   Z      z0_transpile_function_internal.<locals>.<listcomp>zNot supported: {}r   )r   	   c                 3   s    | ]}| v V  qd S r+   r3   )r;   word	deco_coder3   r4   r   d  s    z/_transpile_function_internal.<locals>.<genexpr>)	rawkernel	vectorizez
Decorator z may not supported in JIT.z#`*args` is not supported currently.z4keyword only arguments are not supported currently .z&`**kwargs` is not supported currently.z+Default values are not supported currently.c                 S      g | ]}|j qS r3   )arg)r;   r   r3   r3   r4   rA   u      z	() takes z positional arguments but z were given.c                 S   s   g | ]\}}|t ||fqS r3   r   r;   r   tr3   r3   r4   rA   z  s    T, c                 S   s   g | ]
\}}| |d qS r+   )declvarr   r3   r3   r4   rA   }  r   c                 S   s"   g | ]\}}|j |d d qS )N;)ctyper   )r;   nr   r3   r3   r4   rA   ~  s   " rV   ())"r   itemsrZ   rR   rr   NotImplementedErrorformatr,   rk   ra   sysversion_infounparseanywarningswarnRuntimeWarningrM   vararg
kwonlyargskwargdefaultsrh   zipr   _transpile_stmtsr`   rD   r   r   r   voidr   r.   )rP   r_   r   r   r   r   r   r   deco	argumentsrM   r   r   r`   params_s
local_varsheadr$   r3   r   r4   r   X  sZ   






r   oprM   r   c                 C   sD  t | rtt| }t|dd |D  S t| tjrT|\}}t	||}t	||}t
|jdr@|j|||}|tur@|S t
|jdrT|j|||}|turT|S t| tjr|\}}t	||}t	||}t
|jdr~|j|||}|tur~|S t
|jdr|j|||}|tur|S t|jt| }t||d |S )Nc                 S   r   r3   objr   r3   r3   r4   rA     r   z!_eval_operand.<locals>.<listcomp>_add_radd_sub_rsub)r   r   
get_pyfuncr,   r   rZ   rR   Addr   inithasattrr   r   NotImplementedr   Subr   r   	get_ufuncr   _call_ufunc)r   rM   r   pyfuncr   youtufuncr3   r3   r4   _eval_operand  s:   r   r   dtypec                 C   s  t || jkrtdg }|D ]+}t|tr t|j|jj	}nt|j
tjr,|j
j	}n	td|j
 d|| qt| t|||j}|d u rUtd| j d| |jd ur^|  | jdkrH|jdrHt|jd }|jdd	}	g }
t||jD ]\}}t|t|d
|}t||}|
| qd}t| jD ]}t tt d!||jdkrd}d| d|jv rd}q|j"#| j$ |rt%|
D ]\}}|	d| |j&}	qd|	dt'| d }	nWd(dd t| jD }| j dt't)	|jd  }d(dd t| jD }d| d| d| d| d| d|	 d}|j"#| d(dd |
D }| d| d}	t|	|S t*d| j d )!Nzinvalid number of argumentszcupy.ufunc: z is unsupported"z(" does not support for the input types: r8   zout0 = r   rU   r    Tzin{}Fin_typer   	out0_typer   r   c                 S   s   g | ]}d | dqS )ztypename inr   r3   r;   r<   r3   r3   r4   rA     s    z_call_ufunc.<locals>.<listcomp>rv   c                 S   s   g | ]
}d | d| qS )r   z_type inr3   r   r3   r3   r4   rA     r   z
template <z>
__device__ rV   z) {
    typedef z out0_type;
    return z;
}
c                 S   r   r3   r$   )r;   ar3   r3   r4   rA     r   zufunc `z` is not supported.)+rk   ninre   rZ   r   r   get_ctype_from_scalarr   r   r   r   r   Scalarrh   r   guess_routiner   r_   
error_funcnoutroutine
startswith	out_typesrW   r   r   _astype_scalarr   r   rangelistrefinditerr   r   r   	_preamblerE   r$   r.   rD   numpyr   )r   rM   r   r   r   r   r   r   out_typeexpr	in_paramscan_use_inline_expansionr<   template_typenames
ufunc_namer   
ufunc_codein_params_coder3   r3   r4   r     st   


 
 
r   stmtsis_toplevelc                 C   s$   g }| D ]}| t||| q|S r+   )extend_transpile_stmt)r  r  r   
codeblocksstmtr3   r3   r4   r     s   r   r  c                 C   sN  t | tjr
tdt | tjtjfrtdt | tjrZt| j|}t |t	r0|j
du r0dgS t||}|j}|jdu rB||_n|j|krRtd|j d| d|j dgS t | tjrdtd	t | tjrt| jd
krutdt| j|}| jd }t|rt |tjr|j}t |j
ts|rt || trtd| d||j|< g S tdt||}t||||S t | tjr9t| j|}t| j|}t |tstd|j t||}t|d|j}t| j ||f|}	t |tsJ t |	tsJ t |jt!j"sJ t |	jt!j"sJ t#|	jj$|jj$d d|j%d|j | d |j&||	 d gS t | tj'r	t| j(dkrLtdt | jtjsVJ | jj}t| j)|}
|| }|du r|t| jj|
j}||j*|< ||j+|< n t |t	rtd|j|
jkrtd| d|j d|
j t |
t,j-stdt.| j/d|}|
j d|
j0j d|
j1j d |
j2j }d!}|
j3d"u rd#}n|
j3du rd$}d%| d| d&}t4|| d'g| g}|
j5}|d"u rd(g| }|S |durd)| d*g| }|S t | tj6rtd+t | tj7rOt| j(dkr'tdt| j8|}t9|t!j:d,|}t||}t.| j/d|}d-|j d*}t4||gS t | tj;rt| j8|}t|rq|j
rh| j/n| j(}t.|||S d.|j d*}t.| j/d|}t.| j(d|}t4||t4d/|gS t | tj<tj=frtd0t | tj>tj?frtd1t | tj@rt| j8|}t|r|j
sJ dgS d2| d3 gS t | tjAtjBfrtd4t | tjCtjDfrtd5t | tjErt| j|}t|rdgS |jd gS t | tjFrdgS t | tjGrtdt | tjHr%tdJ )6zaTranspile the statement.

    Returns (list of [CodeBlock or str]): The generated CUDA code.
    z!class is not supported currently.z-Nested functions are not supported currently.Nzreturn;z!Failed to infer the return type: z or zreturn r   z!`del` is not supported currently.r8   Not implemented.r   zType mismatch of variable: ``z.Cannot assign constant value not at top-level.zCannot augassign to _tmp_r    z{ &z; z; }zwhile-else is not supported.z'loop counter must not be constant value!Data type mismatch of variable: ``:  != z.for-loop is supported only for range iterator.Fz __it = z, __stop = z, __step = z+__step >= 0 ? __it < __stop : __it > __stopTz__it < __stopz__it > __stopzfor (z; __it += __step)z = __it;z#pragma unrollz#pragma unroll(r   z`async for` is not allowed.r!   zwhile (zif (elsez#Switching contexts are not allowed.zthrow/catch are not allowed.zassert(z);z0Cannot import modules from the target functions.z3Cannot use global/nonlocal in the target functions.)IrZ   rR   ClassDefr   rr   AsyncFunctionDefrs   _transpile_exprvaluer   r   r   r   r   r   re   r$   DeleteAssignrk   targetsr   Nameid_typeclassesrh   r   _transpile_assign_stmt	AugAssigntargetr   r   r   r   r  r   r   r   assignFororelseiterr   r   r   Ranger   r`   r@   stopstepstep_is_positiver   unrollAsyncForWhiletestr	  bool_IfWith	AsyncWithRaiseTryAssertImport
ImportFromGlobalNonlocalr   PassBreakContinue)r  r  r   r)  r   varr_   r2  tmpr   itersloop_varr`   	init_codecondr   r$   r;  	conditionr  	then_body	else_bodyr3   r3   r4   r    s,  











r  r  c                 C   s,   t | |}t|trt|jtjr|jS |S )z`Transpile the statement.

    Returns (Data): The CUDA code and its type of the expression.
    )_transpile_expr_internalrZ   r   r   r   r   )r  r   resr3   r3   r4   r(    s   
r(  c                    s  t | tjr( fdd| jD }|d }|dd  D ]}t| j||f }q|S t | tjrCt| j }t| j	 }t| j||f S t | tj
rWt| j }t| j|f S t | tjratdt | tjr| jg| j }t|dkrxtd fdd|D }t| jd | S t | tjrt| j }t| j }t| j }	t | tr| jr|S |	S |jjjd	krtd
t||	 t|	| }}	|jj|	jjkrtd|jj d|	jj t|tjd }t d|j! d|j! d|	j! d|jS t | tj"rFt| j# }
 fdd| j$D }i }| j%D ]}|j&d usJ t|j' ||j&< qt(j)}t*|
r;|
j|v r;||
j }
t |
t+j,rN|
j- g|R i |S t |
ts\td|
 d|
j}
t |
t.jj/rs|
j0 g|R i |S t |
t.j1r|
j2std|
j3j4 d fdd|D }t5dd |D }t6|
j3dg j7|d  j8\}}d9dd |D }t | d| d|S t |
t:j;r|<dtd j}t|dkrt=t>|}td| d|
j? t@|
|| S t*g || R  rdd |D }tAdd |B D }t|
|i |S tCD|
r=tE|
tFr=t|dkr/td |
 d!tG|
}t|d |d S td"|
j4 d#| d u rOtd S t | tjr[t| j'S t | tjHrtt| j' }t| jI }tJ|| S t | tjKr | jL }|d u rtMd$| jL |S t | tjNrt| j' }t |trttO|j| jPS t |t rtQ|j| jPrtO|j| jPd }t |tRjSr||S t |t r|S tTd%| jP t | tjUr> fd&d| jVD }tWd'd |D rtt5d(d |D S  fd)d|D }d9d*d |D }t|dkr-t d+| dtUd,d |D S t d-| dtUd.d |D S t | tjXrKt| j' S tYd/Zt[| )0Nc                       g | ]}t | qS r3   r(  r;   r2   r   r3   r4   rA         z,_transpile_expr_internal.<locals>.<listcomp>r   r8   r     z2Comparison of 3 or more values is not implemented.c                    rX  r3   rY  rZ  r[  r3   r4   rA     r\  cz/Complex type value cannot be boolean condition.z*Type mismatch in conditional expression.: r$  r!   r   z ? z : r   c                    rX  r3   rY  r   r[  r3   r4   rA     r\  'z' is not callable.zCalling __global__ function z) from __global__ function is not allowed.c                       g | ]}t | qS r3   r   r   r   r[  r3   r4   rA         c                 S   r   r3   r   r   r3   r3   r4   rA     r   
__device__r   c                 S   r   r3   r   r   r3   r3   r4   rA     r   r   z!' is an invalid keyword to ufunc c                 S   r   r3   r   r   r3   r3   r4   rA     r   c                 S   s   g | ]	\}}||j fqS r3   r   r   r3   r3   r4   rA     s    zfunction takes z invalid number of argumentzInvalid function call 'z'.zUnbound name: zUnknown attribute: c                    rX  r3   rY  r   r[  r3   r4   rA   8  r\  c                 S   s   g | ]}t |tqS r3   r   r   r3   r3   r4   rA   :  r\  c                 S   r   r3   r   r   r3   r3   r4   rA   ;  r   c                    r`  r3   ra  r   r[  r3   r4   rA   =  rb  c                 S   r   r3   r   r   r3   r3   r4   rA   >  r   zSTD::make_pair(c                 S   r   r3   rc  r   r3   r3   r4   rA   C  r   zSTD::make_tuple(c                 S   r   r3   rc  r   r3   r3   r4   rA   G  r   zNot supported: type {})\rZ   rR   BoolOpr   r   r   BinOpr(  leftrightUnaryOpoperandr[   r   Comparecomparatorsrk   opsIfExpr>  r`   r5  r   r   r   r   kindrh   _infer_typer	  r   r?  r   r$   CallrP   rM   keywordsr   r)  r   builtin_functions_dictr   r   BuiltinFunccallr   r   _instantiate_JitRawKernel_device_funcrI   r   r   r   r   rD   r   r   popnextr6  r_   r   r   r   rf   isclass
issubclassr/  r  	Subscriptslice	_indexingr-  r.  	NameError	AttributerC   attrr   types
MethodTypeAttributeErrorr   eltsr   Indexre   r   r,   )r  r   r   r)  rhsrg  rh  rR  r   r   rP   rM   rN   kwbuiltin_funcsr   fnamer%   r  r   r_   r   arrayindexr  r  	elts_coder3   r[  r4   rV    s  
&









rV  lvaluervaluec                 C   s   t | tr	tdt | jtjr t |jtjr t|| jd|}n| j|jkr6td| j d| j d|j | j| |d gS )Nz/lvalue of assignment must not be constant valuer    r"  r#  r$  r   )	rZ   r   rh   r   r   r  r	  r$   r3  )r  r  r   r3   r3   r4   _emit_assign_stmtP  s   

r  r2  r)  depthc                 C   s  t | tjr7| j}|| }|d u r1t||j}||j|< |r,|dkr,|j||d gS ||j|< t	|||S t | tj
rLtt| ||}t	|||S t | tjrt |jtjsat|j dt| j}t|jj|krvtd| dt|jj|k rtd| d|jd| |d g}t|D ]&}	d|	 d	| d}
|jj|	 }t| j|	 |t|
|||d
 }|| qtd|gS J )Nr   r   z cannot be unpackz$too many values to unpack (expected r   z&not enough values to unpack (expected _temp	STD::get<z>(_tempr8   rU   )rZ   rR   r-  r.  r   r   r   r   r   r  r~  r   r(  r   r   re   rk   r  r  r
  r0  r  r   )r2  r   r)  r  r  r_   r  sizer   r<   r$   r   r  r3   r3   r4   r0  c  s>   


r0  r  r  c           	      C   s  t | trt |trt| j|j S tt| j dt| |} t | jtj	rHt |trD|j}| jj
| }td| d| j d|S tdt | jtjrSt||}| jj}t |jtjr|jj}|dkrmtd|dkrt| jj| jjd | jj| jj}t| j d	|j d|S |jd
vrtdt| j d|j d| jjS t |jtj	rE|t|jj
k rtd| |jj
D ]}t |tjstd|jjdvrtdq|t|jj
krt| jj| jjt|jj
 | jj| jj}|j dt|jj
 d}t| j d	| d|S |dkr"t| j d| jjS |dkr6t| j d|j d| jjS t| j d|j d| jjS t |jtjrQtdJ t| j d)Nz) is not subscriptable with non-constants.r  z>(r   z.Tuple is not subscriptable with non-constants.r   z1Scalar indexing is not supported for 0-dim array.r8   z
._slicing(uizArray indices must be integers.[]z+The number of indices is beyond array dim: zArray indices must be scalar.iuzArray indices must be integer.z, Dim<z>()z[0]z[STD::get<0>(z)]z._indexing(z#Advanced indexing is not supported.Fz is not subscriptable.)rZ   r   r   rh   r,   r   r   r   r   r   r  r$   	ArrayBase_ndimr  r   CArray_c_contiguous_index_32_bitsro  
child_typerk   
IndexError)	r  r  r   r<   r   ndimindex_dtype
new_carrayr   r3   r3   r4   r    s   






r  _Tr   r   castingc                 C   s   t | trt | trJ t|j| jS t | jtjs$t	| j
 d| jj}|j}||kr1| S t|d||sIt	d| d| d| d|jdkrk|jdkrk|jdkr^td	t td
| d| j
 d|S td
| d| j
 d|S )Nz is not scalar type.r   zCannot cast from 'z' to z with casting rule .r^  bz:Casting complex values to real discards the imaginary partr   z)(z.real())r   )rZ   r   r   r   r,   r   r   r   r  rh   r$   r  can_castro  r   r   r
   )r   r   r  r   from_tto_tr3   r3   r4   r	    s.   

r	  hintc                 C   s^   t | trt | jtjrt| |S t||}t |jtj	s!J t
| |jd|}t||S )Nr    )rZ   r   r   r  genericr   r   r   r   r  r	  )r   r  r   cast_xr3   r3   r4   rp    s   rp  )r9   )r   )TrR   collectionsrf   ro   numbersr  r   typingr   r   r   r   r   r   r   r	   r   r  r  cupy.exceptionsr
   cupy_backends.cuda.apir   cupy._core._codeblockr   r   
cupy._corer   cupy._core._dtyper   cupyxr   	cupyx.jitr   r   r   cupyx.jit._internal_typesr   r   r   r   r   boolr?  Numberr/  r   r   _CastingTyper.   
namedtupler"   rL   r*   rT   r{   r|   r   r   r   r   r   r   r   r   rS   r   r   r   r   r  r   r  r  r(  rV  r  intr0  r  r  r  r	  rp  r3   r3   r3   r4   <module>   s,   (
K$ 6!0

'
G
 &
 


+

L
"

