[Common] Remove redundant grad_logits zero-initialization in fused router backward kernels (NVIDIA#2745)

Remove redundant grad_logits zero-initialization in fused router backward kernels

Signed-off-by: Sung Hyun Cho <hope5487@gmail.com>
Co-authored-by: Xin Yao <xiny@nvidia.com>

Author: roycho96

transformer_engine/common/fused_router/fused_score_for_moe_aux_loss.cu

@@ -211,10 +211,6 @@ __global__ void fused_score_for_moe_aux_loss_backward_kernel(const CompType *int
          * - Load the dgrad/output_from_fwd to shmem
          */
     int pos_offset = token_offset_cur_warp * num_experts;
-    // Clear the logits_grad in global mem
-    for (int i = lane_id; i < num_experts; i += kThreadsPerWarp) {
-      grad_logits[pos_offset + i] = 0.0f;
-    }
     // Load the dgrad/output_from_fwd to shmem
     for (int i = lane_id; i < num_experts; i += kThreadsPerWarp) {
       local_grad[i] = grad_scores[pos_offset + i];

transformer_engine/common/fused_router/fused_topk_with_score_function.cu

@@ -328,10 +328,6 @@ __global__ void fused_topk_with_score_function_backward_kernel(
          * - Load the dgrad/output_from_fwd to shmem
          */
     int pos_offset = token_offset_cur_warp * num_experts;
-    // Clear the logits_grad in global mem
-    for (int i = lane_id; i < num_experts; i += kThreadsPerWarp) {
-      grad_logits[pos_offset + i] = 0.0f;
-    }
     // Load the dgrad/output_from_fwd to shmem
     for (int i = lane_id; i < num_experts; i += kThreadsPerWarp) {
       local_grad[i] = grad_probs[pos_offset + i];