uki dZddlZddlZddlmZddlmZddl m Z de fdZ ejejgdd d d d dd ej d e de dededej f dZy)zPallas softmax kernel.N)pallas)triton block_rowcH|jd}tj||k}tj|j t jd||td }tj|d}tj||z jtj}tj|d}tj|j t jd|||z j|j|y)Nrinf)maskother)axis)r )shapejnparangeplgpuloadatpldsfloatmaxexpastypefloat32sumstoredtype) input_ref probs_refrrow_lenr rowrow_max numerator denominators b/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/experimental/pallas/ops/gpu/softmax.py_vmappable_softmax_kernelr$s OOB ' I  ($ ll255I&'d5<- # GGCa 'ggsW},,S[[9:) *+++ll255I&';&&y7 )r num_warps interpretdebug)static_argnamesrFxr r&r'r(returnc4|dk\r|nt|j|z}|t|jdz k7r td|jd}tj|}t j |f|j}tjt|}tj|tj|dd||| } tt|jdz D]} t j| } | |S) aComputes the softmax of the input array along the specified axis. Args: x: input array axis: the axis along which to perform the computation num_warps: the number of warps to use for executing the Triton kernel interpret: whether to interpret the kernel using pallas debug: whether to use pallas in debug mode Returns: The result of the softmax operation over the specified axis of x. rz3reductions along non-trailing dimension unsupportedr)r r)r)r& num_stages)compiler_paramsgrid out_shaper(r')lenr NotImplementedErrorrnext_power_of_2jaxShapeDtypeStructr functoolspartialr$ pallas_callrCompilerParamsrangevmap) r+r r&r'r(rrr3kernelf_s r#softmaxrB4s$AGG t 3$ S\A  = ?? GGBK'  ))""'177C)   6) L&nn **!-  ! QWW! "a  A 1+r%)__doc__r9r7 jax.numpynumpyr jax.experimentalrrjax.experimental.pallasrrintr$r:jitArrayboolrBr0r%r#rLs )38377-67"$a5( yy((03(($(( YY(7(r%