CUDA: add COL2IM_1D op#25151
Conversation
github-actions
Bot
added
documentation
|
Hi @Ssamdeman, thanks for your contribution! Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:
Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below. |
|
Thanks. I clarified it more! |
|
remove the ops.csv changes |
Implements GGML_OP_COL2IM_1D on the CUDA backend (inverse of IM2COL), the missing counterpart to the merged CPU reference. Output-centric gather kernel templated over f32/f16/bf16 with fp32 accumulation. Passes 33/33 test-backend-ops cases; full regression suite clean.
Ssamdeman
force-pushed
the
cuda-col2im1d-op
branch
from
7437374 to
56418ba
Compare
|
|
||
|
|
||
|
|
| const int64_t oc = i / T_out; | ||
| const int64_t t_out = i % T_out; |
There was a problem hiding this comment.
not sure how performance critical this is but you can use fast_div
There was a problem hiding this comment.
yes I found it. I digged deeper inside the ggml/src/ggml-cuda/common.cuh. I can implement it new faster but it does not mean it is a good idea. here is my analysis. The fast_div_modulo(it calls fastdiv inside of it) speeds up the kernel by avoiding expensive hardware division, but it requires casting thread IDs to 32-bit, Given our 1D tensor use cases, is this 32-bit ceiling a safe tradeoff for the performance boost, or should I stick to standard 64-bit division?
There was a problem hiding this comment.
Okay not required for this PR anyway
am17an
left a comment
am17an
left a comment
There was a problem hiding this comment.
Just some extra whitespace
| const int64_t oc = i / T_out; | ||
| const int64_t t_out = i % T_out; |
There was a problem hiding this comment.
Okay not required for this PR anyway
| col2im_1d_kernel<float><<<num_blocks, block_size, 0, stream>>>( | ||
| (const float *)src->data, (float *)dst->data, T_out, T_in, K, OC, K_OC, s0, p0); | ||
| } else if (src->type == GGML_TYPE_F16) { | ||
| col2im_1d_kernel<half><<<num_blocks, block_size, 0, stream >>>( |
There was a problem hiding this comment.
| col2im_1d_kernel<half><<<num_blocks, block_size, 0, stream >>>( | |
| col2im_1d_kernel<half><<<num_blocks, block_size, 0, stream>>>( |
| @@ -0,0 +1,79 @@ | |||
| #include "col2im_1d.cuh" | |||
|
|
|||
|
|
|||
JohannesGaessler
left a comment
JohannesGaessler
left a comment
There was a problem hiding this comment.
The performance could likely be optimized but the implementation seems logically correct.
3 participants
Overview
Implements
GGML_OP_COL2IM_1Don the CUDA backend — the inverse of IM2COL and the missing counterpart to the CPU reference merged in #24206. Uses an output-centric gather kernel (no atomics), templated over f32/f16/bf16 with fp32 accumulation. The host wrapper readsop_params = [s0, OC, p0]and derivesK = src->ne[0] / OC.Testing:
test-backend-ops -o COL2IM_1D→ 33/33 OK across all dtype × kernel/stride/padding casestest-backend-opsregression clean (12868/12868)Additional information
The kernel math was verified line-by-line against the CPU
_implreference.col2im_1d.cu/.cuhare self-contained; the siblingim2colfiles are untouched.Requirements
I have read and agree with the contributing guidelines
AI usage disclosure:
AI usage disclosure: YES. I authored the code in this PR myself — I wrote the CUDA kernel, host wrapper, and backend wiring based on my own understanding of the col2im_1d math and the ggml backend structure. I used AI assistively in two ways: (1) to explain ggml architecture and help me verify my indexing/memory-layout assumptions against the merged CPU reference, and (2) when my tests failed, I gave the AI my code and the errors, and it helped me understand why they were failing and what to correct. I reviewed and applied the fixes myself, I understand every line in the final code, and I take full responsibility for all submitted changes.