
    uki7                        d 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(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.l(m2Z2 dd/l(m3Z3 dd0l(m4Z4 dd1l(m5Z5 dd2l(m6Z6 dd3l(m7Z7 dd4l(m8Z8 dd5l(m9Z9 dd6l(m:Z: dd7l(m;Z; dd8l(m<Z< dd9l(m=Z= dd:l(m>Z> dd;l(m?Z? dd<l(m@Z@ dd=l(mAZA dd>l(mBZB dd?l(mCZC dd@l(mDZD ddAl(mEZE ddBl(mFZF ddCl(mGZG ddDl(mHZH ddElImJZJ ddFlKmLZL ddGlMmNZN ddHlMmOZO e
j                  ZPe
j                  ZQe
j                  ZRe
j                  ZSyI)JzExperimental GPU backend for Pallas targeting H100.

These APIs are highly unstable and can change weekly. Use at your own risk.
    )Barrier)	BlockSpec)ClusterBarrier)CompilerParams)kernel)Layout)layout_cast)MemoryRefTransform)MemorySpace)Mesh)multicast_ref)
PeerMemRef)RefUnion)
remote_ref)SemaphoreType)SwizzleTransform)TilingTransform)
TMEMLayout)transform_ref)transpose_ref)TransposeTransform)TryClusterCancelResult)unswizzle_ref)
untile_ref)WarpMesh)WGMMAAccumulatorRef)find_swizzle)format_tcgen05_sparse_metadata)nd_loop)
NDLoopInfo)planar_snake)dynamic_scheduling_loop)emit_pipeline)emit_pipeline_warp_specialized)PipelinePipeline)async_copy_scales_to_tmem)"async_copy_sparse_metadata_to_tmem)async_load_tmem)async_prefetch)async_store_tmem)barrier_arrive)barrier_wait)broadcasted_iota)commit_smem)commit_smem_to_gmem_group)commit_tmem)copy_gmem_to_smem)copy_smem_to_gmem)inline_mgpu)load)multimem_store)multimem_load_reduce)print_layout)query_cluster_cancel)RefType)semaphore_signal_multicast)semaphore_signal_parallel)SemaphoreSignal)set_max_registers)ShapeDtypeStruct)tcgen05_commit_arrive)tcgen05_mma)try_cluster_cancel)wait_load_tmem)wait_smem_to_gmem)wgmma)
wgmma_wait)as_torch_kernel)LoweringSemantics)
Replicated)TilingN)T__doc__jax._src.pallas.mosaic_gpu.corer   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   ACC"jax._src.pallas.mosaic_gpu.helpersr   r   r   r    r!   r"   #jax._src.pallas.mosaic_gpu.pipeliner#   r$   r%   %jax._src.pallas.mosaic_gpu.primitivesr&   r'   r(   r)   r*   r+   r,   r-   r.   r/   r0   r1   r2   r3   r4   r5   r6   r7   r8   r9   r:   r;   r<   r=   r>   r?   r@   rA   rB   rC   rD   rE    jax._src.pallas.mosaic_gpu.torchrF    jax.experimental.mosaic.gpu.corerG   ,jax.experimental.mosaic.gpu.fragmented_arrayrH   rI   GMEMSMEMTMEMREGS     ]/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/experimental/pallas/mosaic_gpu.py<module>rZ      s  
 ? B L L < < F T F 8 J D @ D J P N D J J T \ J D @ F V K o A G K a N p T h z T R V R N V L h L X X L > R ^ N ^ D j h T X V ` L Z R X @ J O S Q I rX   