Fused fprop scale/bias iterator advances without the split-k factor and drifts from activation channels
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
Description
In the fused (scale/bias) fprop mainloop, the scale/bias vector iterator and the activation/filter iterators disagree on the per-R*S-cycle channel stride when split_k_slices > 1, so scale/bias drift to the wrong channel blocks.
- Activation iterator on gemm-k wrap (include/cutlass/conv/threadblock/conv2d_fprop_activation_tile_access_iterator_analytic.h ~line 216,
_optimizedanalogue ~314):filter_c_ += Shape::kColumn * problem_size_.split_k_slices; - Scale/bias iterator on the same wrap (include/cutlass/conv/threadblock/predicated_scale_bias_vector_access_iterator.h:247-262): calls
add_tile_offset(TensorCoord(1, 0)), which advances by a singleThreadblockShape::kContiguous, with no split-k factor.
Both start consistently at channel threadblock_tile_idx.k() * kK (implicit_gemm_convolution_fusion.h ~296-337), so the divergence accumulates once the loop wraps past RS. With interleaved slice partitioning, each additional RS cycle applies bias/scale from one kK-block earlier than the channels the activations actually carry. This needs split_k_slices > 1 and more than one C-channel block per slice (ceil(ceil(C/S)/kK) >= 2) to appear, e.g. split_k_slices=4, C=256, ThreadblockShape::kK=32, RS>=2: from the second cycle on, bias[64+j32] is applied where the activations come from channels 192+. The result is silent wrong numerics, not an OOB.
The same pattern exists in the conv3d fusion headers (conv3d_fprop_activation_tile_access_iterator_analytic/_optimized vs the shared scale/bias iterator). Wgrad fusion is unaffected: it loads scale/bias once, with no per-tile advance.
The device layer explicitly supports split-k for fused fprop (workspace handling in include/cutlass/conv/device/implicit_gemm_convolution_fusion.h ~148-175), so the combination is reachable through the public API.
This finding is from static verification of the two advance paths; I have not built an end-to-end failing kernel, but the stride mismatch is exactly the missing split_k_slices factor and the start offsets agree.
Suggested fix
Multiply the scale/bias iterator's wrap stride by problem_size.split_k_slices in advance(), mirroring the activation iterators.
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start with predicated_scale_bias_vector_access_iterator.h and compare its advance path with the 2D activation iterator headers named in the issue; then inspect the fused fprop device path in include/cutlass/conv/device/implicit_gemm_convolution_fusion.h. Check the corresponding conv3d iterator pair as well. Done means the scale/bias and activation iterators use matching split-k strides and split-k fused fprop numerics no longer drift.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- backend, performance
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 55/100