summaryrefslogtreecommitdiff
path: root/llama.cpp/ggml/src/ggml-sycl/cpy.hpp
diff options
context:
space:
mode:
Diffstat (limited to 'llama.cpp/ggml/src/ggml-sycl/cpy.hpp')
-rw-r--r--llama.cpp/ggml/src/ggml-sycl/cpy.hpp223
1 files changed, 223 insertions, 0 deletions
diff --git a/llama.cpp/ggml/src/ggml-sycl/cpy.hpp b/llama.cpp/ggml/src/ggml-sycl/cpy.hpp
new file mode 100644
index 0000000..3c331f1
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-sycl/cpy.hpp
@@ -0,0 +1,223 @@
+#ifndef GGML_SYCL_CPY_HPP
+#define GGML_SYCL_CPY_HPP
+
+#include "common.hpp"
+#include <float.h>
+
+typedef void (*cpy_kernel_t)(const char * cx, char * cdst);
+
+__dpct_inline__ int best_index_int8(int n, const int8_t * val, float x) {
+ if (x <= val[0]) {
+ return 0;
+ }
+ if (x >= val[n - 1]) {
+ return n - 1;
+ }
+ int ml = 0, mu = n - 1;
+ while (mu - ml > 1) {
+ int mav = (ml + mu) / 2;
+ if (x < val[mav]) {
+ mu = mav;
+ } else {
+ ml = mav;
+ }
+ }
+ return x - val[mu - 1] < val[mu] - x ? mu - 1 : mu;
+}
+
+inline void cpy_blck_f32_q8_0(const char * cxi, char * cdsti) {
+ const float * xi = (const float *) cxi;
+ block_q8_0 * dsti = (block_q8_0 *) cdsti;
+
+ float amax = 0.0f; // absolute max
+
+ for (int j = 0; j < QK8_0; j++) {
+ const float v = xi[j];
+ amax = sycl::fmax(amax, sycl::fabs((float) v));
+ }
+
+ const float d = amax / ((1 << 7) - 1);
+ const float id = d ? 1.0f / d : 0.0f;
+
+ dsti->d = d;
+
+ for (int j = 0; j < QK8_0; ++j) {
+ const float x0 = xi[j] * id;
+
+ dsti->qs[j] = sycl::round((float) x0);
+ }
+}
+
+inline void cpy_blck_f32_q4_0(const char * cxi, char * cdsti) {
+ const float * xi = (const float *) cxi;
+ block_q4_0 * dsti = (block_q4_0 *) cdsti;
+
+ float amax = 0.0f;
+ float vmax = 0.0f;
+
+ for (int j = 0; j < QK4_0; ++j) {
+ const float v = xi[j];
+ if (amax < sycl::fabs((float) v)) {
+ amax = sycl::fabs((float) v);
+ vmax = v;
+ }
+ }
+
+ const float d = vmax / -8;
+ const float id = d ? 1.0f / d : 0.0f;
+
+ dsti->d = d;
+
+ for (int j = 0; j < QK4_0 / 2; ++j) {
+ const float x0 = xi[0 + j] * id;
+ const float x1 = xi[QK4_0 / 2 + j] * id;
+
+ const uint8_t xi0 = dpct::min(15, (int8_t) (x0 + 8.5f));
+ const uint8_t xi1 = dpct::min(15, (int8_t) (x1 + 8.5f));
+
+ dsti->qs[j] = xi0;
+ dsti->qs[j] |= xi1 << 4;
+ }
+}
+
+inline void cpy_blck_f32_q4_1(const char * cxi, char * cdsti) {
+ const float * xi = (const float *) cxi;
+ block_q4_1 * dsti = (block_q4_1 *) cdsti;
+
+ float vmin = FLT_MAX;
+ float vmax = -FLT_MAX;
+
+ for (int j = 0; j < QK4_1; ++j) {
+ const float v = xi[j];
+
+ vmin = sycl::min(v, vmin);
+ vmax = sycl::max(v, vmax);
+ }
+
+ const float d = (vmax - vmin) / ((1 << 4) - 1);
+ const float id = d ? 1.0f / d : 0.0f;
+
+ dsti->dm.x() = d;
+ dsti->dm.y() = vmin;
+
+ for (int j = 0; j < QK4_1 / 2; ++j) {
+ const float x0 = (xi[0 + j] - vmin) * id;
+ const float x1 = (xi[QK4_1 / 2 + j] - vmin) * id;
+
+ const uint8_t xi0 = dpct::min(15, (int8_t) (x0 + 0.5f));
+ const uint8_t xi1 = dpct::min(15, (int8_t) (x1 + 0.5f));
+
+ dsti->qs[j] = xi0;
+ dsti->qs[j] |= xi1 << 4;
+ }
+}
+
+inline void cpy_blck_f32_q5_0(const char * cxi, char * cdsti) {
+ const float * xi = (const float *) cxi;
+ block_q5_0 * dsti = (block_q5_0 *) cdsti;
+
+ float amax = 0.0f;
+ float vmax = 0.0f;
+
+ for (int j = 0; j < QK5_0; ++j) {
+ const float v = xi[j];
+ if (amax < sycl::fabs((float) v)) {
+ amax = sycl::fabs((float) v);
+ vmax = v;
+ }
+ }
+
+ const float d = vmax / -16;
+ const float id = d ? 1.0f / d : 0.0f;
+
+ dsti->d = d;
+
+ uint32_t qh = 0;
+ for (int j = 0; j < QK5_0 / 2; ++j) {
+ const float x0 = xi[0 + j] * id;
+ const float x1 = xi[QK5_0 / 2 + j] * id;
+
+ const uint8_t xi0 = dpct::min(31, (int8_t) (x0 + 16.5f));
+ const uint8_t xi1 = dpct::min(31, (int8_t) (x1 + 16.5f));
+
+ dsti->qs[j] = (xi0 & 0xf) | ((xi1 & 0xf) << 4);
+ qh |= ((xi0 & 0x10u) >> 4) << (j + 0);
+ qh |= ((xi1 & 0x10u) >> 4) << (j + QK5_0 / 2);
+ }
+ memcpy(dsti->qh, &qh, sizeof(qh));
+}
+
+inline void cpy_blck_f32_q5_1(const char * cxi, char * cdsti) {
+ const float * xi = (const float *) cxi;
+ block_q5_1 * dsti = (block_q5_1 *) cdsti;
+
+ float min = xi[0];
+ float max = xi[0];
+
+ for (int j = 1; j < QK5_1; ++j) {
+ const float v = xi[j];
+ min = v < min ? v : min;
+ max = v > max ? v : max;
+ }
+
+ const float d = (max - min) / 31;
+ const float id = d ? 1.0f / d : 0.0f;
+
+ dsti->dm.x() = d;
+ dsti->dm.y() = min;
+
+ uint32_t qh = 0;
+ for (int j = 0; j < QK5_1 / 2; ++j) {
+ const float x0 = (xi[0 + j] - min) * id;
+ const float x1 = (xi[QK5_1 / 2 + j] - min) * id;
+
+ const uint8_t xi0 = (uint8_t) (x0 + 0.5f);
+ const uint8_t xi1 = (uint8_t) (x1 + 0.5f);
+
+ dsti->qs[j] = (xi0 & 0xf) | ((xi1 & 0xf) << 4);
+ qh |= ((xi0 & 0x10u) >> 4) << (j + 0);
+ qh |= ((xi1 & 0x10u) >> 4) << (j + QK5_1 / 2);
+ }
+ memcpy(dsti->qh, &qh, sizeof(qh));
+}
+
+inline void cpy_blck_f32_iq4_nl(const char * cxi, char * cdsti) {
+ const float * xi = (const float *) cxi;
+ block_iq4_nl * dsti = (block_iq4_nl *) cdsti;
+
+ float amax = 0.0f;
+ float vmax = 0.0f;
+
+ for (int j = 0; j < QK4_NL; ++j) {
+ const float v = xi[j];
+ if (amax < sycl::fabs((float) v)) {
+ amax = sycl::fabs((float) v);
+ vmax = v;
+ }
+ }
+
+ float d = vmax / kvalues_iq4nl[0];
+ const float id = d ? 1.0f / d : 0.0f;
+
+ float sumqx = 0, sumq2 = 0;
+ for (int j = 0; j < QK4_NL / 2; ++j) {
+ const float x0 = xi[0 + j] * id;
+ const float x1 = xi[QK4_NL / 2 + j] * id;
+ const uint8_t xi0 = best_index_int8(16, kvalues_iq4nl, x0);
+ const uint8_t xi1 = best_index_int8(16, kvalues_iq4nl, x1);
+ dsti->qs[j] = xi0 | (xi1 << 4);
+ const float v0 = kvalues_iq4nl[xi0];
+ const float v1 = kvalues_iq4nl[xi1];
+ const float w0 = xi[0 + j] * xi[0 + j];
+ const float w1 = xi[QK4_NL / 2 + j] * xi[QK4_NL / 2 + j];
+ sumqx += w0 * v0 * xi[j] + w1 * v1 * xi[QK4_NL / 2 + j];
+ sumq2 += w0 * v0 * v0 + w1 * v1 * v1;
+ }
+
+ dsti->d = sumq2 > 0 ? sumqx / sumq2 : d;
+}
+
+void ggml_sycl_cpy(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1);
+void ggml_sycl_dup(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+#endif // GGML_SYCL_CPY_HPP
generated by cgit v1.2.3 (git 2.46.0) at 2026-04-29 13:29:27 +0000