cuda : add conv3d with implicit GEMM

[leejet]

Overview

Add CUDA support for GGML_OP_CONV_3D, used by ggml_conv_3d_direct, with an F16 implicit-GEMM kernel and a direct fallback for F32 weights and shapes outside the fast path. Inputs and outputs are contiguous F32 tensors; weights can be F16 or F32.

There are already CUDA conv3d proposals, including #16948, #17255, and #24569. This PR offers a compact implementation intended to make the indexing, bounds checks, and performance tradeoffs easier to review and maintain. It reuses mma.cuh, uses one fixed 64x64x64 tile shape, and keeps the convolution indexing and dispatch in one source file. The goal is to lower the review burden around this operation.

Additional information

Tested on Windows with an RTX 4090, CUDA 12.4, driver 596.36, and a Release build:

IC -> OC	Output D x H x W	Kernel D x H x W	im2col (ms)	direct (ms)	im2col/direct
384 -> 384	1 x 38 x 26	3 x 3 x 3	0.2107	0.1753	1.202x
192 -> 192	2 x 76 x 52	3 x 3 x 3	0.6857	0.2645	2.593x
96 -> 96	4 x 152 x 104	3 x 3 x 3	2.9104	0.6659	4.371x
96 -> 3	4 x 304 x 208	3 x 3 x 3	9.9162	1.3327	7.441x
1024 -> 1024	1 x 38 x 26	3 x 3 x 3	0.6331	0.9144	0.692x
256 -> 256	4 x 152 x 104	3 x 3 x 3	7.3020	3.0387	2.403x
384 -> 768	1 x 38 x 26	3 x 1 x 1	0.0634	0.0773	0.819x
192 -> 96	2 x 76 x 52	1 x 1 x 1	0.0563	0.0342	1.644x

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES. Codex assisted with porting the implementation from stable-diffusion.cpp's ggml fork.

[ggml-gh-bot]

Hi @leejet, thanks for your contribution!

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

• Multiple open PRs from a new contributor: We limit new contributors (those without a previously merged PR) to 1 open PR at a time. You currently have 7 open PRs.

---

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

[leejet]

@ggerganov @am17an This is the Conv3D version, following a similar path to the Conv2D implicit GEMM implementation #29135. Could you take a look at this one as well?

[ggerganov]

I can't make a detailed review, but I think it is OK to merge as the chances to impact llama.cpp functionality is zero and it would be useful for downstream projects to get the extra performance.

[ggerganov]

/bot review

[ggml-gh-bot]

Automated code review

Static review of ggml/src/ggml-cuda/conv3d.cu/.cuh, the dispatch hooks in ggml-cuda.cu, and the new test-backend-ops cases. I verified the op_params layout against ggml_conv_3d_direct (src[0]=kernel, src[1]=input, p[0..2]=stride, p[3..5]=pad, p[6..8]=dilation, p[9..11]=IC/N/OC), the index decomposition in all four kernels (direct fallback, pad, implicit GEMM, split-k reduce), and the split-k partial-buffer layout — the writer's blockIdx.z * oc slice placement matches the reduce's n*(split_k-1)*per_batch + i base. The fast-path guards (int32 limits, alignment, 65535 grid caps, padded-bounds preconditions) look sound, and the mma.cuh tile selection matches the established fattn patterns (J_MAJOR accumulator + swapped get_i/get_j on AMD is equivalent to fattn's swapped consumption sites). No blocking issues found.

Will slow the review

(point 1) AMD/HIP path is untested — the PR only reports an RTX 4090 run, but the AMD-specific code is exactly what cannot be validated there: the DATA_LAYOUT_J_MAJOR accumulator transposition, RDNA3 I_MAJOR_MIRRORED A/B fragments, and the warp-64 load_lane/load_row mapping in the B-staging loop (also the SIMT fallback, which only compiles into a real code path on non-Turing NVIDIA / non-WMMA AMD). Please run test-backend-ops -o CONV_3D on the CI AMD runners (gfx1100/gfx1030/gfx90x) and the oldest NVIDIA arch you can, or state in the PR that it was checked.

(point 2) split_k can become 0: in the min-list, (IC* KW* KH* KD + 63) / 64 evaluates to 0 for a zero-extent kernel (e.g. KW == 0 is not rejected by the asserts, only ne[3] == IC*OC is checked). That gives blockIdx.z / split_k (divide by zero in the kernel) and a zero-size partial alloc. Clamp that term with a max(1, ...), or assert KW/KH/KD > 0.

(point 3) docs/ops.md still marks CONV_3D as unsupported on CUDA (line 32), and docs/ops/CUDA.csv has all CONV_3D rows at support=0. Please update the ops.md table in this PR (the CSV is regenerated by the test script, but ops.md is manual).

(point 4) Process note for maintainers: there are competing conv3d implementations in flight (#16948, #17255, #24569, all referenced in the description). Someone should decide which approach is preferred before this gets a deep review, so the review effort is not spent on the wrong one.

Nits

(point 5) #pragma unroll on the A/B staging loops (for (int i = tid; ... i += nthreads)) is a no-op — the trip count and start are runtime values. Drop them.

(point 6) conv3d_params is passed by value (22 x int64_t = 176 B) to every kernel; conv2d.cu passes const conv_params &. Consider a const ref for consistency and marginally smaller kernel args.

(point 7) A one-line comment in conv3d_reduce_split_k explaining the partial layout (the GEMM stores split k of batch n at (n*split_k + k)*per_batch via blockIdx.z, and the reduce's n*(split_k-1)*per_batch + i base is derived from that) would save the next reader a nontrivial derivation — the two kernels' index math is coupled and currently only one side is commented (via the store expression).

The new test cases are well chosen (OC=65 crossing the 64-tile boundary, stride/dilation/padding combos, shapes small enough to trigger split_k > 1, both kernel dtypes), and reuse the existing test_conv_3d infrastructure rather than adding new test files. Perf data is included per the kernel contribution expectations. AI disclosure is filled in.

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

[ggerganov]

@JohannesGaessler About your earlier concern in #24569 (comment) - I think this is a great opportunity to get some progress on the convolution kernels. With this work, the kernels will get heavily exercised in @leejet's https://github.com/leejet/stable-diffusion.cpp project and chances are that any issues with correctness and performance will be quickly resolved over there.

[IMbackK]

an issue i have had trying to optimize the kernels mostly used by sdcpp (im2col) in the past is that sdcpp lacks an easy way to do end to end benchmarks sweeps, ie a llama-bench equivalent.

[JohannesGaessler]

Thank you for sorting this out on your end. The only thing that I think needs to be addressed prior to a merge is the misaligned pointer, the rest I'll leave at your discretion.

[leejet]

Thanks everyone for the reviews! I’ve updated the code based on the feedback, including some of the helpful suggestions from the bot review. When you have a chance, please take another look. Thanks!

[ggerganov]

@leejet Could you take a look at this error: https://github.com/ggml-org/llama.cpp/actions/runs/35969488472/job/107535592917#step:3:9865

[jeffbolznv]

@leejet Could you take a look at this error: https://github.com/ggml-org/llama.cpp/actions/runs/35969488472/job/107535592917#step:3:9865

We need to handle misaligned offsets in the vulkan backend. @0cc4m can you do this?

[0cc4m]

I'll look into it.

[leejet]

@leejet Could you take a look at this error: https://github.com/ggml-org/llama.cpp/actions/runs/35969488472/job/107535592917#step:3:9865

Thanks for the heads-up! Looks like this has already been picked up and there's a fix PR open now.