cuda: support row-contiguous SUM_ROWS

[MohamedElashri]

Overview

This is straightforward, the PR extends CUDA GGML_OP_SUM_ROWS from fully contiguous tensors to F32 row-contiguous tensors. What is does is that it adds a stride-aware kernel that derives the logical i1, i2, and i3 coordinates for each row and calculates its address using nb[1..3]. It also preserve the current behavio by using the existing reduction kernel as the fast path for fully contiguous tensors for the use cases it would be more suitable for.

The validation/testing results is also complete (with focus on related stuff)

• SUM_ROWS: 10/10
• SUM: 7/7
• MEAN: 7/7

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure:

Yes, AI was used. I used an agent to explore the codebase for codebase patterns and inspect possible style or merge-conflict problems. It also was used to see the test log files. But in both cases, I did re-check everything manually.

[ggml-gh-bot]

Hi @MohamedElashri, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[am17an]

/bot review

[ggml-gh-bot]

Automated code review

Reviewing the CUDA GGML_OP_SUM_ROWS row-contiguous extension (ggml/src/ggml-cuda/sumrows.cu and the supports_op change in ggml-cuda.cu).

I read the PR description, the full sumrows.cu, the shared reduce_rows.cu kernel it copies from, the block_reduce/block_reduce_policy helpers, ggml_is_contiguous_rows in ggml.c, and the test_sum_rows cases in tests/test-backend-ops.cpp.

Correctness

The change is correct. Verified points:

• ggml_is_contiguous_rows only requires nb[0] == ggml_type_size(type), which is exactly the property the strided kernel relies on (x[i] stepping consecutive floats within a row). Good guard.
• Row index decomposition i1 = row % ne1; i2 = (row/ne1) % ne2; i3 = row/(ne1*ne2) matches nrows = ggml_nrows = ne1*ne2*ne3 and the row-major flattening of the output [1, ne1, ne2, ne3], so dst[row] lands at the right logical (i1,i2,i3). Consistent with the contiguous reduce_rows_f32 path, which also writes dst[blockIdx.x].
• The F32 src+dst gate added to supports_op is an improvement over the prior code (which claimed any contiguous SUM_ROWS and then asserted F32 inside the op, i.e. would crash on a non-F32 tensor). SUM_ROWS output is always the input's (F32) type, so no behavior regression for legitimate uses; non-F32 cases now correctly fall back instead of asserting.
• block_reduce is only valid for block sizes that are multiples of WARP_SIZE and <= 1024; the launch uses 512/128/32, all fine. shared_vals[32] covers up to 32 warps, adequate for the largest block (512 = 16 warps).
• The contiguous fast path is byte-identical to the prior implementation, so existing behavior is preserved when ggml_is_contiguous holds.

No blocking correctness or security issues.

Will slow the review

• Duplicated reduction body. sum_rows_f32_strided is a near-verbatim copy of reduce_rows_f32 (the num_unroll/temp/sum_temp/block_reduce dance) with only the address computation differing. Keeping two copies means the unroll/reduce logic has to be maintained in two places and a future tweak (e.g. to accumulation order) risks drifting the two paths out of sync. A cleaner approach is to generalize reduce_rows_f32 to take the row base pointer as a per-block argument (or accept byte strides and do x = (const float*)(x_ptr + blockIdx.x*stride) for the strided case), letting norm=false cover both with one kernel and no copy. If you prefer to keep a separate kernel, at minimum add a one-line provenance note ("copied from reduce_rows_f32, address via nb[1..3]") so the lineage is explicit. ggml/src/ggml-cuda/sumrows.cu:19-57.

• Missing perf data. Per the ggml/backend reviewer expectations, a new kernel should come with throughput numbers across realistic shapes (e.g. small ncols / large ne2*ne3 views that motivated this). The PR description only lists test pass counts. Please add even a short table (shapes, GB/s vs the contiguous path) so the cost of carrying a second kernel is justified. ggml/src/ggml-cuda/sumrows.cu:19-57.

Nits

• The contiguous branch inlined in ggml_cuda_op_sum_rows (sumrows.cu:81-95) duplicates the launch logic already encapsulated in sum_rows_f32_cuda. For the contiguous case you could just call sum_rows_f32_cuda(src0_d, dst_d, ncols, nrows, stream) to avoid two copies of the SM-occupancy heuristic diverging over time.
• The two pairs of branches (contiguous small/large and strided small/large) differ only in block_dims. The block_dims selection ((nrows/nsm) < 2 ? 512 : (ncols < 1024 ? 32 : 128)) could be computed once, then dispatched to the two kernels, which would also make the "preserve fast path" intent clearer. ggml/src/ggml-cuda/sumrows.cu:78-106.
• ASCII only: description text has a stray unicode issue ("behavio" truncation is fine, but the description body uses a unicode middle dot in the bullet test list) - not in the diff, so just a heads-up if you amend the description.

Tests: test_sum_rows cases at tests/test-backend-ops.cpp:9387-9396 already include permute (ggml_permute(...,0,2,3,1), keeps dim 0 innermost) and slice (nb[2]*2) variants that exercise the strided path, and the gradient path goes through them too. That coverage is the right shape; please just confirm those cases are actually selected/run on CUDA in CI rather than only on CPU.

This review was generated automatically by pi coding agent using zai-org/GLM-5.2. It may contain mistakes. Maintainers make the final call.

[ORippler]

Thanks for the effort!

1. Please add a test-case if possible (or confirm the tests already contained in the suite according to AI pass locally).
2. IIRC, reduce_rows_f32 (i.e. the contiguous case) was used also for GGML_OP_MEAN. How much work would it be to extend this op to support row-contiguous tensors as well?

[MohamedElashri]

Thanks for the effort!

1. Please add a test-case if possible (or confirm the tests already contained in the suite according to AI pass locally).

2. IIRC, reduce_rows_f32 (i.e. the contiguous case) was used also for GGML_OP_MEAN. How much work would it be to extend this op to support row-contiguous tensors as well?

Hi @ORippler, thanks

1- The existing SUM_ROWS permute/slice tests already covered row-contiguous views, But I added equivalent MEAN permute/slice test cases. so we not have suitable test coverage
2- It would be straightforward, I did extend ML_OP_MEAN to support row-contiguous tensors using the same shared kernel with norm=true

I ran the following

./build/bin/test-backend-ops test -b CUDA0 -o SUM_ROWS
  ./build/bin/test-backend-ops test -b CUDA0 -o MEAN
  ./build/bin/test-backend-ops test -b CUDA0 -o SUM
  ./build/bin/test-backend-ops test -b CUDA0

And the summary of results of the tests tried on RTX 3090 GPU (compute capability 8.6).

SUM_ROWS: 10/10 passed
  MEAN:     10/10 passed
  SUM:      7/7 passed
  Full CUDA0 suite: 13346/13346 passed

I did some performance tests on explicit-stride synthetic cases for contiguous, permute-like, and slice-like row-contiguous layouts.

./build/bin/test-backend-ops perf -b CUDA0 

Logical bandwidth below counts actual reduced row bytes plus output bytes. For sliced views this avoids inflating GB/s by counting skipped backing-buffer holes.

op	shape	layout	us/run	logical GB/s	vs contig
SUM_ROWS	8192 x 16384	contiguous	587.97	850.49	baseline
SUM_ROWS	8192 x 16384	permute view	588.40	849.87	+0.07%
SUM_ROWS	8192 x 16384	slice view	588.20	850.15	+0.04%
SUM_ROWS	128 x 32768	contiguous	23.63	666.40	baseline
SUM_ROWS	128 x 32768	permute view	27.23	578.30	+15.23%
SUM_ROWS	128 x 32768	slice view	26.35	597.61	+11.51%
MEAN	8192 x 16384	contiguous	587.88	850.62	baseline
MEAN	8192 x 16384	permute view	588.25	850.08	+0.06%
MEAN	8192 x 16384	slice view	588.18	850.18	+0.05%
MEAN	128 x 32768	contiguous	23.43	672.09	baseline
MEAN	128 x 32768	permute view	27.45	573.66	+17.16%
MEAN	128 x 32768	slice view	26.60	592.00	+13.53%

If you got sometime, please review it and I would be happy to do extra work to finalize the best implementation.

[ORippler]

Thanks