
    uki
                     f    d Z ddlZddlZddlmZ ddlmZ ddlm	Z
 ddlmZ d ZdefdZdefd	Zy)
zHelpers for Pallas TPU kernels.    N)helpers)
primitives)corec                      t         j                  j                         syt        j                  t
        j                  t        j                  j                  d             fd       }y)zCopies a PyTree of Refs to another PyTree of Refs.

  Args:
    src_ref: A Pytree of source Refs/TransformedRefs.
    dst_ref: A Pytree of destination Refs/TransformedRefs.
  N )semc                       fd}t         j                  j                  t        j                  |d             t         j                  j                  t        j                  |d             y )Nc                     t        j                  ||      }| dk(  r|j                          y | dk(  r|j                          y t	        d|        )NstartwaitzUnknown action: )plm_primitivesmake_async_copyr   r   
ValueError)actionsrc_refdst_ref
descriptorr   s       Y/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/_src/pallas/mosaic/helpers.py_copy_start_or_waitz1sync_copy.<locals>._.<locals>._copy_start_or_wait(   sP    !11'7CHj	7	V+F8455    r   r   )jaxtreemap	functoolspartial)r   r   r   r   s   ` r   _zsync_copy.<locals>._$   sZ    6 HHLL-w7
 HHLL-v6r   )
r   r   leavesr   r   pl_primitives
run_scopedtpu_coreSemaphoreTypeDMA)r   r   r   s   `` r   	sync_copyr#      sQ     
	!
H$:$:$>$>r$Br   core_axis_namec                 `     t         j                  j                         }|dk(  rd S  fd}|S )z2Runs a function on the first core in a given axis.   c                      |        S Nr   fs    r   <lambda>z#run_on_first_core.<locals>.<lambda>A   s    QS r   c                      t         j                  j                        }t        j                  |dk(        t        j                          fd              }y )Nr   c                               S r(   r   r)   s   r   r   z-run_on_first_core.<locals>.wrapped.<locals>._F   s     Sjr   )r   lax
axis_index
pl_helperswhenr   wraps)r*   core_idr   r$   s   `  r   wrappedz"run_on_first_core.<locals>.wrappedC   sG    gg  0G__W\"__Q  #r   )r   r.   	axis_size)r$   	num_coresr4   s   `  r   run_on_first_corer7   =   s0    gg/)!^ 
.r   c                     t         j                  j                  |      t         j                  j                  |      t	        j
                  dkD         fd       }y)z'Synchronizes all cores in a given axis.r&   c                      t        j                  d      5  fd} t              D ]
  } | |        t        j                  dz
         d d d        y # 1 sw Y   y xY w)N
sync_coresc                 J     t        j                   k7         fd       }y )Nc                  6    t        j                  d        y )Nr&   )
core_index)r   semaphore_signal)ir   s   r   r   z7core_barrier.<locals>._.<locals>.signal_core.<locals>._Y   s    

(
(aA
>r   )r0   r1   )r?   r   r3   r   s   ` r   signal_corez,core_barrier.<locals>._.<locals>.signal_coreW   s"    	A	&	? 
'	?r   r&   )r   named_scoperanger   semaphore_wait)r@   r?   r3   r6   r   s     r   r   zcore_barrier.<locals>._S   sZ    		& 
7? Y !A""3	A6
7 
7 
7s   8AA!N)r   r.   r5   r/   r0   r1   )r   r$   r   r3   r6   s   `  @@r   core_barrierrD   N   sJ    gg/)GG~.'??9q=!7 "7r   )__doc__r   r   jax._src.pallasr   r0   r   r   jax._src.pallas.mosaicr   r    r   r#   strr7   rD   r   r   r   <module>rI      s:    &  
 1 7 3 ?!Hc "7 7r   