SYCL: Add mrope kernel (llama/13755)

* SYCL: Add mrope kernel

* feat: Optimize rope operations with vectorization

Uses `sycl::vec` to load and store two elements at a time,
significantly improving performance in `rope_norm`,
`rope_neox`, and `rope_multi`. This reduces the number of memory
accesses and leverages SIMD instructions for faster execution.

* Use ceil_div
This commit is contained in:
Akarshan Biswas
2025-05-30 19:40:57 +05:30
committed by Georgi Gerganov
parent 1893359cfd
commit f7f92d0aab
2 changed files with 118 additions and 19 deletions

View File

@ -4257,14 +4257,6 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_OP_SOFT_MAX:
return true;
case GGML_OP_ROPE:
{
const int mode = ((const int32_t *) op->op_params)[2];
// mode is not used as a bitmask in practice, the various rope type modes are independent implementations
if (mode == GGML_ROPE_TYPE_MROPE) {
return false;
}
return true;
}
case GGML_OP_IM2COL:
return true;
case GGML_OP_UPSCALE: