
    VpfT                     2   d dl Z d dlZd dlZd dlZd dlZd dl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 d dlmZ 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  e            Ze Z!dZ"dZ# G d d          Z$d Z% ej&        d           G d d                      Z' ej&        d           G d d                      Z( ej&        d           G d d                      Z)d Z*d Z+d Z,de-de-d e fd!Z. ej/        d"          Z0e01                    d#            e0j2        d$             Z3d%ej4        fd&Z5e5ej6        e0<   d' Z7 ej/        d(          Z8de8_9        e8j2        d)             Z:d%ej4        fd*Z;e;ej6        e8<   d+ Z<d, Z=d- Z>dS ).    N)AnySequence)lax)core)	tree_util)tpu)lowering)
primitives)
split_listunzip2)ir)arith)func)PassManagera  
#define buf_readers(index, device, core) _buf_readers[(index)*(NDEVICE*NCORE) + (device)*NCORE + core]
#define buf_written(index, device, core) _buf_written[(index)*(NDEVICE*NCORE) + (device)*NCORE + core]
#define sems(index, device, core) _sems[(index)*(NDEVICE*NCORE) + (device)*NCORE + core]
#define barrier_sems(device, core) _barrier_sems[(device)*NCORE + core]

#ifndef NDMA
#define NDMA 2
#endif

mtype = { DMA };
chan dma_queue = [NDMA] of { mtype, int, int, int, int, int, int, int, int, int, int };
a  
active [NDMA] proctype DmaEngine() {
  int src_dev, src_core, src_sem, src_buf_base, src_buf_len;
  int dst_dev, dst_core, dst_sem, dst_buf_base, dst_buf_len;
  do
    :: skip;
       end: dma_queue?DMA(src_dev, src_core, src_sem, src_buf_base, src_buf_len, dst_dev, dst_core, dst_sem, dst_buf_base, dst_buf_len);
       d_step {
         printf("DMA read done: [%d, %d)@{%d, %d} (%d++)\n", src_buf_base, src_buf_base + src_buf_len, src_dev, src_core, src_sem);
         int i;
         for (i : src_buf_base .. src_buf_base + src_buf_len - 1) {
           buf_readers(i, src_dev, src_core)--;
         }
         sems(src_sem, src_dev, src_core)++;
       }  // Read complete
       d_step {
         printf("DMA write done: [%d, %d)@{%d, %d} (%d++)\n", dst_buf_base, dst_buf_base + dst_buf_len, dst_dev, dst_core, dst_sem);
         int i;
         for (i : dst_buf_base .. dst_buf_base + dst_buf_len - 1) {
           buf_written(i, dst_dev, dst_core)--;
         }
         sems(dst_sem, dst_dev, dst_core)++;
       }  // Write complete
  od
}
c            
           e Zd ZdZd Zdee         fdZdee         fdZde	de	fdZ
d	 Zd
 Zej        de	de	fd            Zefdej        defdZdej        defdZdedededee         de	f
dZdS )PrintCtx   c                 H   d| _         d| _        d| _        g | _        t	          j                    | _        i | _        t          d t          t          |                    D                       | _        d| _        |                     d d           |                     d d           |r1|                     d dd                    | j                              |                     d	d
          5  d}g }t!          t#          t%          |                              D ]R\  }}|                    |                     d d| d| d|                      |                     d| d|           }S|                     d d|            d d d            d S # 1 swxY w Y   d S )N   r   c              3       K   | ]	}d | V  
dS )pidN ).0is     c/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/_src/pallas/mosaic/verification.py	<genexpr>z$PrintCtx.__init__.<locals>.<genexpr>^   s(      MM1Y1YYMMMMMM    dev_idzint core_id = 0z
int dev_idzint , d_step {}_pidr    = z % intz / z	dev_id = )levelnum_semaphoresnum_bufferslocals	itertoolscountcounterenvtuplerangelenprogram_ids	device_idemitjoinblockreversedlist	enumerateappend)selfiteration_boundsidxr0   r   bs         r   __init__zPrintCtx.__init__W   s   DJDDDK?$$DL*,DHMMc:J6K6K0L0LMMMMMDDN 	IId%&&&IIdL!!! <
ii:TYYt'788::;;;	J	$	$ ) )ck4	*: ; ;<<== / /$!Q499T+B+B+Bs+B+Bq+B+BCCDDDii#~~!~~..
ii'#''((() ) ) ) ) ) ) ) ) ) ) ) ) ) ) ) ) )s   0BFFFshapec                     d}|r|d         | j         k    r|d         }| j        }| xj        |z  c_        t          ||          S )Nr   r   )MAX_REF_UNROLLr'   GlobalRefModel)r9   r>   slotsbases       r   emit_global_refzPrintCtx.emit_global_refp   sU    E qT000AheD$&&&r   c                 x    t          j        |          }| j        }| xj        |z  c_        t          ||          S N)mathprodr&   GlobalSemaphoreModel)r9   r>   r*   rC   s       r   emit_global_semaphore_refz"PrintCtx.emit_global_semaphore_refx   s=    IeED5 e,,,r   textreturnc                 <    t          j        |d| j        z            S )Nz  )textwrapindentr%   )r9   rK   s     r   _indentzPrintCtx._indent~   s    ?4
!2333r   c                     d }|.dt          t          | j                            z   }| d| d| }| j                            |                     |          dz              |S )Nl r#   z;
)strnextr+   r(   r8   rP   )r9   tyexprnames       r   r2   zPrintCtx.emit   so    D	~3tDL))***d$$T$$d$$dKt||D))E1222Kr   c                 h    | j                             |                     d| d                     d S )Nz/* z */
)r(   r8   rP   )r9   comments     r   rZ   zPrintCtx.comment   s6    Kt||$8'$8$8$899:::::r   beginendc              #     K   | j                             |                     |          dz              | xj        dz  c_        d V  | xj        dz  c_        | j                             |                     |          dz              d S )N
r   )r(   r8   rP   r%   )r9   r[   r\   s      r   r4   zPrintCtx.block   s      Kt||E**T1222JJ!OJJ	EEEJJ!OJJKt||C((4/00000r   valuedefaultc                 d    |t           u r| j        |         S | j                            ||          S rF   )_UNSPECIFIEDr,   get)r9   r_   r`   s      r   rc   zPrintCtx.get   s/    ,Xe_X\\%)))r   model_valuec                     || j         |<   d S rF   )r,   )r9   r_   rd   s      r   setzPrintCtx.set   s    !DHUOOOr   has_barrier_semsnum_devices	num_coresparallel_iteration_boundsc                    t          j                    }|                    d| d           |                    d           |                    d| j         d           |                    d| j         d           |                    d| j         d           |r|                    d           |                    t
                     |                    d           t          j        |          }|                    d	| d
           | j        D ]}|                    |           |                    d           |                    t                     |
                                S )Nz#define NDEVICE r^   z#define NCORE 1
zbyte _buf_readers[z*NDEVICE*NCORE] = 0;
zbool _buf_written[zbyte _sems[z'byte _barrier_sems[NDEVICE*NCORE] = 0;
zactive [NDEVICE*z] proctype Kernel() {
z}
)ioStringIOwriter'   r&   PREAMBLErG   rH   r(   DMA_PROCESSgetvalue)r9   rg   rh   ri   rj   resultparallel_threadsrR   s           r   	get_modelzPrintCtx.get_model   sh    []]F
LL3K333444
LL$%%%
LLNd&6NNNOOO
LLNd&6NNNOOO
LLJt2JJJKKK ?ll=>>>
LL
LLy!:;;
LLN$4NNNOOO[  ll1oooo
LL
LL??r   N)__name__
__module____qualname__r@   r=   r   r$   rD   rJ   rT   rP   r2   rZ   
contextlibcontextmanagerr4   rb   r   Valuer   rc   rf   boolrt   r   r   r   r   r   T   sr       .) ) )2'8C= ' ' ' '-Xc] - - - -4# 4# 4 4 4 4  ; ; ; 1 13 1 1 1 1 1= * *rx *# * * * *"rx "c " " " "  	
 "*#      r   r   c                     | d d g} nt          |           } | d         d| d<   | d         d| d<   t          |           S )Nr   r   r   core_id)r6   r-   locations    r   resolve_locationr      sP    d|HHH~~Ha[HQKa[HQK	xr   T)frozenc                   4    e Zd ZU dZeed<   eed<   d Zd ZdS )rA   zA model of a memory reference.

  When a reference has a small leading dimension, it might be represented by
  multiple slots in the reference array. Its region starts at base (that can be
  dynamic) and has the given length (always static).
  rC   lengthc                 p     t          |          \   fdt           j                  D             S )Nc                 8    g | ]}d j          d| d d d	S )zbuf_readers( + r   )rC   r   r   r   devr9   s     r   
<listcomp>z-GlobalRefModel.readers_at.<locals>.<listcomp>   ?    YYY<49<<<<c<<T<<<YYYr   r   r.   r   r9   r   r   r   s   ` @@r   
readers_atzGlobalRefModel.readers_at   @     **ICYYYYYYeDKFXFXYYYYr   c                 p     t          |          \   fdt           j                  D             S )Nc                 8    g | ]}d j          d| d d d	S )zbuf_written(r   r   r   r   r   s     r   r   z-GlobalRefModel.written_at.<locals>.<listcomp>   r   r   r   r   s   ` @@r   
written_atzGlobalRefModel.written_at   r   r   N)	ru   rv   rw   __doc__r   __annotations__r$   r   r   r   r   r   rA   rA      sZ           	)))
+++Z Z ZZ Z Z Z Zr   rA   c                   .    e Zd ZU dZeed<   eed<   d ZdS )rI   zA model of a semaphore reference.

  Semaphore arrays are always fully unrolled and are represented by a contiguous
  subset of the global semaphore array.
  rC   r   c                 H    t          |          \  }}d| j         d| d| dS )Nzsems(r   r   )r   rC   r9   r   r   r   s       r   atzGlobalSemaphoreModel.at   s6     **IC.49....t....r   N)ru   rv   rw   r   r   r   r$   r   r   r   r   rI   rI      sC          
 	)))
+++/ / / / /r   rI   c                       e Zd Zd ZdS )GlobalBarrierSemaphoreModelc                 8    t          |          \  }}d| d| dS )Nzbarrier_sems(r   r   )r   r   s       r   r   zGlobalBarrierSemaphoreModel.at   s,     **IC)3))$))))r   N)ru   rv   rw   r   r   r   r   r   r      s#        * * * * *r   r   c                 :    |j         xdk    r t           |j                   d S xdk    r  j        S xdk    rW t          j                            |j        j                  r+t          t	          j
        |j                  j                  S d S xdk    r% t                               |j        d                                |j        d           f          }                     |j                  }|                    |          }                     |j                  }t          |t$                    r-                     d d|d          d|d          d| d	           n4                     d d
|j         d|d          d|d          d| d		                                d d| d| d| d           d S xdk    r                      |j                  }|                    d           }                     |j                  }                     d d| d| d| d| d| d           t          |t$                    r                     d d| d	           d S                      d d|j         d| d	           d S xdk    r t                               |j        d                                |j        d           f          }                     |j                  }                     |j                  }                     |j                  }	                     |j                  }
d                    d |                    d           D                       }d                    d t7          j        |	                    |          |	                    |                    D                       }                     d d|j         d|j         d|j         d|	j         d|	j         d|	j         d|d          d|d          d	                                dd          5                       d d| d                                 d d| d!           |                    d           D ]L}                     d | d"           |	                    |          D ]}                     d | d#           M	 d d d            n# 1 swxY w Y                        d d$|j         d|j         d|j         d|d          d|d          d|
j         d|	j         d|	j         d	           d S xd%k    rr                      |j                  }|                    d           }                     d d| d&| d| d'                                d d(|j         d	           d S xd)k    r t%                      S xd*k    r                      |j         d           }|tB          S t	          j"        |j         j                  j#        }t	          j"        |j        j                  j#        }t          j$        %                                }t          |tL                    rd+}tO          ||          D ]%\  }}|dk    r||k    r|rtQ          d,          d-}&g }d}|d d d.         D ]}|)                    |           ||z  }tU          |          } fd/|j+        D             }d                    d0 tO          ||          D                       }tM          |j         d| tY          j-        |          1          S t          |t\                    sJ                      |j+        d         d           }|r%|d         |k    s|d         |k    s|j        dk    s||S t]          |j         d| |d                   S xd2k    r'                      |j/        d           }|tB          n|S xd3k    r'                      |j        d           }|tB          n|S xd4k    r ta           d5d6g|j1        R  S xd7k    r ta           d5d8g|j1        R  S xd9k    r ta           d5d:g|j1        R  S xd;k    r ta           d5d<g|j1        R  S xd=k    r ta           d5d>g|j1        R  S xd?k    r |j2        j        xtf          j4        j5        k    r ta           d@dAg|j1        R  S xtf          j4        j6        k    r ta           d@dBg|j1        R  S xtf          j4        j7        k    r ta           d@dCg|j1        R  S xtf          j4        j8        k    r ta           d@dDg|j1        R  S xtf          j4        j9        k    r ta           d@dEg|j1        R  S tf          j4        j:        k    rta           d@dFg|j1        R  S ta           d@d>g|j1        R  S xdGk    r"  ;                    |j<        j                   d S xd3k    r!                      |j        tB                    S xdHk    r g }|j1        D ]A}                     |d           x}t{          dI|           |)                    |           B                     dd          5  |D ]3}|                    d           D ]}                     d | d"           4	 d d d            n# 1 swxY w Y                        dd          5  |D ]3}|                    d           D ]}                     d | dJ           4	 d d d            d S # 1 swxY w Y   d S xdKk    r:  fdL|j>        D             }|j?        |j@        |jA        f} t           j        |           x\  }!}"}#}$tO          |$|           D ]\  }}%|t{          dM|%                                d5|!          }&                     dNdO          5                       d dP|& dQ|" d            C                    |jD        |&           tO          ||jE        d-R          D ]\  }'}(|' C                    |(|'           t           |j                   |j        jF        t          |j        jF                  dz
           })|)j1        }*                     dd          5  tO          ||*d-R          D ]5\  }'}+|'.                     d |' d                     |+                      6                     d |& d|& d|#            d d d            n# 1 swxY w Y                        d dS           d d d            n# 1 swxY w Y   t          |          dk    r|d         S t          |          S dTk    r|jI        rtP                               |jJ        d           x},t{          dU|jJ                                        dVdW          5                       d dX|, d	           t           |jK                   |jL        d         jM        r,                     d dY           t           |jN                   n                     d dZ           d d d            d S # 1 swxY w Y   d S 	 |jL        stB          S tQ          d[          )\Nz
tpu.regionztpu.device_idzarith.constantztpu.sem_signalz+printf("Signal: BARRIER@{%d, %d} += %d\n", r   r   r   r   z&printf("Signal: %d@{%d, %d} += %d\n", z	d_step { r#   r   z }ztpu.sem_waitr~   z	atomic { z >= z; z - z%printf("Wait done: BARRIER -= %d\n", z printf("Wait done: %d -= %d\n", ztpu.enqueue_dmaz
    && c              3       K   | ]	}|d z   V  
dS z == 0Nr   r   
is_writtens     r   r   z_print_op.<locals>.<genexpr>  s'      %b%bzj7&:%b%b%b%b%b%br   c              3       K   | ]	}|d z   V  
dS r   r   r   s     r   r   z_print_op.<locals>.<genexpr>  s;       $ $ w
$ $ $ $ $ $r   z8printf("DMA: [%d, %d)@{%d, %d} -> [%d, %d)@{%d, %d}\n", z, dev_id, core_id, r    r!   zassert(z );  // Source is not written to.z);  // Destination is unused.z++z = 1zdma_queue!DMA(dev_id, core_id, ztpu.wait_dmaz >= 1; z - 1 }zprintf("Awaited DMA: %d\n", ztpu.sem_barrierztpu.memref_sliceFz)Non-contiguous slices of semaphore arraysTc                 :    g | ]}                     |          S r   )rc   )r   r;   ctxs     r   r   z_print_op.<locals>.<listcomp>L  s#    777C3773<<777r   c              3   *   K   | ]\  }}| d | V  dS )z * Nr   )r   r;   ss      r   r   z_print_op.<locals>.<genexpr>M  s0      "V"Vfc1c>>a>>"V"V"V"V"V"Vr   )rC   r   ztpu.memref_squeezeztpu.assume_multiplez
arith.addir$   +z
arith.subi-z
arith.muli*zarith.remsi%zarith.divsi/z
arith.cmpir{   z==z!=<z<=>z>=ztpu.trace_startverification.pretendzCould not model the read of z--zscf.forc                     g | ]Q}t           j                            |j                  r)                    d                     |                    ndRS )r$   N)r   IntegerType
isinstancetyper2   rc   )r   argr   s     r   r   z_print_op.<locals>.<listcomp>  sb         ^&&sx00;#((5#''#,,
'
'
'6:  r   z$Could not model loop bound or step: doodz:: z < strictz:: else -> breakzscf.ifz"Could not model branch condition: iffiz:: (z:: elsez:: else -> skipz Must handle all ops with regions)OrX   _print_blockbodyr1   r   r   r   rr   r   rT   IntegerAttrr_   r   rc   r}   	semaphorer   amountr   r2   rC   sourcesource_semaphoretargettarget_semaphorer3   r   r)   chainr   r   r4   mem_refNotImplemented
MemRefTyper>   
ShapedTypeget_dynamic_sizerI   zipNotImplementedErrorr8   r5   base_idxrG   rH   rA   inputbin_opoperands	predicater   CmpIPredicateeqnesltslesgtsgerZ   message
ValueErrorinitArgs
lowerBound
upperBoundstepmaprf   induction_variableinner_iter_args
operationsr/   r-   results	condition
then_blockregionsblocks
else_block)-r   opr   	sem_modelsemr   dst_locationsrcsrc_semdstdst_semsrc_readonly
dst_unusedrwrr   	src_shape	dst_shapedynamicseen_nontrivial_unequalr   dstridesstrideindiceslinear_offset	major_idx	read_refsomodelloccarrysboundslowerupperr   bound_modelsvinduction_varcr   
terminator
new_carrysnewr   s-   `                                            r   	_print_opr     sT   
	3     	]							"	"29>	2	2 2>"(++1222					!3772<#>#>
TX@Y@Y"Z[[h''",''iLL""cwwry!!f	I:	;	; ~oQRooW_`aWbooflooopppp|9>||U]^_U`||dlmndo||sy|||}}}	hht=#==#==&===>>>>>	''",''iLL$L''cwwry!!f	hhtP#PP6PPSPPSPPVPPPQQQ	I:	;	; XIIIIJJJJJV9>VVVVVVWWWWW					%swwr|T'B'BCGGBJX\D]D]&^__lGGBIc+,,gGGBIc+,,g %%%b%bS^^\`MaMa%b%b%bbbl## $ $%Onn\**CNN<,H,H $ $ $  j 
hh
!h! !(! !'*z! !h! !(! !'*z! !5A!_! ! 1o! ! !   99Z%% ' 'OOOOPPPJJJJKKK%% 	' 	'A
((4A
"
"
">>,// ' 'aHHTa:::&&&&'	'' ' ' ' ' ' ' ' ' ' ' ' ' ' ' 
hh
8GL 8 8CH 8 8j8 8(O8 8/;A8 8l8 8!h8 8*-*8 8 8     
''",''iLL$L''c	hhtB#BBcBBcBBBCCC	hhtFY^FFFGGGGG					(***					wwrz4((f	-
006i-	//5i..00g 
F0	1	1 L"'	9-- 	+ 	+DAq!VV!VV& U'(STTT&*#44R4 	 	A
..
 
 
 
A+&&7##77772;777

"V"VGW@U@U"V"V"VVV#K33M33DIi<P<P
 
 
 	
 &.11111GGBKND11	 	Yq\W44	!8O8O}!!Y%6-<<<<ilKKK					wwrx&&f%~^^69					wwrx&&f%~^^69	C2bk2222	C2bk2222	C2bk2222	C2bk2222	C2bk2222	L#U ####VT8BK888
8#U ####VT8BK888
8$U $$$$VS72;777
7$U $$$$VT8BK888
8$U $$$$VS72;777
7 $$$VT8BK888
8C3r{3333						kk"*"#####					WWRX~...					i{    !WWQ%%%E.=!==>>
>99Z%% ' ' 	' 	'A\\$'' ' 'cHHTc:::&&&&'	'' ' ' ' ' ' ' ' ' ' ' ' ' ' ' 99Z%% ' ' 	' 	'A\\$'' ' 'cHHTc:::&&&&'	'' ' ' ' ' ' ' ' ' ' ' ' ' ' ' ' ' ' 
    [  f
 r}bg6f*-cgv*>*>>eUD<,// G G(%=E!EEFF
F hhue,,m99T4   + +8]88u888999%}555&""4TBBB 	 	FAs]GGCOOOS"'"""W'BG,>(?(?!(CD
(
YYz3'' 	H 	HFJt<<< 6 6fa}hht44cggcll44555
((4MFFmFFFF
G
G
G		H 	H 	H 	H 	H 	H 	H 	H 	H 	H 	H 	H 	H 	H 	H
 	)***+ + + + + + + + + + + + + + + 
V		ayV}}		 "!!wwr|T22
2)	;LblLLMMM99T4   , ,*i***+++S"-(((:a= 	,
((4
#
#
#
sBM
*
*
*
*
((4*
+
+
+, , , , , , , , , , , , , , , , , , 
Z  BCCCs   5BSS S7kkk17l66l:=l:#Cu/A't"u"t&	&u)t&	*uuu!By22y69y6c                     |                      |d           }|                      |d           }||t          S |                     || d| d|           S )NrS   )rc   r   r2   )r   	result_tyr   lhsrhss        r   r   r     sa    T#T#[CK	)00b00300	1	11r   c                 .   |D ]}	 t          | |          }n%# t          $ r}t          d|           |d }~ww xY w|t          u rC|j        s|J Ot          |j                  dk    rt          |          |                     |j        |           d S )NzFailed to print op: r   )	r  	ExceptionRuntimeErrorr   r   r/   r   rf   rr   )r   r4   r   r   es        r   r   r     s     " "b=#r""gg = = =444551<=.  : "____	RZ1		###	ggbi!!!!" "s   
949rh   num_cores_per_devicerL   c                    | j         5  t          j        | j                  \  }}t          j                            | j                            d                    } ddg}t          j        dd	                    |           d          }|
                    | j                   t          j                            d          }| j        D ]}t          |d	d           |k    r nt          d
          t!          |t"          j                  sJ d}	d|j        v rt	          j        |j        d                   }	t          j                                        t/          fd|	D                       rt          d          t	          j        |j        d                   }
t          j                            d          t/          fd|
D                       rt5          d          d}d|j        v r$t	          j        |j        d                   j        }|j        \  }t;          |	          }t          j                            d          }t          j                            d          }t?          |j         tC          |	          |g          \  }}}tE          ||j#        d          D ]\  }}|$                    ||           ~|D ]}t          j%                            |j&                  rt	          j%        |j&                  }|j'        |k    s|j'        |k    r/|$                    ||(                    |j)                             |$                    ||*                    |j)                             tW          ||           |,                    ||||	          cd d d            S # 1 swxY w Y   d S )NT)binarycanonicalizecsezbuiltin.module(,r   mainrX   zNo main function foundr   r:   c              3   $   K   | ]
}|k    V  d S rF   r   )r   r<   r   s     r   r   z'export_promela_model.<locals>.<genexpr>  s'      44aQ'\444444r   z&Dynamic iteration bounds not supporteddimension_semanticsz"#tpu.dimension_semantics<parallel>c              3   $   K   | ]
}|k    V  d S rF   r   )r   r   parallels     r   r   z'export_promela_model.<locals>.<genexpr>  s'      88qQ(]888888r   z%Non-parallel dimensions not supportedr   scalar_prefetchz!tpu.semaphorez!tpu.dma_semaphorer   )-contextr   private_has_communication	operationr   Moduleparseget_asmr   r3   run
StringAttrrc   r   getattrr   r   r   FuncOp
attributesDenseI64ArrayAttrr   r   any	ArrayAttr	Attributer   r   r_   r   Typer   	argumentsr/   r   r0   rf   r   r   element_typerJ   r>   rD   r   rt   )modulerh   r  _uses_barrier_semaphorespassespipelinemain_str_attrfr:   r  num_scalar_prefetchentry_blockr   sem_ty
dma_sem_typrogram_id_argsprefetch_args
other_argsr   r   rV   r   r   s                         @@r   export_promela_modelrB    s    ~ 4 4!$!>v?O!P!PAY__V-55T5BBCCFe$F !F388F3C3C!F!F!FGGHLL!"""M%%f--M[ 1 1	FD	!	!]	2	2 
3 /000a%%%%%&(Q\))-al;M.NOO..00g	4444#3444	4	4 CABBBL6K)LMM##$HIIh	8888$7888	8	8 K!"IJJJAL((N1<8I+JKKQVN[
#
$
$CW]]+,,F344J1;$4 5 57JK2 2.O]J /3?4HHH  
U	ggc5 6 6		!	!#(	+	+ 6]38$$?f$$:(E(E
''#s44RX>>
?
?
?
?
''#s**2844
5
5
5k"""==.BDT e4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4s   N N77N;>N;assume_for_verificationc                     | S rF   r   xys     r   <lambda>rH    s    q r   c                 ,    |                      |          S rF   )r3   rE  s     r   _assume_abstract_evalrJ    s    	
r   r   c                 "    | j         j        r|n|S rF   )lowering_contextfor_verification)r   rF  rG  s      r   _assume_loweringrN    s    "3	::r   c                8    t                               | |          S rF   )assume_pbindnormallywhen_verifyings     r   assumerU    s    	x	0	00r   pretend_for_verificationc                      ~dS )Nr   r   )r5  paramss     r   _pretend_abstract_evalrY  "  s
    	r   c                @   | j         j        rt          j        ||          \  }}t          j        || j                  \  }}t          j        || j                  \  }}d t          ||||d          D             }t          j        	                    d|           dS )Nc                 R    g | ]$\  }}}}t          j        ||||          d          %S )r   )r	   
_index_ref)r   refavalblock_shapeindexers        r   r   z%_pretend_lowering.<locals>.<listcomp>,  sF       +C{G 	C{G<<Q?  r   Tr   r   )r   r   )
rL  rM  r   tree_unflattenavals_inblock_shapesr   r   	Operationcreate)	r   tree	flat_argsbase_read_refsindexersread_ref_avalsr5  rc  r   s	            r   _pretend_loweringrk  '  s    * 
D!*!9$	!J!J^X!0s|DDNA.tS5EFFOL! /2NL(40
 0
 0
  I L.CCC	r   c                     t          d | D                       \  }}t          j        ||f          \  }}t          j        |d|iS )Nc              3   >   K   | ]}t          j        |          V  d S rF   )r
   _get_ref_and_indexers)r   r   s     r   r   zpretend.<locals>.<genexpr>8  s-      QQ!*:1==QQQQQQr   rf  )r   r   tree_flatten	pretend_prQ  )r   refsri  rg  rf  s        r   pretendrr  7  sP    QQyQQQQQ.$*D(+;<</)T		.	.	..r   c                       fd}|S )z-Skips the verification of the given function.c                  r     t          dd          } t          j        |           fd           d S )Nr   r   rR  c                        i S rF   r   argsr:  kwargss   r   rH  z'skip.<locals>.wrapper.<locals>.<lambda>A  s    qq$'9&'9'9 r   )rU  r   cond)rw  rx  is_not_verifyingr:  s   `` r   wrapperzskip.<locals>.wrapper?  sI    q;;;CH999999:::::r   r   )r:  r{  s   ` r   skipr|  =  s#    ; ; ; ; ; 
.r   c                       fd}|S )zBReplaces a function with its simplified model during verification.c                       fd}|S )Nc                  j     t          j        t          dd           fd fd           d S )Nr   r   rR  c                        i S rF   r   rv  s   r   rH  zBdefine_model.<locals>.decorator.<locals>.wrapper.<locals>.<lambda>K  s    !!T$V$$ r   c                        i S rF   r   )rw  rx  r   s   r   rH  zBdefine_model.<locals>.decorator.<locals>.wrapper.<locals>.<lambda>L  s    %%((( r   )r   ry  rU  )rw  rx  r:  r   s   ``r   r{  z0define_model.<locals>.decorator.<locals>.wrapperH  sW    	h
!A
.
.
.
$
$
$
$
$
$
(
(
(
(
(
(    r   r   )r:  r{  r   s   ` r   	decoratorzdefine_model.<locals>.decoratorG  s)          Nr   r   )r   r  s   ` r   define_modelr  E  s$         
r   )?rx   dataclassesrl   r)   rG   rN   typingr   r   jaxr   jax._srcr   jax_corer   jax._src.libr   jax._src.pallas.mosaicr	   r
   jax._src.utilr   r   jaxlib.mlirr   jaxlib.mlir.dialectsr   r   jaxlib.mlir.passmanagerr   objectrb   rT   Varro   rp   r   r   	dataclassrA   rI   r   r  r   r   r$   rB  	PrimitiverP  def_impldef_abstract_evalrJ  LoweringRuleContextrN  lowering_rulesrU  rp  multiple_resultsrY  rk  rr  r|  r  r   r   r   <module>r     s           				                             % % % % % %             + + + + + + - - - - - - , , , , , , , ,       & & & & & & % % % % % % / / / / / /vxx	6` ` ` ` ` ` ` `F	 	 	 d###Z Z Z Z Z Z Z $#Z& d###/ / / / / / / $#/ d###* * * * * * * $#*MD MD MD`2 2 2" " " 774777 7 7 7t 8788   .. ! ! !	  ;(6 ; ; ; ; %5  !1 1 1 H9::	!	 
  87     &7 	 "/ / /  
 
 
 
 
r   