From 27cf5223e03e4bf8d64a9905e23ebdb3ab259198 Mon Sep 17 00:00:00 2001 From: QIAOQIAN <73099895+huoyushequ@users.noreply.github.com> Date: Wed, 29 May 2024 09:23:46 +0800 Subject: [PATCH] `softmax_autoregressive_backward_kernel` does not use share memory in the kernel we do not need to launch the kernel with 256 bytes share memory, so remove it --- train_gpt2.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index c511b49..1bdf11e 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -1875,7 +1875,7 @@ void attention_backward(floatX* dinp, floatX* dqkvr, floatX* dpreatt, floatX* da // backward into preatt int hs = C / NH; // head size float scale = 1.0f / sqrtf(hs); - softmax_autoregressive_backward_kernel<<>>(dpreatt, datt, att, B, T, C, scale); + softmax_autoregressive_backward_kernel<<>>(dpreatt, datt, att, B, T, C, scale); // backward into q cublasCheck(cublasGemmStridedBatchedEx(cublas_handle, CUBLAS_OP_N, CUBLAS_OP_N, HS, T, T, &alpha, k, CUBLAS_LOWP, HS, T * HS, dpreatt, CUBLAS_LOWP, T, T * T, &beta,