o
    X۷i                     @  s\  d dl m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 d dl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-d Z.ne/Z.G dd de
Z0G dd de1Z2dd Z3dd Z4G dd dZ5d d! Z6d"d# Z7dqdrd*d+Z8dsd/d0Z9G d1d2 d2Z:d3d4 Z;d5d6 Z<dtd=d>Z=dudDdEZ>dvdJdKZ?e3dwdNdOZ@e3dxdRdSZAdydUdVZBdzdYdZZC	 d{d|d_d`ZDd}dddeZEedfe$e#e$e#B ZFd~dldmZGddodpZHdS )    )annotationsN)Any
NamedTupleTypeVar)Sequence)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unsafec                   @  sF   e Zd ZU ded< ded< ded< ded< ded< d	ed
< ded< dS )Resultstr	func_namecode_cuda_types.TypeBasereturn_typeboolenable_cooperative_groupsbackendztuple[str, ...]optionsjitifyN)__name__
__module____qualname____annotations__ r.   r.   H/home/ubuntu/vllm_env/lib/python3.10/site-packages/cupyx/jit/_compile.pyr   *   s   
 r   c                   @  s   e Zd Zdd Zdd ZdS )_JitCompileErrorc                 C  s   t || _t|| _|| _d S N)type
error_typer    mesnode)selfer5   r.   r.   r/   __init__6   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startr.   r/   
<listcomp>>   s    2z,_JitCompileError.reraise.<locals>.<listcomp>z

)r5   linenogetattrjoin	enumeratesplitr3   r4   )r6   pycoder.   rA   r/   reraise;   s   z_JitCompileError.reraiseN)r*   r+   r,   r8   rK   r.   r.   r.   r/   r0   4   s    r0   c                   s   d fdd}|S )Nr5   ast.ASTc              
     sJ   z | g|R i |W S  t y     ty$ } zt || d }~ww r1   )r0   	Exception)r5   argskwargsr7   funcr.   r/   new_funcE   s   
z,transpile_function_wrapper.<locals>.new_func)r5   rL   r.   )rQ   rR   r.   rP   r/   transpile_function_wrapperD   s   rS   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 ) rT   r;   )replacer>   r@   )
num_indentr.   r/   rD   o   s    z*_parse_function_object.<locals>.<listcomp>r;   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 r.   )
isinstanceastLambdarE   )r>   r5   )end_line
start_liner.   r/   rD      s    
z4Multiple callables are found near the definition of z4, and JIT could not identify the source code for it._lambda_kernel)namerN   bodydecorator_listreturnstype_comment)callable
ValueErrorinspectgetsourcefile	TypeErrorRuntimeErrorr*   getsourcelineslenlstriprG   rZ   parserY   Moduler`   	linecachegetlineswalkFunctionDefrN   Return)	rQ   filenamelines_sourcetreefull_sourcenodesr5   r.   )r\   rX   r]   r/   _parse_function_objectP   sV   	






r{   c                   @  s    e Zd Zd
ddZdddZd	S )	GeneratedreturnNonec                 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)   )r6   r.   r.   r/   r8      s   
zGenerated.__init__r"   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   )r6   r"   r.   r.   r/   add_code   s   
zGenerated.add_codeN)r}   r~   )r"   r    r}   r~   )r*   r+   r,   r8   r   r.   r.   r.   r/   r|      s    
r|   c              	   C  s   t  }t|}t| |||||\}}|j| |f \}}	d|j}
d|
 }
|j}|j}|j}|j	}t
r7t|
 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.
    r:   z#include "cupy/float16.cuh"
)r!   r"   r$   r)   r&   r'   r(   )r|   tuple_transpile_func_objr   rG   r   r'   r(   r)   r   _is_debug_modeprintr   )rQ   
attributesmodein_typesret_type	generatedr_   r$   r!   rv   r"   r'   r(   r)   r   r.   r.   r/   	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.rU   r   rv   )rw   r.   )r   re   r{   rf   getclosurevarsdictglobals	nonlocalsbuiltinsrG   r_   rk   r    _transpile_functionr   r   )rQ   r   r   r   r   r   resultrx   rw   cvarsconstsr_   	cuda_codeenvr.   r.   r/   r      s(   





r   r<   ru   	list[str]spacesr    r}   c                   s    fdd| D S )Nc                   s   g | ]} | qS r.   r.   rW   r   r.   r/   rD      s    z_indent.<locals>.<listcomp>r.   )ru   r   r.   r   r/   _indent   s   r   values_internal_types.Exprr%   c                  G  s(   t dd | D sJ t dd | D S )Nc                 s  s    | ]	}t |tjV  qd S r1   )rY   r   Exprr>   xr.   r.   r/   	<genexpr>   s    zis_constants.<locals>.<genexpr>c                 s  s    | ]}t |tV  qd S r1   rY   r   r   r.   r.   r/   r      s    )all)r   r.   r.   r/   is_constants   s   r   c                   @  s2   e Zd ZdZdddZdddZ	dd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    r   dict[str, Constant]paramsdict[str, Data]r   r#   r   r|   c                 C  s4   || _ || _|| _i | _i | _|| _|| _d| _d S )Nr   )r   r   r   localsdeclsr   r   count)r6   r   r   r   r   r   r.   r.   r/   r8     s   
zEnvironment.__init__keyr}   Constant | Data | Nonec                 C  s@   || j v r
| j | S || jv r| j| S || jv r| j| S d S r1   )r   r   r   )r6   r   r.   r.   r/   __getitem__&  s   





zEnvironment.__getitem__rT   prefixsuffixc                 C  s<   |  j d7  _ | | j  | }| | d u r|S | ||S )Nr;   )r   get_fresh_variable_name)r6   r   r   r_   r.   r.   r/   r   /  s
   z#Environment.get_fresh_variable_nameN)
r   r    r   r   r   r   r   r#   r   r|   )r   r    r}   r   )rT   rT   )r   r    r   r    r}   r    )r*   r+   r,   __doc__r8   r   r   r.   r.   r.   r/   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_internalr0   r   rK   )rQ   r_   r   r   r   r   r   r   rw   r7   excr.   r.   r/   r   9  s   


r   c                   s  t dd | D }t| tjstdt| t| j	dkrA| j	D ]}t
| t fdddD s@td  d	t q$| j}	|	jd urMtd
t|	jdkrXtd|	jd uratdt|	jdkrlt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 r.   r   r>   kvr.   r.   r/   rD   \      z0_transpile_function_internal.<locals>.<listcomp>zNot supported: {}r   c                 3  s    | ]}| v V  qd S r1   r.   )r>   word	deco_coder.   r/   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 r.   )arg)r>   r   r.   r.   r/   rD   u      z	() takes z positional arguments but z were given.c                 S  s   g | ]\}}|t ||fqS r.   r   r>   r   tr.   r.   r/   rD   z  s    T, c                 S  s   g | ]
\}}| |d qS r1   )declvarr   r.   r.   r/   rD   }  r   c                 S  s"   g | ]\}}|j |d d qS )N;)ctyper   )r>   nr   r.   r.   r/   rD   ~  s   " rU   ()) r   itemsrY   rZ   rr   NotImplementedErrorformatr2   rk   ra   unparseanywarningswarnRuntimeWarningrN   vararg
kwonlyargskwargdefaultsrh   zipr   _transpile_stmtsr`   rG   r   r   r   voidr	   r    )rQ   r_   r   r   r   r   r   r   deco	argumentsrN   r   r   r`   params_s
local_varsheadr"   r.   r   r/   r   Z  sX   





r   oprL   rN   Sequence[Constant | Data]r   Constant | Datac                 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   r.   objr   r.   r.   r/   rD     r   z!_eval_operand.<locals>.<listcomp>_add_radd_sub_rsub)r   r   
get_pyfuncr2   r   rY   rZ   Addr   inithasattrr   r   NotImplementedr   Subr   r   	get_ufuncr   _call_ufunc)r   rN   r   pyfuncr   youtufuncr.   r.   r/   _eval_operand  s:   r   r   _kernel.ufuncdtypenumpy.dtype | Noner   c                 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: r;   zout0 = r   rT   r   Tzin{}Fin_typer   	out0_typer   r   c                 S  s   g | ]}d | dqS )ztypename inr  r.   r>   r?   r.   r.   r/   rD     s    z_call_ufunc.<locals>.<listcomp>rv   c                 S  s   g | ]
}d | d| qS )r   z_type inr.   r  r.   r.   r/   rD     r   z
template <z>
__device__ rU   z) {
    typedef z out0_type;
    return z;
}
c                 S  r   r.   r"   )r>   ar.   r.   r/   rD     r   zufunc `z` is not supported.)+rk   ninre   rY   r   r   get_ctype_from_scalarr   r   r   r   r   Scalarrh   r   guess_routiner   r_   
error_funcnoutroutine
startswith	out_typesrV   r   r   _astype_scalarr   r   rangelistrefinditerr   r   r   	_preamblerH   r"   r    rG   numpyr   )r   rN   r   r   r   r   r   r   out_typeexpr	in_paramscan_use_inline_expansionr?   template_typenames
ufunc_namer   
ufunc_codein_params_coder.   r.   r/   r     st   


 
 
r   stmtslist[ast.stmt]is_toplevelr
   c                 C  s$   g }| D ]}| t||| q|S r1   )extend_transpile_stmt)r  r   r   
codeblocksstmtr.   r.   r/   r     s   r   r$  ast.stmtc                 C  sJ  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d6gS t | tjHr#d7gS J )8zaTranspile 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.r;   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.zbreak;z	continue;)IrY   rZ   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`   rC   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   r1  r   varr_   r:  tmpr   itersloop_varr`   	init_codecondr   r"   rC  	conditionr  	then_body	else_bodyr.   r.   r/   r"    s,  











r"  r  ast.exprc                 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_internalrY   r   r   r   r   )r  r   resr.   r.   r/   r0    s   
r0  ast.expr | Nonec                   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 r.   r0  r>   r7   r   r.   r/   rD         z,_transpile_expr_internal.<locals>.<listcomp>r   r;   r&     z2Comparison of 3 or more values is not implemented.c                   rb  r.   rc  rd  re  r.   r/   rD     rf  cz/Complex type value cannot be boolean condition.z*Type mismatch in conditional expression.: r,  r   r   z ? z : r   c                   rb  r.   rc  r   re  r.   r/   rD     rf  'z' is not callable.zCalling __global__ function z) from __global__ function is not allowed.c                      g | ]}t | qS r.   r   r   r   re  r.   r/   rD         c                 S  r   r.   r   r   r.   r.   r/   rD     r   
__device__r   c                 S  r   r.   r  r   r.   r.   r/   rD     r   r   z!' is an invalid keyword to ufunc c                 S  r   r.   r   r   r.   r.   r/   rD     r   c                 S  s   g | ]	\}}||j fqS r.   r   r   r.   r.   r/   rD     s    zfunction takes z invalid number of argumentzInvalid function call 'z'.zUnbound name: zUnknown attribute: c                   rb  r.   rc  r   re  r.   r/   rD   8  rf  c                 S  s   g | ]}t |tqS r.   r   r   r.   r.   r/   rD   :  rf  c                 S  r   r.   r   r   r.   r.   r/   rD   ;  r   c                   rj  r.   rk  r   re  r.   r/   rD   =  rl  c                 S  r   r.   r  r   r.   r.   r/   rD   >  r   zSTD::make_pair(c                 S  r   r.   rm  r   r.   r.   r/   rD   C  r   zSTD::make_tuple(c                 S  r   r.   rm  r   r.   r.   r/   rD   G  r   zNot supported: type {})\rY   rZ   BoolOpr   r   r   BinOpr0  leftrightUnaryOpoperandr[   r   Comparecomparatorsrk   opsIfExprF  r`   r=  r   r   r   r   kindrh   _infer_typer  r   rG  r   r"   CallrQ   rN   keywordsr   r1  r   builtin_functions_dictr   r   BuiltinFunccallr   TypeBase_instantiate_JitRawKernel_device_funcr*   r   r   r   r   rG   r   r   popnextr>  r_   r   r   r   rf   isclass
issubclassr7  r  	Subscriptslice	_indexingr5  r6  	NameError	AttributerF   attrr   types
MethodTypeAttributeErrorTupleeltsr   Indexre   r   r2   )r  r   r   r1  rhsrq  rr  rZ  r   r   rQ   rN   rO   kwbuiltin_funcsr   fnamer$   r  r   r_   r   arrayindexr  r  	elts_coder.   re  r/   r_    s  
&









r_  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   )	rY   r   rh   r   r   r  r  r"   r;  )r  r  r   r.   r.   r/   _emit_assign_stmtP  s   

r  r:  r1  depthintc                 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>(_tempr;   rT   )rY   rZ   r5  r6  r   r   r   r   r   r  r  r   r0  r  r   re   rk   r  r  r  r8  r!  r	   )r:  r   r1  r   r  r_   r  sizer   r?   r"   r   r$  r.   r.   r/   r8  c  s>   


r8  r  r  Data | Constantc           	      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.r;   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.)rY   r   r   rh   r2   r   r   r   r   r  r  r"   	ArrayBase_ndimr  r   CArray_c_contiguous_index_32_bitsry  
child_typerk   
IndexError)	r  r  r   r?   r   ndimindex_dtype
new_carrayr   r.   r.   r/   r    s   






r  _Tr   r   _cuda_types.Scalarcasting_CastingTypec                 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 .rh  bz:Casting complex values to real discards the imaginary partr   z)(z.real())r   )rY   r   r   r   r2   r   r   r   r  rh   r"   r  can_castry  r   r   r   )r   r   r  r   from_tto_tr.   r.   r/   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   )rY   r   r   r  genericr   r   r   r   r  r  )r   r  r   cast_xr.   r.   r/   rz    s   rz  )r<   )ru   r   r   r    r}   r   )r   r   r}   r%   )r   rL   rN   r   r   r   r}   r   )
r   r   rN   r   r   r   r   r   r}   r   )r  r  r   r%   r   r   r}   r
   )r$  r%  r   r%   r   r   r}   r
   )r  r^  r   r   r}   r   )r  ra  r   r   r}   r   )r  r   r  r   r   r   r}   r
   )r   )r:  r^  r   r   r1  r   r   r%   r  r  r}   r
   )r  r   r  r   r   r   r}   r  )
r   r  r   r  r  r  r   r   r}   r  )r   r   r  r   r   r   r}   r   )I
__future__r   rZ   rf   ro   numbersr  systypingr   r   r   collections.abcr   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   r%   rG  Numberr7  version_infor   r  r    r   rM   r0   rS   r{   r|   r   r   r   r   r   r   r   r   r   r   r"  r0  r_  r  r8  r  r  r  rz  r.   r.   r.   r/   <module>   sx    


K$ 
6!
.
'
G &
 

+L
"