
    rhJ6                         d dl mZmZmZ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mZ  ed      Zefd	ed
eg ef   fdZ G d dee         Zy)    )SequenceListTypeVarTupleCallable)TritonSemantic   )_core)SliceLayout)GluonOpBuilder)flatten_values_to_irunflatten_ir_valuesTensorTycondmsg_fnc                 "    | s | |             y N )r   r   categorys      /var/www/html/ai-insurance-compliance-backend/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/_semantic.py_checkr      s    vx       c                       e Zd ZU ej                  ZeZeed<   defdZd Z	de
e   de
e   fdZdeded	efd
Zdeded	ef fdZded	eeef   f fdZdedee   d	ef fdZdedee   d	efdZdeded	ef fdZ fdZdede
e   def fdZd Zd Zd Zd Zd Zd Zd  Zd! Zd" Z d# Z!d$ Z"d% Z#d& Z$d' Z%e&d(        Z'd)e(e   ded	eed*f   fd+Z)d,e(e   d-e(e   fd.Z* xZ+S )/GluonSemanticbuilderc                     || _         y r   )r   )selfr   s     r   __init__zGluonSemantic.__init__   s	    r   c                     t        j                  |j                  j                  |j                  | j
                  j                  |j                              }| j                  |j                  |      S r   )	ttgldistributed_typetypescalarshaper   get_gluon_layout_from_tensorhandletensor)r   r'   tys      r   _wrap_tensor_infer_layoutz'GluonSemantic._wrap_tensor_infer_layout   sP    ""6;;#5#5v||#'<<#L#LV]]#[]{{6=="--r   	lhs_shape	rhs_shapec                 X   t        |      t        |      k7  rt        d| d|       g }t        |      D ]q  \  }}||   }|dk(  r|j                  |       "|dk(  s||k(  r|j                  |       >t        dt	        |      z   dz   t	        |      z   dz   t	        |      z          |S )N!Cannot broadcast, rank mismatch: , r	   z?Cannot make_shape_compatible: incompatible dimensions at index : z and )len
ValueError	enumerateappendstr)r   r*   r+   	ret_shapeileftrights          r   _broadcast_shapeszGluonSemantic._broadcast_shapes   s    y>S^+@2i[YZZ	 + 	aGAtaLEqy  '1*%4-  &  "-/21v"68<"=?B4y"IKR"SUXY^U_"` a a	a r   inputaxisreturnc                    j                   D cg c]  }t        j                  |       }}|j                  d       dk  rt	        j                         z  t        t        j                  t        j                        fd       j                  j                  t        t        t              fd       t        j                  k(  fd       t        j                  j                  j                  |j                        }| j                  j                  j                   |j#                  | j                              }| j%                  ||      S c c}w )Nr	   r   c                  "    d j                   S Nz=expected expand_dims input to be a distributed_type but got: r"   r:   s   r   <lambda>z+GluonSemantic.expand_dims.<locals>.<lambda>6       VW\WaWaVde r   c                      d  S )Nz;expected expand_dims input to have a SliceLayout, but got: r   )layouts   r   rB   z+GluonSemantic.expand_dims.<locals>.<lambda>9   s    TU[T\] r   c                  (    d  dj                    S )Nz7expected expand_dims input layout to be sliced in axis z	 but got )dim)r;   rE   s   r   rB   z+GluonSemantic.expand_dims.<locals>.<lambda>;   s     PQUPVV_`f`j`j_kl r   )r$   r    _unwrap_if_constexprinsertr0   r   
isinstancer"   r!   rE   r   rG   r#   parentr   create_expand_dimsr&   to_irr'   )r   r:   r;   x	dst_shaperet_tyr&   rE   s    ``    @r   expand_dimszGluonSemantic.expand_dims.   s   ;@;;GaT..q1G	Gq!!8C$$Dz%**d&;&;<e	g""z&+.]	_vzzT!l	n &&uzz'8'8)V]]S00tV\\RVR^R^E_`{{66**! Hs   E'abc                     | j                  ||      \  }}t        |j                  g k7  d       t        |   ||      }| j                  |      S )NzCannot join scalars in gluon)broadcast_impl_valuer   r$   superjoinr)   )r   rR   rS   value	__class__s       r   rW   zGluonSemantic.joinA   sM    ((A.1qww"}<=Q"--e44r   c                 l    t         |   |      \  }}| j                  |      | j                  |      fS r   )rV   splitr)   )r   rR   lhsrhsrY   s       r   r[   zGluonSemantic.splitG   s7    7=#S--c2D4R4RSV4WWWr   dimsc                 F    t         |   ||      }| j                  |      S r   )rV   permuter)   )r   r:   r^   rX   rY   s       r   r`   zGluonSemantic.permuteK   s$    t,--e44r   r$   c                    t        t        j                  t        j                        fd       j                  j                         t        t              t              k(  fd       k(  rS t              D ]0  \  }}|   |k7  s|dk7  st        d|    d| d| d d 
       t        j                  j                  j                  j                  j                        }| j                  j                  j                  |j                  | j                              }| j                  ||      S )	Nc                  "    d j                   S r?   r@   rA   s   r   rB   z4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>Q   rC   r   c                      d d  S )Nr-   r.   r   )r$   	src_shapes   r   rB   z4GluonSemantic.broadcast_impl_shape.<locals>.<lambda>S   s    7XYbXccefkel5m r   r	   z3Cannot broadcast, the expanded size of the tensor (z ) must match the existing size (z) at non-singleton dimension r/   r.   )r   rJ   r"   r    r!   get_block_shapesr0   r2   r1   r#   rE   r   create_broadcastr&   rM   r'   )r   r:   r$   r6   itemrP   r&   rd   s    ``    @r   broadcast_impl_shapez"GluonSemantic.broadcast_impl_shapeO   s-   z%**d&;&;<e	gJJ//1	s9~U+-mnIL + 	@GAtQx4DAI #VW\]^W_V` aCCG& I%%&Cr)Bug"? @ @	@
 &&uzz'8'8%ARARS..u||V\\$,,=WX{{66**r   r\   r]   c                 d   |j                   |j                   j                         rj                         st        |   ||      S t	        t        t        j                        fd       t	        t        t        j                        fd       j                         }j                         }| j                  ||      }j                  j                  k7  r%t        dj                   dj                         | j                  ||      }| j                  ||      }||fS )Nc                      d S )Nz@expected broadcast left input to be a distributed_type but got: r   )lhs_tys   r   rB   z4GluonSemantic.broadcast_impl_value.<locals>.<lambda>g   s    YZ`Ycd r   c                      d S )NzAexpected broadcast right input to be a distributed_type but got: r   )rhs_tys   r   rB   z4GluonSemantic.broadcast_impl_value.<locals>.<lambda>i   s    Z[aZde r   zLayout mismatch in broadcast: z vs )r"   is_blockrV   rU   r   rJ   r    r!   re   r9   rE   r1   rh   )	r   r\   r]   r*   r+   r5   rk   rm   rY   s	         @@r   rU   z"GluonSemantic.broadcast_impl_value_   s    (97/S99z&$"7"78d	fz&$"7"78e	g ++-	++-	**9i@	==FMM)=fmm_DQWQ^Q^P_`aa''Y7''Y7Cxr   c                 ~    ||z
  g}t        j                  t         j                  ||      }t        |   |||      S )N)rP   )r    r!   int32rV   arange)r   startendrE   r$   rP   rY   s         r   rq   zGluonSemantic.arangeu   s;    u&&tzz5&Aw~eS~88r   rO   can_reorderc                 b    t        | d       t        | 	  |||      }| j                  |      S )Nz%can_reorder is not supported in gluon)r   rV   reshaper)   )r   r:   rO   rt   rX   rY   s        r   rv   zGluonSemantic.reshapez   s3    ; GHy+>--e44r   c                     t        j                  |j                  ||      }| j                  j	                  |j                  | j                        |j                        }t        j                  ||      S r   )r    r!   dtyper   create_splatrM   r&   r'   )r   rX   r$   rE   rP   r&   s         r   splatzGluonSemantic.splat   sR    &&u{{E6B**6<<+Eu||T{{66**r   c                 L    | j                  ||      }| j                  |||      S r   )make_scalarrz   )r   r$   rX   rx   rE   r#   s         r   fullzGluonSemantic.full   s'    !!%/zz&%00r   c                 j   |j                   t        t        t        j                        fd       t        j                  j
                  j                  |      }| j                  j                  |j                  | j                        |j                        }t        j                  ||      S )Nc                      d S )Nz@expected convert_layout input to be a distributed_type but got: r   )r(   s   r   rB   z.GluonSemantic.convert_layout.<locals>.<lambda>   s    YZ\Y_` r   )r"   r   rJ   r    r!   
element_tyr$   r   create_convert_layoutrM   r&   r'   )r   rX   rE   rP   r&   r(   s        @r   convert_layoutzGluonSemantic.convert_layout   s~    ZZz"d334`	b&&r}}bhhG33FLL4NPUP\P\]{{66**r   c                 P   t        j                  ||||      }|@| j                  j                  |j	                  | j                        |j
                        }n4| j                  j                  |j	                  | j                              }t        j                  |||||      S r   )r    shared_memory_descriptor_typer   create_local_allocrM   r&   shared_memory_descriptor)r   r   r$   rE   rX   r(   r&   s          r   allocate_sharedzGluonSemantic.allocate_shared   s    //
E65Q\\44RXXdll5KU\\ZF\\44RXXdll5KLF,,VZPUVVr   c                    t        j                  |j                  |j                  |      }| j                  j                  |j                  | j                        |j                        }t        j                  ||      S r   )	r    r!   rx   r$   r   create_local_loadrM   r&   r'   )r   mem_descrE   rP   r&   s        r   shared_loadzGluonSemantic.shared_load   sV    &&x~~x~~vN//T\\0JHOO\{{66**r   c                 d    | j                   j                  |j                  |j                         y r   )r   create_local_storer&   )r   r   rX   s      r   shared_storezGluonSemantic.shared_store   s    ''Fr   c                 N    | j                   j                  |j                         y r   )r   create_local_deallocr&   )r   r   s     r   shared_dealloczGluonSemantic.shared_dealloc   s    ))(//:r   c                 6   |j                   }t        j                  |j                  |||j                  j
                        }| j                  }|j                  |j                  |      |j                  |      }t        j                  |fi |j                  S r   )rE   r    r   rx   r"   alloc_shaper   create_memdesc_subviewrM   r&   r   __dict__)r   r   offsetsr$   rE   r(   r   r&   s           r   _memdesc_subviewzGluonSemantic._memdesc_subview   st    //vx}}OhOhi,,//0A8??T[\,,VCr{{CCr   c                     | j                   j                  d      g|j                  z  }| j                  |      j                  ||<   t        |j                        }|||<   | j                  |||      S )Nr   )r   	get_int32rank	to_tensorr&   listr$   r   )r   r   rr   lengthrG   r   r$   s          r   memdesc_slicezGluonSemantic.memdesc_slice   sf    <<))!,-=~~e,33X^^$c
$$Xw>>r   c                     |j                   dd  }| j                  j                  d      g|j                  z  }| j	                  |      j
                  |d<   | j                  |||      S )Nr	   r   )r$   r   r   r   r   r&   r   )r   r   indexr$   r   s        r   memdesc_indexzGluonSemantic.memdesc_index   s_    qr"<<))!,-=^^E*11
$$Xw>>r   c                 \   t        |      t        |j                        k(  s!J d|j                   dt        |       d       |D cg c]  }|j                  |    }}|j                  j                  }|d t        |      |j                  z
   }||D cg c]   }|t        |      |j                  z
  d  |   " c}z  }| j
                  j                  |j                  |      }| j
                  j                  |      }t        j                  ||j                  |||      S c c}w c c}w )Nzsource rank (z) and order length (z) must match)r   r$   r   rE   )r0   r$   r   r"   r   r   create_memdesc_transr&   get_gluon_layout_from_memdescr    r   rx   )	r   r   orderr6   r$   r   new_alloc_shaper&   rE   s	            r   memdesc_transzGluonSemantic.memdesc_trans   s%   5zSNN  	i,X]]O;OPSTYPZ|[gh	i  -22q"22mm//%&Gs;'7(--'GHW\]RSKK(88==(H(IJ1M]]228??EJ;;FC,,VV[9HQWY 	Y 3 ^s   D$%D)c                 ,   t        j                  |j                  |||j                  j                        }| j
                  j                  |j                  | j
                        |j                        }t        j                  |fi |j                  S r   )r    r   rx   r"   r   r   create_memdesc_reshaperM   r&   r   r   )r   r   r$   rE   r(   r&   s         r   memdesc_reshapezGluonSemantic.memdesc_reshape   sg    //vx}}OhOhi44RXXdll5KX__],,VCr{{CCr   c                     t        j                  ||||      }| j                  j                  |j	                  | j                        |j
                        }t        j                  |fi |j                  S r   )r    r   r   create_memdesc_reinterpretrM   r&   r   r   )r   r   rx   r$   rE   r(   r&   s          r   memdesc_reinterpretz!GluonSemantic.memdesc_reinterpret   s\    //ufeL88$,,9OQYQ`Q`a,,VCr{{CCr   c                 ^    |rt        j                  |||      }n|}| j                  ||      S r   )r    r!   r'   )r   rN   	scalar_tyr5   rE   res_tys         r   wrap_tensorzGluonSemantic.wrap_tensor   s/    **9iHFF{{1f%%r   c                    | D ]3  t        t        j                  t        j                        fd       5 | D cg c]  }|j                  j
                   c}d   t        t        fddd  D              fd       y c c}w )Nc                  "    d j                   S )Nz#expected distributed_type but got: r@   )rN   s   r   rB   z2GluonSemantic._check_same_layout.<locals>.<lambda>   s    HklmlrlrkuFv r   r   c              3   (   K   | ]	  }|k(    y wr   r   ).0ll0s     r   	<genexpr>z3GluonSemantic._check_same_layout.<locals>.<genexpr>   s     0q170s   r	   c                      d  S )Nz3Expected inputs to have matching layouts, but got: r   )layoutss   r   rB   z2GluonSemantic._check_same_layout.<locals>.<lambda>   s    LWIV r   )r   rJ   r"   r    r!   rE   all)xsrN   r   r   s    `@@r   _check_same_layoutz GluonSemantic._check_same_layout   su     	xA:affd&;&;<>vw	x*,-Q166==-QZs0GABK00V	X .s    B
inputs.c                    	
 t        d ud        d   j                  j                  t              t        dcxk  xr k  nc fd        j	                         t              D cg c]  \  }}|k7  s| c}}
t        d   j                  j                        	t        fdD              sJ d        j                  j                  D cg c]  }|j                   c}       |       j                         sJ t        	
 fdt        t                    D              S c c}}w c c}w )Nc                       y)Nz*All-reduce is not yet implemented in gluonr   r   r   r   rB   z)GluonSemantic.reduction.<locals>.<lambda>   s    r   r   c                      d d  S )Nz/expected reduction axis to be in the range [0, z
) but got r   )r;   r   s   r   rB   z)GluonSemantic.reduction.<locals>.<lambda>   s    +Z[_Z``jkojp)q r   c              3   P   K   | ]  }|j                   j                  k(    y wr   )r"   r$   )r   tr$   s     r   r   z*GluonSemantic.reduction.<locals>.<genexpr>   s     9Q166<<5(9s   #&z-all reduction inputs must have the same shapec              3      K   | ]>  }j                  j                  |      |   j                  j                         @ y wr   )r   
get_resultr"   r#   )r   r6   r   	reduce_op
ret_layoutr5   r   s     r   r   z*GluonSemantic.reduction.<locals>.<genexpr>   sD      ) Y11!4fQinn6K6KYXbc)s   AA)r   r"   r$   r0   r   r2   r   rE   r   r   create_reducer&   verifytuplerange)r   r   r;   region_builder_fnr6   sr   r   r   r   r5   r$   s   ```    @@@@@r   	reductionzGluonSemantic.reduction   s   t4!UVq	$$5zqD4!qr'#,U#3A41aqDyQA	 vay~~'<'<=
9&99j;jj9LL..&/IQ/I4P	)$!!! )3v;') ) 	) B 0Js   9EE&Eworker_num_warpsworker_num_regsc           	         t        |      }|t        |      k(  sJ d| dt        |       d       |t        |      k(  sJ d| dt        |       d       | j                  }|j                         }	|j                         }
|j	                  |
       |j                  ||i       }g }|t        |      }|j                  |       |D cg c]  }|j                          }}|j                  |	       t        |      }|j                  |||      }|j                         j                  |
       |j                  |       |j                  |j                         g        |j!                  |      }|D cg c]  }|j                          }}t#        |      D ]  }|j                  |j%                  |      |      }t#        t        |            D cg c]  }|j'                  |       }}t)        ||D cg c]  }|j*                   c}      }|j                  ||   |i        |j-                           |j/                  |j1                                t#        t        |            D cg c]  }|j3                  |       }}|y t5        t)        ||D cg c]  }|j*                   c}            S c c}w c c}w c c}w c c}w c c}w c c}w )Nzwarp specialize got z partitions but z warp countsz register counts)kwargs)r0   r   get_insertion_point	new_blockset_insertion_point_to_startcall_JitFunctionr   create_warp_yieldget_typerestore_insertion_pointcreate_warp_specializeget_default_region	push_backset_requested_registerscreate_block_with_parentget_partition_op_holder!create_warp_specialize_partitionsr   
get_regionget_argumentr   r"   create_warp_returnset_insertion_point_afterget_operationr   r   )r   argsdefault_partitionworker_partitionsr   r   	generatornum_partitionsr   	insert_ptdefault_blockdefault_resultsmlir_resultsrresult_types	mlir_argsws_oppartitions_oparg	arg_typesr6   blockj
block_argss                           r   warp_specializezGluonSemantic.warp_specialize   s   ./"
 
 	f!.!11A#FVBWAXXde	f 
 "
 
 	i!.!11A#oBVAWWgh	i 
 ,,//1	  ))+,,];#445FUW4X&/@L!!,/.:;

;; 	''	2(.	..|YHXY  ",,];%%o6 	(()F)F)H"MAA.Q/89S\\^9	9~& 	)A44]5M5Ma5PR[\E9>s9~9NOA%,,Q/OJO,Zd9Ss#((9STJ&&'8';ZPR&S&&(	) 	))%*=*=*?@5:3|;L5MN((+NN"(7X17XYZZ1 < : P9S
 O 8Ys$   J<8KK:K2KK),__name__
__module____qualname__r    r'   langr   __annotations__r   r)   r   intr9   r   rQ   rW   r   r[   r`   rh   rU   rq   boolrv   rz   r}   r   r   r   r   r   r   r   r   r   r   r   r   staticmethodr   r   r   r   __classcell__)rY   s   @r   r   r      s   [[FD .
49 c  + + + +&5h 58 5 5Xx XE(H*<$= X5X 5U3Z 5H 5+( +5: +( +  x H ,9
5X 5$s) 5$ 5
+
1+W+
G;D??YD
D
& X X) 2 )# )UZ[ceh[hUi )&-[\deh\i -[)1#-[r   r   N)typingr   r   r   r   r   triton.language.semanticr    r
   r    _layoutsr   triton._C.libtriton.gluon_irr   triton.compiler.code_generatorr   r   r   r1   r   r4   r   r   r   r   r   <module>r	     s[    ; ; 3  ! 7 T: <F ! !xC0 !
O[N8, O[r   