From 8a2cab343153bf964fbc8fde8c813dd222e6e47c Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Fri, 19 Apr 2024 00:27:28 +0000 Subject: [PATCH] make default dloss be 1/BT --- dev/cuda/classifier_fused.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/dev/cuda/classifier_fused.cu b/dev/cuda/classifier_fused.cu index b9d84a1..479662e 100644 --- a/dev/cuda/classifier_fused.cu +++ b/dev/cuda/classifier_fused.cu @@ -244,9 +244,10 @@ __global__ void fused_classifier_kernel2(float* dlogits, float* losses, float* p losses[idx] = -logf(prob); } + // very sensible default for dlosses is 1/(B*T), which is the uniform loss + float dloss = dlosses != NULL ? dlosses[idx] : 1.0f / (B*T); // calculate the gradients directly, saves bandwidth from probs during training // but also supports writing probs for inference-only and debugging - float dloss = dlosses ? dlosses[idx] : 0.f; const float4* logits_vec4 = reinterpret_cast(logits + idx * P); for (int i = threadIdx.x; i < (V+3)/4; i += blockDim.x) { // this is the 2nd read of logits after the one in prepare_softmax2