From 354fb07525eacfc98d508ca945bd09289aa0954b Mon Sep 17 00:00:00 2001 From: Zeeshan Siddiqui Date: Sun, 17 May 2020 23:32:15 -0700 Subject: [PATCH] Add comments for zero valued normalization factor in SoftmaxCrossEntropyLossGrad CUDA kernel. --- .../cuda/loss/softmax_cross_entropy_loss_impl.cu | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/orttraining/orttraining/training_ops/cuda/loss/softmax_cross_entropy_loss_impl.cu b/orttraining/orttraining/training_ops/cuda/loss/softmax_cross_entropy_loss_impl.cu index 5f93231c9baa9..aeda6fee4cd5d 100644 --- a/orttraining/orttraining/training_ops/cuda/loss/softmax_cross_entropy_loss_impl.cu +++ b/orttraining/orttraining/training_ops/cuda/loss/softmax_cross_entropy_loss_impl.cu @@ -118,6 +118,9 @@ __global__ void _WeightedSoftmaxCrossEntropyLossGrad( int d = i % C; CUDA_KERNEL_ASSERT(weight[row] == 0 || (label[row] >= 0 && label[row] < C)); if(0 == *normalize_factor){ + // normalize_factor is sum of labels' weights. Because zero + // sum implies all weights are 0, the loss function should + // be constant 0 and its corresponding gradient should be 0 as well. output_data[i] = 0; } else { output_data[i] = (*dY) * weight[row] * (_Exp(log_prob[i]) - 1.0 * (d == label[row])) / (*normalize_factor); @@ -140,6 +143,9 @@ __global__ void _WeightedReductionNoneSoftmaxCrossEntropyLossGrad( int d = i % C; CUDA_KERNEL_ASSERT(weight[row] == 0 || (label[row] >= 0 && label[row] < C)); if(0 == *normalize_factor){ + // normalize_factor is sum of labels' weights. Because zero + // sum implies all weights are 0, the loss function should + // be constant 0 and its corresponding gradient should be 0 as well. output_data[i] = 0; } else { output_data[i] = dY[row] * weight[row] * (_Exp(log_prob[i]) - 1.0 * (d == label[row])) / (*normalize_factor);