[TUTORIALS] resolve uninitialized memory access when M< GROUP_SIZE_M in LN tutorial (triton-lang#2837)

👋 There's currently a (minor) issue in the LN tutorial where if M is
smaller than GROUP_SIZE_M parts of _dw/_db will contain uninitialised
memory which is later summed into the final dw/db.

Author: DanielHesslow

python/tutorials/05-layer-norm.py

@@ -283,7 +283,7 @@ def backward(ctx, dy):
         grid = lambda meta: [triton.cdiv(N, meta['BLOCK_SIZE_N'])]
         # accumulate partial sums in separate kernel
         _layer_norm_bwd_dwdb[grid](
-            _dw, _db, dw, db, GROUP_SIZE_M, N,  #
+            _dw, _db, dw, db, min(GROUP_SIZE_M, M), N,  #
             BLOCK_SIZE_M=32,  #
             BLOCK_SIZE_N=128, num_ctas=1)
         return dx, None, dw, db, None