
    uki                        d 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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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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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 m0Z0 dd+l m1Z1 dd,l2m3Z3 dd-l2m4Z4 dd.l2m5Z5 dd/l2m6Z6 dd0l2m7Z7 dd1l2m8Z8 dd2l9m:Z: dd3l9m;Z; dd4l9m<Z< dd5l9m=Z= dd6l>m?Z@ dd7lAmBZB ddlAmZC dd8lDmEZE dd9lDmFZF dd:lDmGZG dd;lDmHZH ej                  ZIej                  ZJej                  ZKej                  ZLej                  ZMej                  ZNej                  ZOej                  ZPej                  ZQej                  ZRej                  ZSd<e@j                  fd=eCj                  fd>ZVej                  re@j                  ZTeCj                  ZU[[@[Cydd?lXmYZZ  eZe[eV      Z\[Z[[@[Cy)@zMosaic-specific Pallas APIs.    N)core)create_tensorcore_mesh)dma_semaphore)GridDimensionSemantics)
KernelType)PrefetchScalarGridSpec)SemaphoreType)SideEffectType)MemorySpace)CompilerParams)	sync_copy)core_barrier)run_on_first_core)InterpretParams)force_tpu_interpret_mode)reset_tpu_interpret_mode_state)set_tpu_interpret_mode)LoweringException)BufferedRef)BufferedRefBase)emit_pipeline)emit_pipeline_with_allocations)get_pipeline_schedule)make_pipeline_allocations)
async_copy)async_remote_copy)bitcast)get_barrier_semaphore)load)make_async_copy)make_async_remote_copy)pack_elementwise)prng_random_bits)	prng_seed)repeat)roll)stochastic_round)store)touch)unpack_elementwise)with_memory_space_constraint)sample_block)stateful_bernoulli)stateful_bits)stateful_normal)stateful_uniform)to_pallas_key)ChipVersion)get_tpu_info)is_tpu_device)TpuInfo)
primitives)	semaphore)DeviceIdType)semaphore_read)semaphore_signal)semaphore_waitz0pltpu.delay is deprecated, use pl.delay instead.z,pltpu.ANY is deprecated, use pl.ANY instead.)delayANY)deprecation_getattr)]__doc__typingjax._src.pallas.mosaicr   jax._src.pallas.mosaic.corer   r   r   r   r   r	   r
   r   r   jax._src.pallas.mosaic.helpersr   r   r   6jax._src.pallas.mosaic.interpret.interpret_pallas_callr   r   r   r   jax._src.pallas.mosaic.loweringr   jax._src.pallas.mosaic.pipeliner   r   r   r   r   r   !jax._src.pallas.mosaic.primitivesr   r   r   r   r   r    r!   r"   r#   r$   r%   r&   r'   r(   r)   r*   r+   jax._src.pallas.mosaic.randomr,   r-   r.   r/   r0   r1   jax._src.pallas.mosaic.tpu_infor2   r3   r4   r5   jax._src.pallasr6   pl_primitivesjax._src.pallas.corer7   GeneralMemorySpacejax._src.pallas.primitivesr8   r9   r:   r;   PARALLELCORE_PARALLELSUBCORE_PARALLEL	ARBITRARYCMEMSMEMVMEMVMEM_SHAREDHBMHOST	SEMAPHOREr<   r=   _deprecationsTYPE_CHECKINGjax._src.deprecationsr>   _deprecation_getattr__name____getattr__     V/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/experimental/pallas/tpu.py<module>rc      s   #  / X F X @ X F H B H A G Q e w D s R F N J l Z b F T @ \ : P ^ R R D > : R < < V j F R H L N H F H J > 8 7 B C G K G!**&44):: ",,	%%oo!!	
 9 	7 



%#
  P$X}=+
ra   