SYCL: Implement few same quantized type copy kernels #13739
Conversation
There was a problem hiding this comment.
This may be a good opportunity to revisit the copy kernel when the source and destination are of the same type and don't require any casting. These should use the same function that copies a number of byte without depending on the type itself. It would reduce the number of kernels and simplify the code. This can be done with sycl::queue::memcpy. What do you think?
|
@Rbiessy I think |
|
You're right, that wouldn't work if there are permutations. I didn't think too much about what would a permuted tensor with quantized type look like? Is it by design only permuting the blocks and never the values inside a block?
|
|
Yeah. Generally permuted tensors are non contiguous. If I find time, I will see if I can use memcpy to copy contiguous quantized tensors or not. (I think it is possible). I am marking this PR draft for now. |
1f6cc01 to
3cdc64b
Compare
This is for supporting kv cache defragmentation when quatized kv cache is used.
test-backend-opsseems to pass with this change.Need further testing before we can merge.