Vpf dZddlZddlZddlmZddlmZde fdZ ej ej gddd d d dd ej d e d e dededej f dZdS)zPallas softmax kernel.N)pallas block_rowc|jd}tj||k}tj|tjd|f|t d }tj|d}tj||z tj }tj |d}tj |tjd|f||z |j |dS)Nrinf)maskother)axis)r)shapejnparangeplloaddslicefloatmaxexpastypefloat32sumstoredtype) input_ref probs_refrrow_lenrrowrow_max numerator denominators g/var/www/html/nettyfy-visnx/env/lib/python3.11/site-packages/jax/experimental/pallas/ops/gpu/softmax.py_vmappable_softmax_kernelr!s OB ' I   ($ ")Ay))+$uU||m   # GCa 'gsW},,S[99::) ***+(")Ay))+;&&y77 )r num_warps interpretdebug)static_argnamesrFxr r#r$r%returnc >|dkr|nt|j|z}|t|jdz krtd|jd}tj|}t j|f|j}tj t|}tj |tt|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)triton)compiler_paramsgrid out_shaper%r$)lenr NotImplementedErrorrnext_power_of_2jaxShapeDtypeStructr functoolspartialr! pallas_calldictrangevmap) r(r r#r$r%rrr1kernelf_s r softmaxr@3s$AG t 3$ S\\A  = ? ?? GBK' )))"'17CCC)  6) L L L&n $q"I"I"IJJJ   ! QW! " "a  AA 1+r")__doc__r7r5 jax.numpynumpyr jax.experimentalrrintr!r8jitArrayboolr@r.r"r rIs ))))))837-6-6-6777"$a5''' y''03''$(' Y'''77'''r"