@@ -48,7 +48,7 @@ __global__ void lowrank_bw_v_kernel(
4848 const int64_t * __restrict__ nids, const int64_t * __restrict__ xv,
4949 float * __restrict__ grad_ext, float * __restrict__ p_lp, float * __restrict__ p_lq,
5050 int batch_size, int num_eblks, int block_size, int rank,
51- int n_ntiles, int tile_n, long ext_base, int TB ) {
51+ int n_ntiles, int tile_n, long ext_base, int TB , int accumulate ) {
5252
5353 const int re = blockIdx .x ; // row * num_eblks + edge block
5454 const int nt = blockIdx .y ; // node tile
@@ -118,11 +118,21 @@ __global__ void lowrank_bw_v_kernel(
118118 const long vo = (xv_v + (long )n * rank + r) * batch_size + b;
119119 const float v = ext[vo];
120120
121- // Each (n, r, b) is touched by exactly one thread (nids partition the nodes), so a plain store
122- // is safe and avoids reading a gradient buffer that is as large as `V` itself. Kept in the
123- // exponent rather than forming `f` first, so a tiny flow does not lose the factor.
121+ // Kept in the exponent rather than forming `f` first, so a tiny flow does not lose the factor.
124122 const float lfv = lf + v;
125- grad_ext[vo] = safe_exp_diff (lfv + lw, lt) - safe_exp_diff (lfv + la, lz);
123+ const float g = safe_exp_diff (lfv + lw, lt) - safe_exp_diff (lfv + la, lz);
124+
125+ // A plain STORE is correct when this node owns its factors: `nids` partition the nodes, so each
126+ // (n, r, b) is touched exactly once, and it avoids reading a gradient buffer as large as `V`.
127+ // When several nodes SHARE one factor pair, every one of them contributes to the same gradient,
128+ // so it has to accumulate. Layers run sequentially and each element is still written once per
129+ // layer, so a plain read-modify-write suffices -- no atomics -- and the buffer is zeroed once
130+ // per backward pass, which is the precondition this relies on.
131+ if (accumulate) {
132+ grad_ext[vo] += g;
133+ } else {
134+ grad_ext[vo] = g;
135+ }
126136
127137 lse_add (acc_p[w], safe_log_diff (lfv, lt));
128138 lse_add (acc_q[w], safe_log_diff (lfv, lz));
@@ -181,7 +191,7 @@ __global__ void lowrank_bw_u_kernel(
181191 const int64_t * __restrict__ cids, const int64_t * __restrict__ xu,
182192 float * __restrict__ grad_ext, float * __restrict__ element_flows,
183193 int batch_size, int num_edges, int num_eblks, int ch_block_size, int rank,
184- int tile_c, long ext_base, int TB ) {
194+ int tile_c, long ext_base, int TB , int accumulate ) {
185195
186196 const int re = blockIdx .x ;
187197 const int ct = blockIdx .y ;
@@ -214,8 +224,16 @@ __global__ void lowrank_bw_u_kernel(
214224 const float u = ext[uo];
215225
216226 // The positive term is both dLL/dU's first part and this child's share of the flow
217- grad_ext[uo] = expf (u + e + lp) - expf (u + lq);
218- lse_add (log_corr, u + e + lp);
227+ const float log_t1 = u + e + lp; // LOG space -- `log_corr` accumulates logs
228+ const float g = expf (log_t1) - expf (u + lq);
229+
230+ if (accumulate) { // see the note in pass A
231+ grad_ext[uo] += g;
232+ } else {
233+ grad_ext[uo] = g;
234+ }
235+
236+ lse_add (log_corr, log_t1);
219237 }
220238
221239 // Atomic because several parent blocks may share a child block: unlike `node_mars` in the forward,
@@ -266,7 +284,7 @@ void lowrank_backward(torch::Tensor node_flows, torch::Tensor element_flows,
266284 torch::Tensor p_lp, torch::Tensor p_lq,
267285 torch::Tensor log_p, torch::Tensor log_q,
268286 int64_t block_size, int64_t ch_block_size, int64_t rank, int64_t ext_base,
269- int64_t tile_n, int64_t tile_c, int64_t tb) {
287+ int64_t tile_n, int64_t tile_c, int64_t tb, bool accumulate ) {
270288
271289 const int batch_size = node_flows.size (1 );
272290 const int rows = xu.size (0 );
@@ -285,7 +303,7 @@ void lowrank_backward(torch::Tensor node_flows, torch::Tensor element_flows,
285303 nids.data_ptr <int64_t >(), xv.data_ptr <int64_t >(),
286304 grad_ext.data_ptr <float >(), p_lp.data_ptr <float >(), p_lq.data_ptr <float >(),
287305 batch_size, num_eblks, (int )block_size, (int )rank,
288- n_ntiles, (int )tile_n, (long )ext_base, (int )tb);
306+ n_ntiles, (int )tile_n, (long )ext_base, (int )tb, accumulate ? 1 : 0 );
289307
290308 const int n_outputs = rows * num_eblks * (int )rank * batch_size;
291309 lowrank_bw_pq_reduce_kernel<<<bw_cdiv(n_outputs * 32 , 256 ), 256 , 0 , stream>>> (
@@ -301,5 +319,5 @@ void lowrank_backward(torch::Tensor node_flows, torch::Tensor element_flows,
301319 cids.data_ptr <int64_t >(), xu.data_ptr <int64_t >(),
302320 grad_ext.data_ptr <float >(), element_flows.data_ptr <float >(),
303321 batch_size, num_edges, num_eblks, (int )ch_block_size, (int )rank,
304- (int )tile_c, (long )ext_base, (int )tb);
322+ (int )tile_c, (long )ext_base, (int )tb, accumulate ? 1 : 0 );
305323}
0 commit comments