aparler/c-llm
1
0
Fork 0
mirror of https://github.com/karpathy/llm.c.git synced 2026-07-28 20:35:09 -04:00
Code Issues Projects Releases Packages Wiki Activity Actions

set stream for attention softmax

Browse source
This commit is contained in:
Erik Schultheis 2025-05-04 22:10:38 +02:00
parent 082d9fa78e
commit a860922827
1 changed files with 1 additions and 1 deletions
Download patch file Download diff file Expand all files Collapse all files

2
llmc/attention.cuh
View file

@ -263,7 +263,7 @@ void attention_backward(floatX* dinp, floatX* dqkvr, floatX* datt, floatX* scrat
matmul_cublaslt(dv, scratch, att, nullptr, HS, T, T, stream, false, true, B * NH, T * HS, T * T, T * HS);
const float scale = 1.0f / sqrtf((float)HS);
// backward into preatt. this is an in-place operation; datt turns into dpreatt here
softmax_autoregressive_backward_inplace_kernel<<<dim3(T / 4, B * NH), 256>>>(datt, att, B, T, C, scale);
softmax_autoregressive_backward_inplace_kernel<<<dim3(T / 4, B * NH), 256, 0, stream>>>(datt, att, B, T, C, scale);
const floatX* dpreatt = datt;
// backward into q
matmul_cublaslt(dq, k, dpreatt, nullptr, HS, T, T, stream, false, false, B * NH, T * HS, T * T, T * HS);