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
CMakePresets.json
Outdated
| { "name": "x64-linux-gcc-debug", "inherits": [ "base", "x64-linux-gcc", "debug" ] }, | ||
| { "name": "x64-linux-gcc-release", "inherits": [ "base", "x64-linux-gcc", "release" ] }, | ||
| { "name": "x64-linux-gcc-reldbg", "inherits": [ "base", "x64-linux-gcc", "reldbg" ] }, | ||
| { "name": "x64-linux-gcc+static-release", "inherits": [ "base", "x64-linux-gcc", "release", "static" ] }, |
There was a problem hiding this comment.
No. Should I separate it another PR c4ecdef
I am fine with removing it, but I did not see a preset that fit my use case and decided to add.
There was a problem hiding this comment.
Maybe easier to merge if you separate into another PR
There was a problem hiding this comment.
Please put it into a separate PR.
CMakePresets.json
Outdated
| { "name": "x64-linux-gcc-debug", "inherits": [ "base", "x64-linux-gcc", "debug" ] }, | ||
| { "name": "x64-linux-gcc-release", "inherits": [ "base", "x64-linux-gcc", "release" ] }, | ||
| { "name": "x64-linux-gcc-reldbg", "inherits": [ "base", "x64-linux-gcc", "reldbg" ] }, | ||
| { "name": "x64-linux-gcc+static-release", "inherits": [ "base", "x64-linux-gcc", "release", "static" ] }, |
There was a problem hiding this comment.
Please put it into a separate PR.
docs/ops/CUDA.csv
Outdated
There was a problem hiding this comment.
What is this file? Did you add it by accident?
There was a problem hiding this comment.
@JohannesGaessler recent merge #14598, in subsequent PRs we'll work out how to have such a huge diff when merging. Currently it records the timestamp, device etc so it becomes an entirely new file
There was a problem hiding this comment.
@YavorGIvanov For now don't commit the docs/ops/CUDA.csv and docs/ops.md. I'll make a follow-up PR after this gets merged to update the ops table.
There was a problem hiding this comment.
I am fine with improving and simplifying process of generating the docs/ops.md to not produce huge diffs myself.
ggml/src/ggml-cuda/unary.cu
Outdated
| const int k) { | ||
|
|
||
| const int i = blockDim.x*blockIdx.x + threadIdx.x; |
There was a problem hiding this comment.
| const int k) { | |
| const int i = blockDim.x*blockIdx.x + threadIdx.x; | |
| const int64_t k) { | |
| const int64_t i = blockDim.x*blockIdx.x + threadIdx.x; |
There was a problem hiding this comment.
Thanks. Applied as part of other PR review changes.
| 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.
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. |
No description provided.