
    9i                     d    S r SSKJr  \R                  SS j5       r\R                  SS j5       rg)a  
Grid Dependency Control (GDC) is a mechanism used when enabling programmatic dependent launch to launch and
synchronize grids. These APIs expose GDC to the programmer.

Programmatic dependent launch is supported on SM90 (Hopper) and beyond.
For PTX reference on grid dependency control see https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-griddepcontrol.
    )coreNc           
      T    [         R                  " SS/ [         R                  SSU S9  g)a>  
GDC wait is a blocking instruction that waits for all instructions in a prior kernel to complete before continuing.
This ensures all memory operations happening before the wait is visible to instructions after it,
e.g. if the prior kernel writes to address "x" the new values will be visible in this kernel after the wait.

This instruction is also safe to execute when programmatic dependent launch is disabled.

See https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-griddepcontrol for more details.
z griddepcontrol.wait; // dummy $0=rF   dtypeis_purepack	_semanticNr   inline_asm_elementwiseint32r   s    ^/var/www/html/land-doc-ocr/venv/lib/python3.13/site-packages/triton/language/extra/cuda/gdc.pygdc_waitr      s+     	 BD"TXT^T^hmtu*35    c           
      T    [         R                  " SS/ [         R                  SSU S9  g)ac  
This operation when launched with programmatic dependent launch signals that
the next program may launch once all programs in the current kernel
call this function or complete.

Repeated calls to this function have no effect past the first call, and the first call should be
treated by the programmer as a hint to the runtime system to launch the next kernel.

This instruction is also safe to execute when programmatic dependent launch is disabled.

See https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-griddepcontrol for more details.
z-griddepcontrol.launch_dependents; // dummy $0r   Fr   r   Nr   r   s    r   gdc_launch_dependentsr      s,     	 OQUWYaeakak(-ALr   )N)__doc__triton.languager   externr   r    r   r   <module>r      sA    ! 5 5 L Lr   