o
    X۷i                     @  sb  d dl mZ d dlmZ d dlmZ d dlmZ dd ZG dd dej	Z
dddZdddZdddZdddZdddZede
 Zede
 Zede
 ZedgdddZed gd!d" Zed#gd$d% Zed&gd'd( Zed)gdd*d+Zed gdd,d-Zed.g	dd/d0Zed.g	dd1d2Zed3gd4d5 Zed6gd7d8 Zed9gd:d; Zed.g	dd<d=Zed.g	dd>d?Zed@g	ddAdBZ edCgddDdEZ!edCgddFdGZ"ed gdHdI Z#G dJdK dKej$Z%edLgddMdNZ&G dOdP dPej$Z'edQgddRdSZ(edTgddUdVZ)edWgddXdYZ*edWg	ddZd[Z+ed\gd]d^ Z,ed\gd_d` Z-edagdbdc Z.edagddde Z/edfgdgdh Z0edfgdidj Z1edkgdldm Z2edngddodpZ3edqg	ddrdsZ4edqg	ddtduZ5edqg	ddvdwZ6edqg	ddxdyZ7edqg	ddzd{Z8edqg	dd|d}Z9edqg	dd~dZ:edqg	dddZ;edCgdddZ<edCg	dddZ=edCgdddZ>edCg	dddZ?edgdd Z@edgdddZAedg	dddZBed gdd ZCdS )    )annotations)_internal_types)_cuda_types)Datac                   s    fdd}|S )Nc                   s   G  fdddt j}| S )Nc                      s   e Zd Z fddZdS )z7_wrap_thrust_func.<locals>.wrapper.<locals>.FuncWrapperc                   sz   D ]} j d| d q j d  j d d j _ fdd|D } fdd	D } g|R i |S )
Nz
#include <>z$#include <thrust/execution_policy.h>z#include <thrust/functional.h>nvccc                   s   g | ]}t | qS  _Datainit.0a)envr   F/home/ubuntu/vllm_env/lib/python3.10/site-packages/cupyx/jit/thrust.py
<listcomp>   s    zP_wrap_thrust_func.<locals>.wrapper.<locals>.FuncWrapper.call.<locals>.<listcomp>c                   s   i | ]}|t |  qS r   r	   )r   kr   kwargsr   r   
<dictcomp>   s    zP_wrap_thrust_func.<locals>.wrapper.<locals>.FuncWrapper.call.<locals>.<dictcomp>)	generatedadd_codebackend)selfr   argsr   header	data_argsdata_kwargsfuncheadersr   r   call   s   z<_wrap_thrust_func.<locals>.wrapper.<locals>.FuncWrapper.callN)__name__
__module____qualname__r!   r   r   r   r   FuncWrapper
   s    r%   )r   BuiltinFunc)r   r%   r    )r   r   wrapper	   s   
z"_wrap_thrust_func.<locals>.wrapperr   )r    r(   r   r'   r   _wrap_thrust_func   s   r)   c                   @  s   e Zd ZdS )_ExecPolicyTypeN)r"   r#   r$   r   r   r   r   r*      s    r*   exec_policyr
   c                 C  s    t | jtst| j dd S )Nz must be execution policy type)
isinstancectyper*   	TypeErrorcode)r+   r   r   r   _assert_exec_policy_type   s   r0   r   returnNonec                 C  s,   t | jtjstd| j d| j dd S )N`z` must be of pointer type: `)r,   r-   r   PointerBaser.   r/   )r   r   r   r   _assert_pointer_type!   s   r5   bc              
   C  s:   | j |j krtd| j d|j d| j  d|j  d	d S )Nr3   ` and `` must be of the same type: `` != `)r-   r.   r/   r   r6   r   r   r   _assert_same_type'   s   r;   c              
   C  sR   t |  t | | jj|jjkr'td| j d|j d| jj d|jj d	d S )Nr3   r7   z%` must be of the same pointer type: `r9   )r5   r-   
child_typer.   r/   typer:   r   r   r   _assert_same_pointer_type.   s   r>   c              
   C  sF   t |  | jj|jkr!td| j d|j d| jj d|j d	d S )Nz`*r7   r8   r9   r3   )r5   r-   r<   r.   r/   r:   r   r   r   _assert_pointer_of8   s   r?   zthrust::hostzthrust::devicezthrust::seqzthrust/adjacent_difference.hNc                 C  `   t | t|| t|| |durtd||||g}ddd |D }td| d|jS )z3Computes the differences of adjacent elements.
    N!binary_op option is not supported, c                 S     g | ]}|j qS r   r/   r   r   r   r   r   O       z'adjacent_difference.<locals>.<listcomp>zthrust::adjacent_difference()r0   r;   r>   NotImplementedErrorjoinr
   r-   r   r+   firstlastresult	binary_opr   paramsr   r   r   adjacent_differenceE      

rP   zthrust/binary_search.hc                 G  s"  t | t| t|| dt|  krdkr4n n|d }t|dkr)|d nd}t|| tj}n;dt|  kr@dkrkn t	d|d }|d }	|d }
t|dkr[|d nd}t|| t||	 |
j}nt	d|durwt
d|||g|}d	d
d |D }td| d|S );Attempts to find the element value with binary search.
          r   N      z0Invalid number of inputs of thrust.binary_searchcomp option is not supportedrB   c                 S  rC   r   rD   r   r   r   r   r   r   rE   z!binary_search.<locals>.<listcomp>zthrust::binary_search(rF   )r0   r5   r;   lenr?   r   bool_r>   r-   r.   rH   rI   r
   r   r+   rK   rL   r   valuecompresult_ctypevalue_first
value_lastrM   rO   r   r   r   binary_searchW   .   

	

r`   zthrust/copy.hc                 C  sP   t | t|| t|| ||||g}ddd |D }td| d|jS )zCopies the elements.
    rB   c                 S  rC   r   rD   r   r   r   r   r      rE   zcopy.<locals>.<listcomp>zthrust::copy(rF   )r0   r;   r>   rI   r
   r-   r   r+   rK   rL   rM   r   rO   r   r   r   copyv   s   

rc   zthrust/count.hc                 C  sf   t | t|jtjstd|j|jkrtd||||g}ddd |D }td| dtjS )zMCounts the number of elements in [first, last) that equals to ``value``.
    z`first` must be of pointer typez+`first` and `last` must be of the same typerB   c                 S  rC   r   rD   r   r   r   r   r      rE   zcount.<locals>.<listcomp>zthrust::count(rF   )	r0   r,   r-   r   r4   r.   rI   r
   int32r   r+   rK   rL   r[   r   rO   r   r   r   count   s   rf   zthrust/equal.hc                 C  s`   t | t|| t|| |durtd||||g}ddd |D }td| dtjS )z2Returns true if the two ranges are identical.
    N#binary_pred option is not supportedrB   c                 S  rC   r   rD   r   r   r   r   r      rE   zequal.<locals>.<listcomp>zthrust::equal(rF   )r0   r;   r>   rH   rI   r
   r   rY   )r   r+   first1last1first2binary_predr   rO   r   r   r   equal   rQ   rl   c                 C  sj   t | t| t|| |durtd||||g}ddd |D }td| dt|j|jgS )z<Attempts to find the element value in an ordered range.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r      rE   zequal_range.<locals>.<listcomp>zthrust::equal_range(rF   	r0   r5   r;   rH   rI   r
   r   Tupler-   )r   r+   rK   rL   r[   r\   r   rO   r   r   r   equal_range   s   

ro   zthrust/scan.hc           	      C  sr   t | t|| t|| |durtd||||g}|dur%|| ddd |D }td| d|jS )z0Computes an exclusive prefix sum operation.
    NrA   rB   c                 S  rC   r   rD   r   r   r   r   r      rE   z"exclusive_scan.<locals>.<listcomp>zthrust::exclusive_scan(rF   )r0   r;   r>   rH   appendrI   r
   r-   )	r   r+   rK   rL   rM   r   rN   r   rO   r   r   r   exclusive_scan   s   


rq   c	                 C  s   t | t|| t|| |durtd|durtd|||||g}	|dur.|	| ddd |	D }
td|
 d|jS )	zaComputes an exclusive prefix sum operation by key.
    _assert_exec_policy_type(exec_policy)
    Nrg   rA   rB   c                 S  rC   r   rD   r   r   r   r   r      rE   z)exclusive_scan_by_key.<locals>.<listcomp>zthrust::exclusive_scan_by_key(rF   )r5   r;   r>   rH   rp   rI   r
   r-   )r   r+   rh   ri   rj   rM   r   rk   rN   r   rO   r   r   r   exclusive_scan_by_key   s   


rr   zthrust/fill.hc                 C  sF   t | t|| ||||g}ddd |D }td| dtjS )z5Assigns the value to every element in the range.
    rB   c                 S  rC   r   rD   r   r   r   r   r      rE   zfill.<locals>.<listcomp>zthrust::fill(rF   )r0   r;   rI   r
   r   voidre   r   r   r   fill   s
   
rt   zthrust/find.hc                 C  N   t | t| t|| ||||g}ddd |D }td| d|jS )z>Finds the first iterator whose value equals to ``value``.
    rB   c                 S  rC   r   rD   r   r   r   r   r      rE   zfind.<locals>.<listcomp>zthrust::find(rF   r0   r5   r;   rI   r
   r-   re   r   r   r   find   s   
rw   zthrust/gather.hc                 C  sZ   t | t| t|| t|| |||||g}ddd |D }td| d|jS )zFCopies elements from source into destination  according to a map.
    rB   c                 S  rC   r   rD   r   r   r   r   r     rE   zgather.<locals>.<listcomp>zthrust::gather(rF   )r0   r5   r;   r>   rI   r
   r-   )r   r+   	map_firstmap_lastinput_firstrM   r   rO   r   r   r   gather      

r{   c                 C  r@   )z0Computes an inclusive prefix sum operation.
    NrA   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   z"inclusive_scan.<locals>.<listcomp>zthrust::inclusive_scan(rF   rG   rJ   r   r   r   inclusive_scan  s   

r}   c           
      C  sz   t | t| t|| t|| |durtd|dur"td|||||g}ddd |D }	td|	 d|jS )	z7Computes an inclusive prefix sum operation by key.
    Nrg   rA   rB   c                 S  rC   r   rD   r   r   r   r   r   .  rE   z)inclusive_scan_by_key.<locals>.<listcomp>zthrust::inclusive_scan_by_key(rF   )r0   r5   r;   r>   rH   rI   r
   r-   )
r   r+   rh   ri   rj   rM   rk   rN   r   rO   r   r   r   inclusive_scan_by_key  s   

r~   zthrust/inner_product.hc           
      C  sr   t | t|| t|| |durtd|durtd|||||g}ddd |D }	td|	 d|jS )	z/Calculates an inner product of the ranges.
    Nz"binary_op1 option is not supportedz"binary_op2 option is not supportedrB   c                 S  rC   r   rD   r   r   r   r   r   @  rE   z!inner_product.<locals>.<listcomp>zthrust::inner_product(rF   rG   )
r   r+   rh   ri   rj   r   
binary_op1
binary_op2r   rO   r   r   r   inner_product2  s   

r   zthrust/sort.hc                 C  \   t | t| t|| |durtd|||g}ddd |D }td| dtjS )z<Returns true if the range is sorted in ascending order.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r   Q  rE   zis_sorted.<locals>.<listcomp>zthrust::is_sorted(rF   )r0   r5   r;   rH   rI   r
   r   rY   r   r+   rK   rL   r\   r   rO   r   r   r   	is_sortedG     

r   c                 C  \   t | t| t|| |durtd|||g}ddd |D }td| d|jS )z=Returns the last iterator for which the range is sorted.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r   _  rE   z#is_sorted_until.<locals>.<listcomp>zthrust::is_sorted_until(rF   r0   r5   r;   rH   rI   r
   r-   r   r   r   r   is_sorted_untilU  r   r   c                 G  "  t | t| t|| dt|  krdkr4n n|d }t|dkr)|d nd}t|| |j}n;dt|  kr@dkrkn td|d }|d }	|d }
t|dkr[|d nd}t|| t||	 |
j}ntd|durwtd|||g|}d		d
d |D }t
d| d|S )rR   rS   rT   r   NrU   rV   z.Invalid number of inputs of thrust.lower_boundrW   rB   c                 S  rC   r   rD   r   r   r   r   r   ~  rE   zlower_bound.<locals>.<listcomp>zthrust::lower_bound(rF   r0   r5   r;   rX   r?   r-   r>   r.   rH   rI   r
   rZ   r   r   r   lower_boundc  ra   r   c                   @     e Zd ZdddZdS )_ConstantIteratorr1   strc                 C     | j }d| dS )Nzthrust::constant_iterator<r   r<   r   
value_typer   r   r   __str__     z_ConstantIterator.__str__Nr1   r   r"   r#   r$   r   r   r   r   r   r         r   z#thrust/iterator/constant_iterator.hc                 C  B   |durt d|g}ddd |D }td| dt|jS )3Finds the first positions whose values differ.
    Nindex_type is not supportedrB   c                 S  rC   r   rD   r   r   r   r   r     rE   z*make_constant_iterator.<locals>.<listcomp>zthrust::make_constant_iterator(rF   )rH   rI   r
   r   r-   r   xir   rO   r   r   r   make_constant_iterator     
r   c                   @  r   )_CountingIteratorr1   r   c                 C  r   )Nzthrust::counting_iterator<r   r   r   r   r   r   r     r   z_CountingIterator.__str__Nr   r   r   r   r   r   r     r   r   z#thrust/iterator/counting_iterator.hc                 C  r   )r   Nr   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   z*make_counting_iterator.<locals>.<listcomp>zthrust::make_counting_iterator(rF   )rH   rI   r
   r   r-   r   r   r   r   make_counting_iterator  r   r   zthrust/mismatch.hc                 C  sl   t | t|| t|| |durtd||||g}ddd |D }td| dt|j|jgS )r   Nzpred option is not supportedrB   c                 S  rC   r   rD   r   r   r   r   r     rE   zmismatch.<locals>.<listcomp>zthrust::mismatch(rF   )	r0   r;   r>   rH   rI   r
   r   rn   r-   )r   r+   rh   ri   rj   predr   rO   r   r   r   mismatch  s   


r   zthrust/reduce.hc           	      C  s|   t | t| t|| |||g}|dur|| |j}n|jj}|dur+tdddd |D }td| d|S )z!Generalization of summation.
    NrA   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   zreduce.<locals>.<listcomp>zthrust::reduce(rF   )	r0   r5   r;   rp   r-   r<   rH   rI   r
   )	r   r+   rK   rL   r   rN   r   return_typerO   r   r   r   reduce  s   


r   c	                 C  s   t | t| t|| t| t| t| ||||||g}	|dur)td|dur1tdddd |	D }
td|
 dt|j|jgS )	z1Generalization of reduce to key-value pairs.
    Nrg   rA   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   z!reduce_by_key.<locals>.<listcomp>zthrust::reduce_by_key(rF   rm   )r   r+   
keys_first	keys_lastvalues_firstkeys_outputvalues_outputrk   rN   r   rO   r   r   r   reduce_by_key  s$   

r   zthrust/remove.hc                 C  ru   )ARemoves from the range all elements that are equal to value.
    rB   c                 S  rC   r   rD   r   r   r   r   r     rE   zremove.<locals>.<listcomp>zthrust::remove(rF   rv   re   r   r   r   remove  s   
r   c                 C  sX   t | t| t|| t| |||||g}ddd |D }td| d|jS )r   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   zremove_copy.<locals>.<listcomp>zthrust::remove_copy(rF   rv   )r   r+   rK   rL   rM   r[   r   rO   r   r   r   remove_copy  s   
r   zthrust/replace.hc                 C  sZ   t | t| t|| t|| |||||g}ddd |D }td| dtjS )KReplaces every element in the range equal to old_value with new_value.
    rB   c                 S  rC   r   rD   r   r   r   r   r     rE   zreplace.<locals>.<listcomp>zthrust::replace(rF   r0   r5   r;   rI   r
   r   rs   )r   r+   rK   rL   	old_value	new_valuer   rO   r   r   r   replace  r|   r   c           	      C  sd   t | t| t|| t| t|| ||||||g}ddd |D }td| d|jS )r   rB   c                 S  rC   r   rD   r   r   r   r   r   +  rE   z replace_copy.<locals>.<listcomp>zthrust::replace_copy(rF   rv   )	r   r+   rK   rL   rM   r   r   r   rO   r   r   r   replace_copy!  s   

r   zthrust/reverse.hc                 C  sL   t | t| t|| |||g}ddd |D }td| dtjS )Reverses a range.
    rB   c                 S  rC   r   rD   r   r   r   r   r   ;  rE   zreverse.<locals>.<listcomp>zthrust::reverse(rF   r   )r   r+   rK   rL   r   rO   r   r   r   reverse3  s   

r   c                 C  V   t | t| t|| t| ||||g}ddd |D }td| d|jS )r   rB   c                 S  rC   r   rD   r   r   r   r   r   H  rE   z reverse_copy.<locals>.<listcomp>zthrust::reverse_copy(rF   rv   rb   r   r   r   reverse_copy?  s   
r   zthrust/scatter.hc                 C  s`   t | t| t|| t| t| |||||g}ddd |D }td| dtjS )zMCopies elements from source range into an output range according to map.
    rB   c                 S  rC   r   rD   r   r   r   r   r   V  rE   zscatter.<locals>.<listcomp>zthrust::scatter(rF   r   )r   r+   rK   rL   maprM   r   rO   r   r   r   scatterL  s   
r   zthrust/sequence.hc                 C  sz   t | t| t|| |||g}|dur|| |dur)t|| || ddd |D }td| dtjS )z0Fills the range with a sequence of numbers.
    NrB   c                 S  rC   r   rD   r   r   r   r   r   j  rE   zsequence.<locals>.<listcomp>zthrust::sequence(rF   )r0   r5   r;   rp   rI   r
   r   rs   )r   r+   rK   rL   r   stepr   rO   r   r   r   sequence]  s   




r   zthrust/set_operations.hc           
      C  l   t | t|| t|| t| ||||||g}|dur"tdddd |D }	td|	 d|jS )zKConstructs a sorted range that is the set difference of sorted inputs.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r   {  rE   z"set_difference.<locals>.<listcomp>zthrust::set_difference(rF   r0   r;   r5   rH   rI   r
   r-   
r   r+   rh   ri   rj   last2rM   r\   r   rO   r   r   r   set_differencen     

r   c              	   C     t | t|| t|| t| t| t| t|	 |||||||||	g	}|
dur1tdddd |D }td| dt|j|	jgS )z>Constructs the key-value set difference of sorted inputs.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   z)set_difference_by_key.<locals>.<listcomp>zthrust::set_difference_by_key(rF   	r0   r;   r5   rH   rI   r
   r   rn   r-   r   r+   keys_first1
keys_last1keys_first2
keys_last2values_first1values_first2keys_resultvalues_resultr\   r   rO   r   r   r   set_difference_by_key  "   



r   c           
      C  r   )zMConstructs a sorted range that is the set intersection of sorted inputs.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   z$set_intersection.<locals>.<listcomp>zthrust::set_intersection(rF   r   r   r   r   r   set_intersection  r   r   c
                 C  s   t | t|| t|| t| t| t| ||||||||g}
|	dur,tdddd |
D }td| dt|j|jgS )z@Constructs the key-value set intersection of sorted inputs.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   z+set_intersection_by_key.<locals>.<listcomp>z thrust::set_intersection_by_key(rF   r   )r   r+   r   r   r   r   r   r   r   r\   r   rO   r   r   r   set_intersection_by_key  s    



r   c           
      C  r   )z@Constructs a sorted range that is the symmetric difference.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   z,set_symmetric_difference.<locals>.<listcomp>z!thrust::set_symmetric_difference(rF   r   r   r   r   r   set_symmetric_difference  r   r   c              	   C  r   )zDConstructs the key-value symmetric difference of sorted inputs.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   z3set_symmetric_difference_by_key.<locals>.<listcomp>z(thrust::set_symmetric_difference_by_key(rF   r   r   r   r   r   set_symmetric_difference_by_key  r   r   c           
      C  r   )zFConstructs a sorted range that is the set union of sorted inputs.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   zset_union.<locals>.<listcomp>zthrust::set_union(rF   r   r   r   r   r   	set_union  r   r   c              	   C  r   )z5Constructs the key-value union of sorted inputs.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   z$set_union_by_key.<locals>.<listcomp>zthrust::set_union_by_key(rF   r   r   r   r   r   set_union_by_key  r   r   c                 C  r   )>Sorts the elements in [first, last) into ascending order.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r   $  rE   zsort.<locals>.<listcomp>zthrust::sort(rF   r0   r5   r;   rH   rI   r
   r   rs   r   r   r   r   sort     

r   c                 C  f   t | t| t|| t| |durtd||||g}ddd |D }td| dtjS )Performs key-value sort.
    NrW   rB   c                 S  rC   r   rD   r   r   r   r   r   5  rE   zsort_by_key.<locals>.<listcomp>zthrust::sort_by_key(rF   r   r   r+   r   r   r   r\   r   rO   r   r   r   sort_by_key(     
r   c                 C  r   )r   NrW   rB   c                 S  rC   r   rD   r   r   r   r   r   H  rE   zstable_sort.<locals>.<listcomp>zthrust::stable_sort(rF   r   r   r   r   r   stable_sort=  r   r   c                 C  r   )r   NrW   rB   c                 S  rC   r   rD   r   r   r   r   r   Y  rE   z&stable_sort_by_key.<locals>.<listcomp>zthrust::stable_sort_by_key(rF   r   r   r   r   r   stable_sort_by_keyL  r   r   zthrust/swap.hc                 C  r   )z-Swaps each of the elements in the range.
    rB   c                 S  rC   r   rD   r   r   r   r   r   g  rE   zswap_ranges.<locals>.<listcomp>zthrust::swap_ranges(rF   rv   )r   r+   rh   ri   rj   r   rO   r   r   r   swap_ranges]  s   
r   zthrust/unique.hc                 C  r   )z4Removes all but the first element of the group.
    Nrg   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   zunique.<locals>.<listcomp>zthrust::unique(rF   r   )r   r+   rK   rL   rk   r   rO   r   r   r   uniquew  r   r   c                 C  sr   t | t| t|| t| ||||g}|durtdddd |D }td| dt|j|jgS )zUniques key-value pairs.
    Nrg   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   z!unique_by_key.<locals>.<listcomp>zthrust::unique_by_key(rF   rm   )r   r+   r   r   r   rk   r   rO   r   r   r   unique_by_key  s   

r   c                 G  r   )rR   rS   rT   r   NrU   rV   z.Invalid number of inputs of thrust.upper_boundrW   rB   c                 S  rC   r   rD   r   r   r   r   r     rE   zupper_bound.<locals>.<listcomp>zthrust::upper_bound(rF   r   rZ   r   r   r   upper_bound  ra   r   )r+   r
   )r   r
   r1   r2   )r   r
   r6   r
   r1   r2   )N)NN)NNN)D
__future__r   	cupyx.jitr   r   cupyx.jit._internal_typesr   r
   r)   TypeBaser*   r0   r5   r;   r>   r?   hostdeviceseqrP   r`   rc   rf   rl   ro   rq   rr   rt   rw   r{   r}   r~   r   r   r   r   r4   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   r   r   <module>   s    




















