jit def _layer_norm_bwd_dx_fused ( DX, # pointer to the input gradient DY, # pointer to the output gradient DW, # pointer to the partial sum of weights gradient DB, # pointer to the partial sum of biases gradient X, # pointer to the input W, # pointer to the weights B, # pointer to the biases Mean, # pointer to the mean Rstd, # pointer to the 1/std Lock, # pointer to the lock stride, # how much to increase the pointer when moving by 1 row N, # number of columns in X eps, # epsilon to avoid division by zero GROUP_SIZE_M : tl. The backward pass for the layer normalization operator is a bit more involved than the forward pass. store ( Y + cols, y, mask = mask ) Backward pass ¶ float32 ) x_hat = ( x - mean ) * rstd y = x_hat * w + b # Write output tl. load ( X + cols, mask = mask, other = 0. arange ( 0, BLOCK_SIZE ) mask = cols < N w = tl. store ( Rstd + row, rstd ) # Normalize and apply linear transformation for off in range ( 0, N, BLOCK_SIZE ): cols = off + tl.

sqrt ( var + eps ) # Write mean / rstd tl. sum ( _var, axis = 0 ) / N rstd = 1 / tl. load ( X + cols, mask = cols < N, other = 0. float32 ) for off in range ( 0, N, BLOCK_SIZE ): cols = off + tl. sum ( _mean, axis = 0 ) / N # Compute variance _var = tl. program_id ( 0 ) Y += row * stride X += row * stride # Compute mean mean = 0 _mean = tl. constexpr, ): # Map the program id to the row of X and Y it should compute. jit def _layer_norm_fwd_fused ( X, # pointer to the input Y, # pointer to the output W, # pointer to the weights B, # pointer to the biases Mean, # pointer to the mean Rstd, # pointer to the 1/std stride, # how much to increase the pointer when moving by 1 row N, # number of columns in X eps, # epsilon to avoid division by zero BLOCK_SIZE : tl. import apex HAS_APEX = True except ModuleNotFoundError : HAS_APEX = False. Import torch import triton import triton.language as tl try : # This is, NOT the apex on PyPi, so it # should not be added to extras_require in setup.py.

Torch nn sequential pointer