Add CUDA non-contiguous Unary Ops support #14639
Conversation
github-actions
bot
added
documentation
YavorGIvanov
force-pushed
the
feature/cuda-non-cont-unary-support
branch
from
c44bfde to
919ce38
Compare
| if (ggml_is_contiguous(src) && ggml_is_contiguous(dst_tensor)) { | ||
| unary_op_kernel<op><<<num_blocks, CUDA_NEG_BLOCK_SIZE, 0, stream>>>(x, dst, k); | ||
| } else { |
There was a problem hiding this comment.
Remove the contiguous path, it's no longer needed.
There was a problem hiding this comment.
I kept it as the performance of the simple cont kernel is obviously better. I thought you may prefer to still use the most optimal path in this case. I know in the big scheme of things these unary operations are a very small part of the inference time, but think it is good idea to not degrade cont perf in this case.
ABS(type=f32,ne_a=[256,256,3,1],v=0): 532415 runs - 1.88 us/run - 1536 kB/run - 778.95 GB/s
ABS(type=f32,ne_a=[256,256,3,1],v=1): 311220 runs - 3.24 us/run - 3070 kB/run - 903.14 GB/s
Here is example perf test using test-backend-ops on a H100 SXM5.
v=0 meaning contiguous and v=1 meaning non-contiguous.
Let me know whether you still want the cont path removed or you agree I should keep it for now.
There was a problem hiding this comment.
Sorry for the late reply, if you want to keep the contiguous path, add a template parameter to the noncontiguous kernel where you return early.
More generally, if you're concerned about the performance one thing you can try is replace the byte offsets with logical offsets (calculate these in host code and pass to the kernel). But I expect the impact on end-to-end performance to be negligible.
YavorGIvanov
force-pushed
the
feature/cuda-non-cont-unary-support
branch
from
1174a95 to
1752873
Compare
YavorGIvanov
force-pushed
the
feature/cuda-non-cont-unary-support
branch
from
1752873 to
64be8c5
Compare
|
@JohannesGaessler @am17an Tried to address all comments. |
|
@YavorGIvanov gentle ping |
No description provided.