sycl : kernel placeholders to make the CI green

ggerganov · Mar 12, 2024 · af93a92 · af93a92
1 parent 4f74b87
commit af93a92
Showing 1 changed file with 2 additions and 42 deletions.
diff --git a/ggml-sycl.cpp b/ggml-sycl.cpp
@@ -4894,27 +4894,8 @@ static void dequantize_block_iq1_s(const void * __restrict__ vx, dst_t * __restr
  const uint32_t *iq1s_grid,
  const uint8_t *ksigns_iq2xs,
  const uint8_t *kmask_iq2xs) {
- const int i = item_ct1.get_group(2);
- const block_iq1_s * x = (const block_iq1_s *) vx;
-
- const int tid = item_ct1.get_local_id(2);
-#if QK_K == 256
- const int il = tid/8; // 0...3
- const int ib = tid%8; // 0...7
- dst_t * y = yy + i*QK_K + 32*ib + 8*il;
- const float delta = x[i].qh[ib] & 0x8000 ? -1 - IQ1S_DELTA : -1 + IQ1S_DELTA;
- const float d = (float)x[i].d * (2*((x[i].qh[ib] >> 12) & 7) + 1);
- uint32_t grid32[2]; const int8_t * q = (const int8_t *)grid32;
- grid32[0] = iq1s_grid_gpu[x[i].qs[4*ib+il] | (((x[i].qh[ib] >> 3*il) & 7) << 8)];
- grid32[1] = (grid32[0] >> 4) & 0x0f0f0f0f;
- grid32[0] &= 0x0f0f0f0f;
- for (int j = 0; j < 4; ++j) {
- y[j] = d * (q[j] + delta);
- }
-#else
+ // TODO: implement
  assert(false);
-#endif
-
 }
 
 /*
@@ -7807,30 +7788,9 @@ static __dpct_inline__ float
 vec_dot_iq1_s_q8_1(const void *__restrict__ vbq,
  const block_q8_1 *__restrict__ bq8_1, const int &iqs,
  const uint32_t *iq1s_grid, const uint64_t *ksigns64) {
-#if QK_K == 256
- const block_iq1_s * bq1 = (const block_iq1_s *) vbq;
-
- const int ib32 = iqs;
- const uint8_t * qs = bq1->qs + 4*ib32;
- const int8_t * q8 = bq8_1[ib32].qs;
- int sumi = 0;
- for (int l = 0; l < 4; ++l) {
- const uint32_t * grid = (const uint32_t *)(iq1s_grid + qs[l]);
- const uint32_t * signs = (const uint32_t *)(ksigns64 + (qs[l] >> 8));
- const int grid_l = dpct::vectorized_binary<sycl::uchar4>(
- grid[0] ^ signs[0], signs[0], std::minus<>());
- const int grid_h = dpct::vectorized_binary<sycl::uchar4>(
- grid[1] ^ signs[1], signs[1], std::minus<>());
- sumi = dpct::dp4a(grid_l, *((int *)q8 + 0), sumi);
- sumi = dpct::dp4a(grid_h, *((int *)q8 + 1), sumi);
- q8 += 8;
- }
- const float d = (float)bq1->d * bq8_1[ib32].ds[0] * 0.25f;
- return d * sumi;
-#else
+ // TODO: implement
  assert(false);
  return 0.f;
-#endif
 }
 
 template <int qk, int qr, int qi, bool need_sum, typename block_q_t, int mmq_x,