
    xKg                         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mZ  G d d      Zd	efd
Ze
d        Z eed      d        Ze
d        Z eedd      d        Zy)    )types)overloadoverload_method)	signature)	nvvmutils)	intrinsic)
grid_group	GridGroupc                       e Zd ZdZddZy)r
   z0A cooperative group representing the entire gridNc                       y)zSynchronize this grid groupN r       Q/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/cg.pysynczGridGroup.sync   s    r   )returnN)__name__
__module____qualname____doc__r   r   r   r   r
   r
   	   s
    :*r   r
   r   c                      t               S )zGet the current grid group.)r
   r   r   r   	this_gridr      s
    ;r   c                 .    t        t              }d }||fS )Nc                     | j                  t        j                  d      }|j                  }|j	                  t        j                  |      |f      S )N   )get_constantr   int32modulecallr    declare_cudaCGGetIntrinsicHandle)contextbuildersigargsonemods         r   codegenz_this_grid.<locals>.codegen   sE    ""5;;2nn||66s;F 	r   )r   r	   )	typingctxr"   r&   s      r   
_this_gridr(      s    
J
C <r   cuda)targetc                      d } | S )Nc                      t               S N)r(   r   r   r   implz_ol_this_grid.<locals>.impl%   s
    |r   r   )r.   s    r   _ol_this_gridr/   #   s     Kr   c                 D    t        t        j                  |      }d }||fS )Nc                     | j                  t        j                  d      }|j                  }|j	                  t        j                  |      g ||      S )Nr   )r   r   r   r   r   r   declare_cudaCGSynchronize)r    r!   r"   r#   flagsr%   s         r   r&   z!_grid_group_sync.<locals>.codegen/   sK    $$U[[!4nn||//4dNEN 	r   )r   r   r   )r'   groupr"   r&   s       r   _grid_group_syncr5   +   s#    
EKK
'C <r   r   c                     d }|S )Nc                     t        |       S r-   )r5   )r4   s    r   r.   z!_ol_grid_group_sync.<locals>.impl;   s    &&r   r   )r4   r.   s     r   _ol_grid_group_syncr8   9   s    ' Kr   N)
numba.corer   numba.core.extendingr   r   numba.core.typingr   
numba.cudar   numba.cuda.extendingr   numba.cuda.typesr	   r
   GridGroupClassr   r(   r/   r5   r8   r   r   r   <module>r@      s     : '   * D* *9 
 
 
 
)F# $ 
 
 7 8r   