
    VpfO                        d Z ddlmZ ddlm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Zd!dZ ej        d          Zde_        ej        d"d            Z ej        e          d#d            ZdS )$z'Module for GPU-specific JAX primitives.    )annotations)SequenceN)core)dialect)lowering)mlirx	jax.Arrayreturnc           
     H   | j         t          j        k    rd}d}nL| j         t          j        k    rd}d}n2| j         t          j        k    rd}d}nt          d| j          d          t          || gd| d	| d
t          j        | j	        | j                   g          \  }|S )zElementwise approximate hyperbolic tangent: :math:`\mathrm{tanh}(x)`.

  See
  https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-tanh.
  ztanh.approx.f16 $0, $1;hztanh.approx.bf16 $0, $1;ztanh.approx.f32 $0, $1;fzapprox_tanh does not accept z arrays=,   )argsconstraintspackresult_shape_dtypes)
dtypejnpfloat16bfloat16float32	TypeErrorelementwise_inline_asmjaxShapeDtypeStructshape)r	   asm
constraintresults       a/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/_src/pallas/triton/primitives.pyapprox_tanhr$      s     W
#CJJw#,
$CJJw#+
#CJJ
C17CCC
D
DD#	3/j//:///AAB  (6 
-    r    strr   Sequence[jax.Array]r   r   intr   Sequence[jax.ShapeDtypeStruct]c               *    t          j        || |||dS )a!  Inline assembly applying an elementwise operation.

  Args:
    asm: The assembly code to run.
    args: The arguments to pass to the assembly code.
    constraints: LLVM inline assembly `constraints
      <https://llvm.org/docs/LangRef.html#inline-asm-constraint-string>`_.
    pack: The number of elements from each argument expected by a single
      instance of the assembly code.
    result_shape_dtypes: The shapes and dtypes of the results produced by the
      assembly code.

  Returns:
    The results produced by the assembly code.
  )r    r   r   r   )elementwise_inline_asm_pbind)r    r   r   r   r   s        r#   r   r   9   s-    . 
"	&
-
 
 
 r%   r+   Tavalsjax_core.ShapedArraySequence[jax_core.ShapedArray]c           	         ~t          d t          ||dd                    D                       st          d          d | D             S )Nc              3  <   K   | ]\  }}|j         |j         k    V  d S )N)r   ).0r	   ys      r#   	<genexpr>z8_elementwise_inline_asm_abstract_eval.<locals>.<genexpr>b   s/      BBDAqQWBBBBBBr%   r   z@All arguments of elementwise_inline_asm must have the same shapec                L    g | ]!}t          j        |j        |j                  "S  )jax_coreShapedArrayr   r   )r2   ss     r#   
<listcomp>z9_elementwise_inline_asm_abstract_eval.<locals>.<listcomp>f   s)    	N	N	NQ(
qw
0
0	N	N	Nr%   )allzip
ValueError)r   r-   kwargss      r#   %_elementwise_inline_asm_abstract_evalr?   ]   se     	BBCuQRRy,A,ABBB	B	B 
J   
O	N:M	N	N	NNr%   ctxlowering.LoweringRuleContextc                   ~t          j        g t          t          j        | j                  ||d||          j        S )NT)r   purepacked_elementr   )
tt_dialectElementwiseInlineAsmOpmapr   aval_to_ir_type	avals_outr"   )r@   r    r   r   r   r   s         r#    _elementwise_inline_asm_loweringrJ   i   sL     		*1D #-001	
 
 
 r%   )r	   r
   r   r
   )r    r&   r   r'   r   r&   r   r(   r   r)   r   r'   )r-   r.   r   r/   )r@   rA   )__doc__
__future__r   collections.abcr   r   r   r7   jax._src.lib.tritonr   rE   jax._src.pallas.tritonr   jax.interpretersr   	jax.numpynumpyr   r$   r   	Primitiver+   multiple_resultsdef_abstract_evalr?   register_loweringrJ   r6   r%   r#   <module>rW      s=   . - " " " " " " $ $ $ $ $ $ 



             5 5 5 5 5 5 + + + + + + ! ! ! ! ! !         8   @ .8-.HII ,0  ) +O O O ,+O 455   65  r%   