ukitUdZddlmZddlmZddlmZddlZddlm Z ddlm Z ddl m Zdd lmZdd lmZdd lmZdd lmZddlmZe j8e j:zZd ed<ddZ ddZ!e jDdZ#de#_$e#jJ d dZ&ejNe# d!dZ(d"dZ)e jDdZ*de*_$e*jJd#dZ+ejNe*d!dZ,dddddd d$dZ-ddd d%dZ.y)&z'Module for GPU-specific JAX primitives.) annotations)Sequence) TypeAliasN)core)state)gpu)dialect) primitives)lowering)mlirrRefc |jtjk(rd}d}n]|jtjk(rd}d}n;|jtjk(rd}d}nt d|jdt ||gd|d |d tj|j|jg \}|S) zElementwise approximate hyperbolic tangent: :math:`\mathrm{tanh}(x)`. See https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#floating-point-instructions-tanh. ztanh.approx.f16 $0, $1;hztanh.approx.bf16 $0, $1;ztanh.approx.f32 $0, $1;fzapprox_tanh does not accept z arrays=,)args constraintspackresult_shape_dtypes) dtypejnpfloat16bfloat16float32 TypeErrorelementwise_inline_asmjaxShapeDtypeStructshape)xasm constraintresults \/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/_src/pallas/triton/primitives.py approx_tanhr'$s WW  #CJww#,, $CJww#++ #CJ 2177)7C DD # 3j\:,/ //AB (6 -c Dtj||||t|dS)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 `_. 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#rrr)elementwise_inline_asm_pbindtuple)r#rrrrs r&rr@s,. " & &   34  r(r*Tc ~tdt||ddDs td|Dcgc],}tj|j |j .c}Scc}w)Nc3TK|] \}}|j|jk("ywN)r!).0r"ys r& z8_elementwise_inline_asm_abstract_eval..is! BDAqQWW  Bs&(rz@All arguments of elementwise_inline_asm must have the same shape)allzip ValueErrorjax_core ShapedArrayr!r)ravalskwargsss r&%_elementwise_inline_asm_abstract_evalr;ds^ BCuQRy,A B B J ;N NQ(  qww 0 NN Ns1A$c~tjgttj|j ||d||j S)NT)rpurepacked_elementr) tt_dialectElementwiseInlineAsmOpmapr aval_to_ir_type avals_outr%)ctxr#rrrrs r& _elementwise_inline_asm_loweringrEpsI  * *1D #--01      F r(c*tjS)z/Synchronizes all kernel executions in the grid.)debug_barrier_pr+r(r& debug_barrierrIs    r(rGcy)NrHrHrHr(r&_debug_barrier_abstract_evalrKs r(c0~tjgSr/) gpu_dialectbarrier)rDs r&_debug_barrier_loweringrOs   )r(Fmaskothercache_modifiereviction_policyvolatilec :tj|d|||||S)a:Loads an array from the given ref. If neither ``mask`` nor ``other`` is specified, this function has the same semantics as ``ref[idx]`` in JAX. Args: ref: The ref to load from. mask: An optional boolean mask specifying which indices to load. If mask is ``False`` and ``other`` is not given, no assumptions can be made about the value in the resulting array. other: An optional value to use for indices where mask is ``False``. cache_modifier: TO BE DOCUMENTED. eviction_policy: TO BE DOCUMENTED. volatile: TO BE DOCUMENTED. NrP)pallas_primitivesload)refrQrRrSrTrUs r&rXrXs,0      #% r(rQrTc6tj|d|||S)zrStores a value to the given ref. See :func:`~jax.experimental.pallas.load` for the meaning of the arguments. NrZ)rWstore)rYvalrQrTs r&r\r\s&     %  r()r" jax.Arrayreturnr^) r#strrSequence[jax.Array]rr`rintrzSequence[jax.ShapeDtypeStruct]r_ra)r8zjax_core.ShapedArrayr_Sequence[jax_core.ShapedArray])rDzlowering.LoweringRuleContext)r_None)r_rc)rYr rQjax.Array | NonerRzjax.typing.ArrayLike | NonerS str | NonerTrfrUboolr_r^) rYr r]r^rQrerTrfr_rd)/__doc__ __future__rcollections.abcrtypingrrjax._srcrr6rjax._src.lib.mlir.dialectsrrMjax._src.lib.tritonr r?jax._src.pallasr rWjax._src.pallas.tritonr jax.interpretersr jax.numpynumpyr AbstractRefTransformedRefr __annotations__r'r Primitiver*multiple_resultsdef_abstract_evalr;register_loweringrErIrGrKrOrXr\rHr(r&r{s."$ %95;+!""U%9%99Y98      8 @.8--.HI,0)++O O#O,O45 % 6 & %($$%67#' "" # O, - ")-!%"&      '       N""&         r(