a zd@sddlZddlZddlmZddZddZedddiedd diejej d d d Z edd diedddiejej d ddZ Gdddej j ZejZdS)NcCsP|d8}||d?O}||d?O}||d?O}||d?O}||d?O}|d7}|S)N)nrra/var/www/html/stable-diffusion-webui/venv/lib/python3.9/site-packages/triton/ops/cross_entropy.pynext_power_of_2s     r cCs|dkr dS|dkrdSdS)Niri rrr)Nrrr num_warpss r cCs t|dSNr r nargsrrr rBLOCKcCs t|dSr r rrrr rr)rc Cstd}td|}t||}||||}||||} ||||} tj|||ktd d} | tj} | t| d} tt t | d| } tj | | ||kdt t| } t ||| dSNrinf)maskother)r) tl program_idarangeloadfloattofloat32maxlogsumexpstoreZ debug_barrier) ZLOGITSPROBSIDXZLOSSr rrowcolsidxZ WRIT_PROBSZ READ_PROBSlogitsprobsrrr _forwards    r,cCs t|dSr rrrrr r3rcCs t|dSr rrrrr r4rc Cstd}td|}t||}||||}tj|||ktdd }t|tj}||k} t||} || | } tj|| |j j ||kddSr) rrrrrr#rrr$dtypeZ element_ty) r%r&ZDPROBSr rr'r(r)r+deltaZdoutZdinrrr _backward3s   r/c@s$eZdZeddZeddZdS)_cross_entropyc s~|jtjksJdjj}}jdtj|||d}tj||d}fdd}t||||||||S)Nz(Indices are expected to be of type long.)r-devicecsfSNZnumeloptr*n_colsrr rRrz(_cross_entropy.forward..)r-torchint64r2shapeZ empty_liker,Zsave_for_backward) clsctxr*indicesr2r-result neg_logprobsgridrr7r forwardHs  z_cross_entropy.forwardcs<|j\}jdfdd}t|||dfS)aWe know d(-log(p[i])/dlogit[k] = -id_mat[i,k] + p[k] so we initialize the gradient as neg_logprobs, so we can just exponentiate to get p[k], which is most of what we need... neg_logprobs will be modified in place to become the gradient we want r1csfSr3r4r5r8r@rr rdrz)_cross_entropy.backward..N)Z saved_tensorsr;r/)r<r=Z dneg_logprobsr>rArrCr backwardXs   z_cross_entropy.backwardN)__name__ __module__ __qualname__ classmethodrBrDrrrr r0Gs r0)r9ZtritonZtriton.languagelanguagerr r Z heuristicsZjitZ constexprr,r/ZautogradZFunctionr0applyZ cross_entropyrrrr s  "