opencl: add initial iq3 support (#31338)
* opencl: add iq3_s AOS mv Assisted-by: Codex * opencl: add iq3_s SOA mv and tiled mm Assisted-by: Codex * opencl: add iq3_s noshuffle gemm/gemv for adreno Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com> Assisted-by: Codex * opencl: add iq3_s dp4a for adreno Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com> Assisted-by: Codex * opencl: optimize iq3_s adreno decode and narrow batches Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com> Assisted-by: Codex * opencl: add iq3_xxs AOS mv Assisted-by: Codex * opencl: add iq3_xxs SOA mv and tiled mm Assisted-by: Codex * opencl: add iq3_xxs noshuffle gemm/gemv for adreno Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com> Assisted-by: Codex * opencl: add iq3_xxs dp4a for adreno Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com> Assisted-by: Codex * opencl: optimize iq3_xxs adreno decode and narrow batches Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com> Assisted-by: Codex --------- Co-authored-by: Li He <lih@qti.qualcomm.com>
This commit is contained in:
shaofeiqiandLi He authored and 
GitHub committed
1 parent
a91eab67d8
commit
060b2fea4e
17 files changed
+3216
-25
No files matched your search
@@ -124,6 +124,12 @@ set(GGML_OPENCL_KERNELS
|
||||
gemm_noshuffle_iq1_s_f32
|
||||
gemm_noshuffle_iq1_s_q8_1_dp4a
|
||||
gemv_noshuffle_iq1_s_f32
|
||||
gemm_noshuffle_iq3_s_f32
|
||||
gemm_noshuffle_iq3_s_q8_1_dp4a
|
||||
gemv_noshuffle_iq3_s_f32
|
||||
gemm_noshuffle_iq3_xxs_f32
|
||||
gemm_noshuffle_iq3_xxs_q8_1_dp4a
|
||||
gemv_noshuffle_iq3_xxs_f32
|
||||
gemm_noshuffle_q3_k_q8_1_dp4a
|
||||
gemv_noshuffle_q3_k_f32
|
||||
mul_mv_q2_k_f32
|
||||
@@ -150,6 +156,14 @@ set(GGML_OPENCL_KERNELS
|
||||
mul_mv_iq1_s_f32
|
||||
mul_mv_iq1_s_f32_flat
|
||||
mul_mm_iq1_s_f32_l4_lm
|
||||
iq3_xxs_grid
|
||||
mul_mv_iq3_xxs_f32
|
||||
mul_mv_iq3_xxs_f32_flat
|
||||
mul_mm_iq3_xxs_f32_l4_lm
|
||||
iq3_s_grid
|
||||
mul_mv_iq3_s_f32
|
||||
mul_mv_iq3_s_f32_flat
|
||||
mul_mm_iq3_s_f32_l4_lm
|
||||
mul_mv_q3_k_f32_flat
|
||||
mul_mm_q3_k_f32_l4_lm
|
||||
mul_mv_q2_k_f32_flat
|
||||
|
||||
File diff suppressed because it is too large.
Load diff
@@ -3178,3 +3178,108 @@ kernel void kernel_restore_block_iq1_m(
|
||||
dst[i].scales[2*j + 1] = (uchar)(sc >> 8);
|
||||
}
|
||||
}
|
||||
|
||||
struct block_iq3_s {
|
||||
half d;
|
||||
uchar qs[QK_K/4];
|
||||
uchar qh[QK_K/32];
|
||||
uchar signs[QK_K/8];
|
||||
uchar scales[QK_K/64];
|
||||
};
|
||||
|
||||
kernel void kernel_convert_block_iq3_s(
|
||||
global struct block_iq3_s * src0,
|
||||
global uchar * dst_qs,
|
||||
global uchar * dst_sg,
|
||||
global uchar * dst_qh,
|
||||
global uchar * dst_sc,
|
||||
global half * dst_d,
|
||||
ulong n_blocks
|
||||
) {
|
||||
const ulong i = get_global_id(0);
|
||||
if (i >= n_blocks) { return; }
|
||||
dst_d[i] = src0[i].d;
|
||||
for (int j = 0; j < 64; ++j) {
|
||||
dst_qs[64*i + j] = src0[i].qs[j];
|
||||
}
|
||||
for (int j = 0; j < 32; ++j) {
|
||||
dst_sg[32*i + j] = src0[i].signs[j];
|
||||
}
|
||||
for (int j = 0; j < 8; ++j) {
|
||||
dst_qh[8*i + j] = src0[i].qh[j];
|
||||
}
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
dst_sc[4*i + j] = src0[i].scales[j];
|
||||
}
|
||||
}
|
||||
|
||||
kernel void kernel_restore_block_iq3_s(
|
||||
global uchar * src_qs,
|
||||
global uchar * src_sg,
|
||||
global uchar * src_qh,
|
||||
global uchar * src_sc,
|
||||
global half * src_d,
|
||||
global struct block_iq3_s * dst,
|
||||
ulong n_blocks
|
||||
) {
|
||||
const ulong i = get_global_id(0);
|
||||
if (i >= n_blocks) { return; }
|
||||
dst[i].d = src_d[i];
|
||||
for (int j = 0; j < 64; ++j) {
|
||||
dst[i].qs[j] = src_qs[64*i + j];
|
||||
}
|
||||
for (int j = 0; j < 32; ++j) {
|
||||
dst[i].signs[j] = src_sg[32*i + j];
|
||||
}
|
||||
for (int j = 0; j < 8; ++j) {
|
||||
dst[i].qh[j] = src_qh[8*i + j];
|
||||
}
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
dst[i].scales[j] = src_sc[4*i + j];
|
||||
}
|
||||
}
|
||||
|
||||
struct block_iq3_xxs {
|
||||
half d;
|
||||
uchar qs[3*QK_K/8];
|
||||
};
|
||||
|
||||
kernel void kernel_convert_block_iq3_xxs(
|
||||
global struct block_iq3_xxs * src0,
|
||||
global uchar * dst_qs,
|
||||
global uint * dst_sas,
|
||||
global half * dst_d,
|
||||
ulong n_blocks
|
||||
) {
|
||||
const ulong i = get_global_id(0);
|
||||
if (i >= n_blocks) { return; }
|
||||
dst_d[i] = src0[i].d;
|
||||
for (int j = 0; j < 64; ++j) {
|
||||
dst_qs[64*i + j] = src0[i].qs[j];
|
||||
}
|
||||
for (int j = 0; j < 8; ++j) {
|
||||
global ushort * q = (global ushort *)(src0[i].qs + 64 + 4*j);
|
||||
dst_sas[8*i + j] = (uint)q[0] | ((uint)q[1] << 16);
|
||||
}
|
||||
}
|
||||
|
||||
kernel void kernel_restore_block_iq3_xxs(
|
||||
global uchar * src_qs,
|
||||
global uint * src_sas,
|
||||
global half * src_d,
|
||||
global struct block_iq3_xxs * dst,
|
||||
ulong n_blocks
|
||||
) {
|
||||
const ulong i = get_global_id(0);
|
||||
if (i >= n_blocks) { return; }
|
||||
dst[i].d = src_d[i];
|
||||
for (int j = 0; j < 64; ++j) {
|
||||
dst[i].qs[j] = src_qs[64*i + j];
|
||||
}
|
||||
for (int j = 0; j < 8; ++j) {
|
||||
global ushort * q = (global ushort *)(dst[i].qs + 64 + 4*j);
|
||||
uint sas = src_sas[8*i + j];
|
||||
q[0] = (ushort)sas;
|
||||
q[1] = (ushort)(sas >> 16);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,92 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#ifdef cl_qcom_reqd_sub_group_size
|
||||
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
|
||||
__attribute__((qcom_reqd_sub_group_size("full")))
|
||||
#endif
|
||||
kernel void kernel_gemm_noshuffle_iq3_s_f32(
|
||||
global const uchar * src0_qs,
|
||||
global const uchar * src0_sg,
|
||||
global const uchar * src0_qh,
|
||||
global const uchar * src0_sc,
|
||||
global const half * src0_d,
|
||||
read_only image1d_buffer_t src1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int m,
|
||||
int n,
|
||||
int k,
|
||||
int n_no_padding
|
||||
) {
|
||||
dst = (global float *)((global char *)dst + offsetd);
|
||||
const int gy = get_global_id(0);
|
||||
const int row = get_global_id(1) * 4;
|
||||
const int n_4 = n / 4;
|
||||
if (row >= m) {
|
||||
return;
|
||||
}
|
||||
|
||||
half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0;
|
||||
for (int i = 0; i < k; i += 16) {
|
||||
const uchar4 sc = vload4(0, src0_sc + row + (i / 64) * m);
|
||||
const half4 d = vload4(0, src0_d + row + (i / 256) * m);
|
||||
const float4 scale_f = convert_float4(d) * (1.0f + 2.0f*convert_float4((sc >> (uchar4)(4*((i / 32) & 1))) & (uchar4)(15)));
|
||||
const half4 scale = convert_half4(scale_f);
|
||||
|
||||
for (int l = 0; l < 16; l += 4) {
|
||||
const int ki = i + l;
|
||||
const uint4 qs = convert_uint4(vload4(0, src0_qs + row + (ki / 4) * m));
|
||||
const uint4 qh = convert_uint4(vload4(0, src0_qh + row + (ki / 32) * m));
|
||||
const uchar4 sg = vload4(0, src0_sg + row + (ki / 8) * m);
|
||||
const uint4 gi = qs | (((qh >> (uint4)(((ki / 4) & 7))) & (uint4)(1)) << (uint4)(8));
|
||||
const uint4 grid = (uint4)(iq3_s_grid[gi.s0], iq3_s_grid[gi.s1],
|
||||
iq3_s_grid[gi.s2], iq3_s_grid[gi.s3]);
|
||||
#pragma unroll
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
half8 b;
|
||||
b.s0123 = read_imageh(src1, gy * 2 + (ki + j) * n_4);
|
||||
b.s4567 = read_imageh(src1, gy * 2 + 1 + (ki + j) * n_4);
|
||||
const int4 magnitude = convert_int4((grid >> (uint4)(8*j)) & (uint4)(255));
|
||||
const int4 signs = convert_int4((sg >> (uchar4)((ki & 7) + j)) & (uchar4)(1));
|
||||
const half4 w = convert_half4(magnitude * ((int4)(1) - 2*signs)) * scale;
|
||||
c0 += b * w.s0;
|
||||
c1 += b * w.s1;
|
||||
c2 += b * w.s2;
|
||||
c3 += b * w.s3;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int idx = gy * 8 * m + row;
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s0, c1.s0, c2.s0, c3.s0), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s1, c1.s1, c2.s1, c3.s1), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s2, c1.s2, c2.s2, c3.s2), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s3, c1.s3, c2.s3, c3.s3), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s4, c1.s4, c2.s4, c3.s4), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s5, c1.s5, c2.s5, c3.s5), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s6, c1.s6, c2.s6, c3.s6), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s7, c1.s7, c2.s7, c3.s7), 0, dst + idx);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,154 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
|
||||
#ifdef cl_khr_integer_dot_product
|
||||
#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable
|
||||
#endif
|
||||
|
||||
#define QK_K 256
|
||||
|
||||
#ifndef TILESIZE_N
|
||||
#define TILESIZE_N 32
|
||||
#endif
|
||||
|
||||
inline uint iq3_s_pack(uint gv, uint sg, uint base) {
|
||||
int v0 = (int)((gv >> 0) & 0xFF); if (sg & (1u << (base + 0))) { v0 = -v0; }
|
||||
int v1 = (int)((gv >> 8) & 0xFF); if (sg & (1u << (base + 1))) { v1 = -v1; }
|
||||
int v2 = (int)((gv >> 16) & 0xFF); if (sg & (1u << (base + 2))) { v2 = -v2; }
|
||||
int v3 = (int)((gv >> 24) & 0xFF); if (sg & (1u << (base + 3))) { v3 = -v3; }
|
||||
return ((uint)v0 & 0xFFu) | (((uint)v1 & 0xFFu) << 8)
|
||||
| (((uint)v2 & 0xFFu) << 16) | (((uint)v3 & 0xFFu) << 24);
|
||||
}
|
||||
|
||||
#define IQ3_S_GEMM_GRID(i) iq3_s_grid[(i)]
|
||||
|
||||
inline uint iq3_s_load(__global const uchar * qs, uint qsb, uint qhb, uint m,
|
||||
uint u, uint sgv, uint base) {
|
||||
const uint gidx = (uint)qs[qsb + u * m] | (((qhb >> u) & 1u) << 8);
|
||||
return iq3_s_pack(IQ3_S_GEMM_GRID(gidx), sgv, base);
|
||||
}
|
||||
|
||||
// Pass staged operands by value; local dp4a operands miscompile on X2.
|
||||
inline int dot8_q8a_v(uint8 qw, uint4 a0, uint4 a1) {
|
||||
int r = 0;
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s0, a0.x, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s1, a0.y, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s2, a0.z, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s3, a0.w, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s4, a1.x, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s5, a1.y, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s6, a1.z, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s7, a1.w, r);
|
||||
return r;
|
||||
}
|
||||
|
||||
__attribute__((qcom_wave_pair_mode(1)))
|
||||
kernel void kernel_gemm_noshuffle_iq3_s_q8_1_dp4a(
|
||||
__global const uchar * src0_qs,
|
||||
__global const uchar * src0_sg,
|
||||
__global const uchar * src0_qh,
|
||||
__global const uchar * src0_sc,
|
||||
__global const half * src0_d,
|
||||
__global const uint * src1_qa,
|
||||
__global const half * src1_da,
|
||||
__global float * dst,
|
||||
ulong offsetd,
|
||||
int m,
|
||||
int n_no_padding,
|
||||
int k
|
||||
) {
|
||||
dst = (global float *)((global char *)dst + offsetd);
|
||||
|
||||
const uint lid = get_local_id(0);
|
||||
const uint block_id_m = get_global_id(1);
|
||||
const uint block_id_n = get_global_id(2);
|
||||
|
||||
const uint row = block_id_m * 64 + lid;
|
||||
const uint col_base = block_id_n * TILESIZE_N;
|
||||
const bool row_valid = row < (uint)m;
|
||||
const uint rrow = row_valid ? row : 0;
|
||||
|
||||
const uint k_u = (uint)k >> 2;
|
||||
const uint k_b = (uint)k >> 5;
|
||||
|
||||
__local uint4 sh_qa4[TILESIZE_N][2];
|
||||
__local half sh_d[TILESIZE_N];
|
||||
|
||||
#define NGROUPS (TILESIZE_N / 4)
|
||||
float4 acc[NGROUPS];
|
||||
#pragma unroll
|
||||
for (int g = 0; g < NGROUPS; ++g) acc[g] = (float4)(0.0f);
|
||||
|
||||
for (uint step = 0; step < (uint)k; step += 32) {
|
||||
const uint sub = step >> 5;
|
||||
|
||||
const uint ib = sub >> 3;
|
||||
const uint sb = sub & 7u;
|
||||
|
||||
const uint scb = (uint)src0_sc[rrow + (ib*(QK_K/64) + (sb >> 1)) * (uint)m];
|
||||
const uint nib = (sb & 1u) ? (scb >> 4) : (scb & 0xFu);
|
||||
const float d_w = (float)src0_d[rrow + ib * (uint)m] * (float)(1 + 2*nib);
|
||||
|
||||
const uint qhb = (uint)src0_qh[rrow + (ib*(QK_K/32) + sb) * (uint)m];
|
||||
const uint qsb = rrow + (ib*(QK_K/4) + 8*sb) * (uint)m;
|
||||
const uint sgb = rrow + (ib*(QK_K/8) + 4*sb) * (uint)m;
|
||||
|
||||
const uint sg0 = (uint)src0_sg[sgb + 0 * (uint)m];
|
||||
const uint sg1 = (uint)src0_sg[sgb + 1 * (uint)m];
|
||||
const uint sg2 = (uint)src0_sg[sgb + 2 * (uint)m];
|
||||
const uint sg3 = (uint)src0_sg[sgb + 3 * (uint)m];
|
||||
|
||||
uint8 qw;
|
||||
qw.s0 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 0u, sg0, 0u);
|
||||
qw.s1 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 1u, sg0, 4u);
|
||||
qw.s2 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 2u, sg1, 0u);
|
||||
qw.s3 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 3u, sg1, 4u);
|
||||
qw.s4 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 4u, sg2, 0u);
|
||||
qw.s5 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 5u, sg2, 4u);
|
||||
qw.s6 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 6u, sg3, 0u);
|
||||
qw.s7 = iq3_s_load(src0_qs, qsb, qhb, (uint)m, 7u, sg3, 4u);
|
||||
|
||||
for (uint idx = lid; idx < TILESIZE_N * 2; idx += 64) {
|
||||
const uint t = idx >> 1;
|
||||
const uint v = idx & 1;
|
||||
const uint c = col_base + t;
|
||||
sh_qa4[t][v] = (c < (uint)n_no_padding)
|
||||
? vload4(0, src1_qa + c * k_u + (step >> 2) + (v << 2))
|
||||
: (uint4)(0u);
|
||||
}
|
||||
if (lid < TILESIZE_N) {
|
||||
const uint c = col_base + lid;
|
||||
sh_d[lid] = (c < (uint)n_no_padding) ? src1_da[c * k_b + sub] : (half)0;
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
#define LD4(arr, b) ((float4)((float)arr[(b)+0], (float)arr[(b)+1], (float)arr[(b)+2], (float)arr[(b)+3]))
|
||||
#pragma unroll
|
||||
for (int g = 0; g < NGROUPS; ++g) {
|
||||
const int b = g * 4;
|
||||
float4 rf;
|
||||
#define TDOT4(T) dot8_q8a_v(qw, (uint4)(sh_qa4[T][0]), (uint4)(sh_qa4[T][1]))
|
||||
rf.s0 = (float)TDOT4(b+0); rf.s1 = (float)TDOT4(b+1);
|
||||
rf.s2 = (float)TDOT4(b+2); rf.s3 = (float)TDOT4(b+3);
|
||||
acc[g] += d_w * LD4(sh_d, b) * rf;
|
||||
}
|
||||
#undef TDOT4
|
||||
#undef LD4
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
if (!row_valid) {
|
||||
return;
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for (int g = 0; g < NGROUPS; ++g) {
|
||||
const uint b = (uint)(g * 4);
|
||||
const float4 a = acc[g];
|
||||
const uint c0 = col_base + b;
|
||||
if (c0 + 0 < (uint)n_no_padding) dst[(c0 + 0) * (uint)m + row] = a.s0;
|
||||
if (c0 + 1 < (uint)n_no_padding) dst[(c0 + 1) * (uint)m + row] = a.s1;
|
||||
if (c0 + 2 < (uint)n_no_padding) dst[(c0 + 2) * (uint)m + row] = a.s2;
|
||||
if (c0 + 3 < (uint)n_no_padding) dst[(c0 + 3) * (uint)m + row] = a.s3;
|
||||
}
|
||||
#undef NGROUPS
|
||||
}
|
||||
@@ -0,0 +1,89 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#ifdef cl_qcom_reqd_sub_group_size
|
||||
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
|
||||
__attribute__((qcom_reqd_sub_group_size("full")))
|
||||
#endif
|
||||
kernel void kernel_gemm_noshuffle_iq3_xxs_f32(
|
||||
global const uchar * src0_qs,
|
||||
global const uint * src0_sas,
|
||||
global const half * src0_d,
|
||||
read_only image1d_buffer_t src1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int m,
|
||||
int n,
|
||||
int k,
|
||||
int n_no_padding
|
||||
) {
|
||||
dst = (global float *)((global char *)dst + offsetd);
|
||||
const int gy = get_global_id(0);
|
||||
const int row = get_global_id(1) * 4;
|
||||
const int n_4 = n / 4;
|
||||
if (row >= m) {
|
||||
return;
|
||||
}
|
||||
|
||||
half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0;
|
||||
for (int i = 0; i < k; i += 16) {
|
||||
const uint4 sas = vload4(0, src0_sas + row + (i / 32) * m);
|
||||
const half4 d = vload4(0, src0_d + row + (i / 256) * m);
|
||||
const float4 scale_f = convert_float4(d) * 0.5f * (0.5f + convert_float4(sas >> (uint4)(28)));
|
||||
const half4 scale = convert_half4(scale_f);
|
||||
|
||||
for (int l = 0; l < 16; l += 4) {
|
||||
const int ki = i + l;
|
||||
const uint4 gi = convert_uint4(vload4(0, src0_qs + row + (ki / 4) * m));
|
||||
const uint4 code = (sas >> (uint4)(7*((ki / 8) & 3))) & (uint4)(127);
|
||||
const uint4 sg = code | ((popcount(code) & (uint4)(1)) << (uint4)(7));
|
||||
const uint4 grid = (uint4)(iq3_xxs_grid[gi.s0], iq3_xxs_grid[gi.s1],
|
||||
iq3_xxs_grid[gi.s2], iq3_xxs_grid[gi.s3]);
|
||||
#pragma unroll
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
half8 b;
|
||||
b.s0123 = read_imageh(src1, gy * 2 + (ki + j) * n_4);
|
||||
b.s4567 = read_imageh(src1, gy * 2 + 1 + (ki + j) * n_4);
|
||||
const int4 magnitude = convert_int4((grid >> (uint4)(8*j)) & (uint4)(255));
|
||||
const int4 signs = convert_int4((sg >> (uint4)((ki & 7) + j)) & (uint4)(1));
|
||||
const half4 w = convert_half4(magnitude * ((int4)(1) - 2*signs)) * scale;
|
||||
c0 += b * w.s0;
|
||||
c1 += b * w.s1;
|
||||
c2 += b * w.s2;
|
||||
c3 += b * w.s3;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int idx = gy * 8 * m + row;
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s0, c1.s0, c2.s0, c3.s0), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s1, c1.s1, c2.s1, c3.s1), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s2, c1.s2, c2.s2, c3.s2), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s3, c1.s3, c2.s3, c3.s3), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s4, c1.s4, c2.s4, c3.s4), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s5, c1.s5, c2.s5, c3.s5), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s6, c1.s6, c2.s6, c3.s6), 0, dst + idx);
|
||||
idx += m;
|
||||
}
|
||||
if (idx + 3 < m * n_no_padding) {
|
||||
vstore4((float4)(c0.s7, c1.s7, c2.s7, c3.s7), 0, dst + idx);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,137 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
|
||||
#ifdef cl_khr_integer_dot_product
|
||||
#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable
|
||||
#endif
|
||||
|
||||
#define QK_K 256
|
||||
|
||||
#ifndef TILESIZE_N
|
||||
#define TILESIZE_N 32
|
||||
#endif
|
||||
|
||||
inline uint iq3_xxs_pack(uint gv, uint sg, uint base) {
|
||||
int v0 = (int)((gv >> 0) & 0xFF); if (sg & (1u << (base + 0))) { v0 = -v0; }
|
||||
int v1 = (int)((gv >> 8) & 0xFF); if (sg & (1u << (base + 1))) { v1 = -v1; }
|
||||
int v2 = (int)((gv >> 16) & 0xFF); if (sg & (1u << (base + 2))) { v2 = -v2; }
|
||||
int v3 = (int)((gv >> 24) & 0xFF); if (sg & (1u << (base + 3))) { v3 = -v3; }
|
||||
return ((uint)v0 & 0xFFu) | (((uint)v1 & 0xFFu) << 8)
|
||||
| (((uint)v2 & 0xFFu) << 16) | (((uint)v3 & 0xFFu) << 24);
|
||||
}
|
||||
|
||||
// Pass staged operands by value; local dp4a operands miscompile on X2.
|
||||
inline int dot8_q8a_v(uint8 qw, uint4 a0, uint4 a1) {
|
||||
int r = 0;
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s0, a0.x, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s1, a0.y, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s2, a0.z, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s3, a0.w, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s4, a1.x, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s5, a1.y, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s6, a1.z, r);
|
||||
r = dot_acc_sat_4x8packed_ss_int(qw.s7, a1.w, r);
|
||||
return r;
|
||||
}
|
||||
|
||||
__attribute__((qcom_wave_pair_mode(1)))
|
||||
kernel void kernel_gemm_noshuffle_iq3_xxs_q8_1_dp4a(
|
||||
__global const uchar * src0_qs,
|
||||
__global const uint * src0_sas,
|
||||
__global const half * src0_d,
|
||||
__global const uint * src1_qa,
|
||||
__global const half * src1_da,
|
||||
__global float * dst,
|
||||
ulong offsetd,
|
||||
int m,
|
||||
int n_no_padding,
|
||||
int k
|
||||
) {
|
||||
dst = (global float *)((global char *)dst + offsetd);
|
||||
|
||||
const uint lid = get_local_id(0);
|
||||
const uint block_id_m = get_global_id(1);
|
||||
const uint block_id_n = get_global_id(2);
|
||||
|
||||
const uint row = block_id_m * 64 + lid;
|
||||
const uint col_base = block_id_n * TILESIZE_N;
|
||||
const bool row_valid = row < (uint)m;
|
||||
const uint rrow = row_valid ? row : 0;
|
||||
|
||||
const uint k_u = (uint)k >> 2;
|
||||
const uint k_b = (uint)k >> 5;
|
||||
|
||||
__local uint4 sh_qa4[TILESIZE_N][2];
|
||||
__local half sh_d[TILESIZE_N];
|
||||
|
||||
#define NGROUPS (TILESIZE_N / 4)
|
||||
float4 acc[NGROUPS];
|
||||
#pragma unroll
|
||||
for (int g = 0; g < NGROUPS; ++g) acc[g] = (float4)(0.0f);
|
||||
|
||||
for (uint step = 0; step < (uint)k; step += 32) {
|
||||
const uint sub = step >> 5;
|
||||
|
||||
const uint ib = sub >> 3;
|
||||
const uint aux = src0_sas[rrow + sub*(uint)m];
|
||||
const float d_w = (float)src0_d[rrow + ib*(uint)m] * (0.5f + (float)(aux >> 28)) * 0.5f;
|
||||
const uint sg0 = iq3_xxs_signs( aux & 127u);
|
||||
const uint sg1 = iq3_xxs_signs((aux >> 7) & 127u);
|
||||
const uint sg2 = iq3_xxs_signs((aux >> 14) & 127u);
|
||||
const uint sg3 = iq3_xxs_signs((aux >> 21) & 127u);
|
||||
const uint qsb = rrow + sub*8u*(uint)m;
|
||||
|
||||
uint8 qw;
|
||||
qw.s0 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 0u*(uint)m]], sg0, 0u);
|
||||
qw.s1 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 1u*(uint)m]], sg0, 4u);
|
||||
qw.s2 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 2u*(uint)m]], sg1, 0u);
|
||||
qw.s3 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 3u*(uint)m]], sg1, 4u);
|
||||
qw.s4 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 4u*(uint)m]], sg2, 0u);
|
||||
qw.s5 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 5u*(uint)m]], sg2, 4u);
|
||||
qw.s6 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 6u*(uint)m]], sg3, 0u);
|
||||
qw.s7 = iq3_xxs_pack(iq3_xxs_grid[src0_qs[qsb + 7u*(uint)m]], sg3, 4u);
|
||||
|
||||
for (uint idx = lid; idx < TILESIZE_N * 2; idx += 64) {
|
||||
const uint t = idx >> 1;
|
||||
const uint v = idx & 1;
|
||||
const uint c = col_base + t;
|
||||
sh_qa4[t][v] = (c < (uint)n_no_padding)
|
||||
? vload4(0, src1_qa + c * k_u + (step >> 2) + (v << 2))
|
||||
: (uint4)(0u);
|
||||
}
|
||||
if (lid < TILESIZE_N) {
|
||||
const uint c = col_base + lid;
|
||||
sh_d[lid] = (c < (uint)n_no_padding) ? src1_da[c * k_b + sub] : (half)0;
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
#define LD4(arr, b) ((float4)((float)arr[(b)+0], (float)arr[(b)+1], (float)arr[(b)+2], (float)arr[(b)+3]))
|
||||
#pragma unroll
|
||||
for (int g = 0; g < NGROUPS; ++g) {
|
||||
const int b = g * 4;
|
||||
float4 rf;
|
||||
#define TDOT4(T) dot8_q8a_v(qw, (uint4)(sh_qa4[T][0]), (uint4)(sh_qa4[T][1]))
|
||||
rf.s0 = (float)TDOT4(b+0); rf.s1 = (float)TDOT4(b+1);
|
||||
rf.s2 = (float)TDOT4(b+2); rf.s3 = (float)TDOT4(b+3);
|
||||
acc[g] += d_w * LD4(sh_d, b) * rf;
|
||||
}
|
||||
#undef TDOT4
|
||||
#undef LD4
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
if (!row_valid) {
|
||||
return;
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for (int g = 0; g < NGROUPS; ++g) {
|
||||
const uint b = (uint)(g * 4);
|
||||
const float4 a = acc[g];
|
||||
const uint c0 = col_base + b;
|
||||
if (c0 + 0 < (uint)n_no_padding) dst[(c0 + 0) * (uint)m + row] = a.s0;
|
||||
if (c0 + 1 < (uint)n_no_padding) dst[(c0 + 1) * (uint)m + row] = a.s1;
|
||||
if (c0 + 2 < (uint)n_no_padding) dst[(c0 + 2) * (uint)m + row] = a.s2;
|
||||
if (c0 + 3 < (uint)n_no_padding) dst[(c0 + 3) * (uint)m + row] = a.s3;
|
||||
}
|
||||
#undef NGROUPS
|
||||
}
|
||||
@@ -0,0 +1,128 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#define QK_K 256
|
||||
|
||||
#ifndef IQ3_S_MV_NSG
|
||||
#define IQ3_S_MV_NSG 4
|
||||
#endif
|
||||
|
||||
inline float4 iq3_s_vals(uint gv, uint sg, uint base) {
|
||||
|
||||
const uint s = sg >> base;
|
||||
const uint4 sgn = (uint4)(s << 31, s << 30, s << 29, s << 28) & 0x80000000u;
|
||||
return as_float4(as_uint4(convert_float4(as_uchar4(gv))) ^ sgn);
|
||||
}
|
||||
|
||||
kernel void kernel_gemv_noshuffle_iq3_s_f32(
|
||||
__read_only image1d_buffer_t grid_img,
|
||||
global const uchar * src0_qs,
|
||||
global const uchar * src0_sg,
|
||||
global const uchar * src0_qh,
|
||||
global const uchar * src0_sc,
|
||||
global const half * src0_d,
|
||||
global const float * src1,
|
||||
ulong offset1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
int ne10,
|
||||
int ne0
|
||||
) {
|
||||
src1 = (global const float *)((global const char *)src1 + offset1);
|
||||
dst = (global float *)((global char *)dst + offsetd);
|
||||
|
||||
const uint m = (uint)ne01;
|
||||
const uint K = (uint)ne00;
|
||||
const uint nsb = K / QK_K;
|
||||
|
||||
const uint lid = get_local_id(0);
|
||||
const uint sgi = get_local_id(1);
|
||||
const uint col = get_group_id(1);
|
||||
|
||||
global const float * y = src1 + (ulong)col * (uint)ne10;
|
||||
|
||||
#define IQ3_S_GRID(i) (read_imageui(grid_img, (int)(i)).x)
|
||||
|
||||
const uint mh = m >> 1;
|
||||
const uint j = get_group_id(0) * 64u + lid;
|
||||
const uint row = j << 1;
|
||||
|
||||
float sumf = 0.f;
|
||||
float sumf1 = 0.f;
|
||||
|
||||
if (j < mh) {
|
||||
global const ushort * qsu = (global const ushort *)src0_qs;
|
||||
global const ushort * qhu = (global const ushort *)src0_qh;
|
||||
global const ushort * sgu = (global const ushort *)src0_sg;
|
||||
global const ushort * scu = (global const ushort *)src0_sc;
|
||||
|
||||
for (uint ib = sgi; ib < nsb; ib += IQ3_S_MV_NSG) {
|
||||
const half2 dh = vload2(j + ib * mh, src0_d);
|
||||
const float d0 = (float)dh.s0;
|
||||
const float d1 = (float)dh.s1;
|
||||
|
||||
const uint scbase = j + ib * 4u * mh;
|
||||
ushort sc4[4];
|
||||
sc4[0] = scu[scbase + 0u * mh];
|
||||
sc4[1] = scu[scbase + 1u * mh];
|
||||
sc4[2] = scu[scbase + 2u * mh];
|
||||
sc4[3] = scu[scbase + 3u * mh];
|
||||
|
||||
float acc0 = 0.f, acc1 = 0.f;
|
||||
for (uint sb = 0; sb < 8u; ++sb) {
|
||||
const uint scv = (uint)sc4[sb >> 1];
|
||||
const uint nib0 = (sb & 1u) ? ((scv >> 4) & 0xFu) : ( scv & 0xFu);
|
||||
const uint nib1 = (sb & 1u) ? ((scv >> 12) & 0xFu) : ((scv >> 8) & 0xFu);
|
||||
|
||||
const uint qhv = (uint)qhu[j + (ib * 8u + sb) * mh];
|
||||
const uint qh0 = qhv & 0xFFu;
|
||||
const uint qh1 = qhv >> 8;
|
||||
|
||||
const uint grp = ib * 64u + sb * 8u;
|
||||
const uint qsb = j + grp * mh;
|
||||
const uint sgb = j + (ib * 32u + sb * 4u) * mh;
|
||||
|
||||
float a0 = 0.f, a1 = 0.f;
|
||||
for (uint u = 0; u < 8u; ++u) {
|
||||
const uint qsv = (uint)qsu[qsb + u * mh];
|
||||
const uint sgv = (uint)sgu[sgb + (u >> 1) * mh];
|
||||
const uint g0 = ( qsv & 0xFFu) | (((qh0 >> u) & 1u) << 8);
|
||||
const uint g1 = ((qsv >> 8) & 0xFFu) | (((qh1 >> u) & 1u) << 8);
|
||||
const uint base = (u & 1u) * 4u;
|
||||
const float4 yv = vload4(grp + u, y);
|
||||
const uint gw0 = IQ3_S_GRID(g0);
|
||||
const uint gw1 = IQ3_S_GRID(g1);
|
||||
a0 += dot(yv, iq3_s_vals(gw0, sgv & 0xFFu, base));
|
||||
a1 += dot(yv, iq3_s_vals(gw1, (sgv >> 8) & 0xFFu, base));
|
||||
}
|
||||
acc0 += (float)(1u + 2u * nib0) * a0;
|
||||
acc1 += (float)(1u + 2u * nib1) * a1;
|
||||
}
|
||||
sumf += d0 * acc0;
|
||||
sumf1 += d1 * acc1;
|
||||
}
|
||||
}
|
||||
|
||||
#if IQ3_S_MV_NSG > 1
|
||||
__local float2 part[IQ3_S_MV_NSG][64];
|
||||
part[sgi][lid] = (float2)(sumf, sumf1);
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
if (sgi != 0) {
|
||||
return;
|
||||
}
|
||||
for (uint s = 1; s < IQ3_S_MV_NSG; ++s) {
|
||||
const float2 p = part[s][lid];
|
||||
sumf += p.s0;
|
||||
sumf1 += p.s1;
|
||||
}
|
||||
#endif
|
||||
|
||||
if (j < mh) {
|
||||
vstore2((float2)(sumf, sumf1), 0, dst + (ulong)col * (uint)ne0 + row);
|
||||
}
|
||||
}
|
||||
|
||||
kernel void kernel_iq3_s_grid_export(global uint * out) {
|
||||
out[get_global_id(0)] = iq3_s_grid[get_global_id(0)];
|
||||
}
|
||||
@@ -0,0 +1,120 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#define QK_K 256
|
||||
|
||||
#ifndef IQ3_XXS_MV_NSG
|
||||
#define IQ3_XXS_MV_NSG 4
|
||||
#endif
|
||||
|
||||
inline float4 iq3_xxs_vals(uint gv, uint sgv, uint base) {
|
||||
const uint s = sgv >> base;
|
||||
const uint4 sgn = (uint4)(s << 31, s << 30, s << 29, s << 28) & 0x80000000u;
|
||||
return as_float4(as_uint4(convert_float4(as_uchar4(gv))) ^ sgn);
|
||||
}
|
||||
|
||||
kernel void kernel_gemv_noshuffle_iq3_xxs_f32(
|
||||
__read_only image1d_buffer_t grid_img,
|
||||
global const uchar * src0_qs,
|
||||
global const uint * src0_sas,
|
||||
global const half * src0_d,
|
||||
global const float * src1,
|
||||
ulong offset1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
int ne10,
|
||||
int ne0
|
||||
) {
|
||||
src1 = (global const float *)((global const char *)src1 + offset1);
|
||||
dst = (global float *)((global char *)dst + offsetd);
|
||||
|
||||
const uint m = (uint)ne01;
|
||||
const uint K = (uint)ne00;
|
||||
const uint nsb = K / QK_K;
|
||||
|
||||
const uint lid = get_local_id(0);
|
||||
const uint sgi = get_local_id(1);
|
||||
const uint col = get_group_id(1);
|
||||
|
||||
global const float * y = src1 + (ulong)col * (uint)ne10;
|
||||
|
||||
#define IQ3_XXS_GRID(i) (read_imageui(grid_img, (int)(i)).x)
|
||||
|
||||
const uint mh = m >> 1;
|
||||
const uint j = get_group_id(0) * 64u + lid;
|
||||
const uint row = j << 1;
|
||||
|
||||
float sumf = 0.f;
|
||||
float sumf1 = 0.f;
|
||||
|
||||
if (j < mh) {
|
||||
global const ushort * qsu = (global const ushort *)src0_qs;
|
||||
|
||||
for (uint ib = sgi; ib < nsb; ib += IQ3_XXS_MV_NSG) {
|
||||
const half2 dh = vload2(j + ib * mh, src0_d);
|
||||
const float d0 = (float)dh.s0;
|
||||
const float d1 = (float)dh.s1;
|
||||
|
||||
float acc0 = 0.f, acc1 = 0.f;
|
||||
for (uint sb = 0; sb < 8u; ++sb) {
|
||||
const uint sub = ib * 8u + sb;
|
||||
const uint2 aux = vload2(j + sub * mh, src0_sas);
|
||||
|
||||
const uint grp = sub * 8u;
|
||||
const uint qsb = j + grp * mh;
|
||||
|
||||
float a0 = 0.f, a1 = 0.f;
|
||||
for (uint u = 0; u < 4u; ++u) {
|
||||
const uint qsv0 = (uint)qsu[qsb + (2u*u)*mh];
|
||||
const uint qsv1 = (uint)qsu[qsb + (2u*u + 1u)*mh];
|
||||
const uint sg0 = iq3_xxs_signs((aux.s0 >> (7u * u)) & 127u);
|
||||
const uint sg1 = iq3_xxs_signs((aux.s1 >> (7u * u)) & 127u);
|
||||
|
||||
const uint g00 = qsv0 & 0xFFu;
|
||||
const uint g01 = qsv1 & 0xFFu;
|
||||
const uint g10 = (qsv0 >> 8) & 0xFFu;
|
||||
const uint g11 = (qsv1 >> 8) & 0xFFu;
|
||||
|
||||
const float4 yl = vload4(sub * 8u + u * 2u + 0u, y);
|
||||
const float4 yh = vload4(sub * 8u + u * 2u + 1u, y);
|
||||
|
||||
a0 += dot(yl, iq3_xxs_vals(IQ3_XXS_GRID(g00), sg0, 0u))
|
||||
+ dot(yh, iq3_xxs_vals(IQ3_XXS_GRID(g01), sg0, 4u));
|
||||
a1 += dot(yl, iq3_xxs_vals(IQ3_XXS_GRID(g10), sg1, 0u))
|
||||
+ dot(yh, iq3_xxs_vals(IQ3_XXS_GRID(g11), sg1, 4u));
|
||||
}
|
||||
acc0 += (0.5f + (float)(aux.s0 >> 28)) * a0;
|
||||
acc1 += (0.5f + (float)(aux.s1 >> 28)) * a1;
|
||||
}
|
||||
sumf += d0 * 0.5f * acc0;
|
||||
sumf1 += d1 * 0.5f * acc1;
|
||||
}
|
||||
}
|
||||
|
||||
#if IQ3_XXS_MV_NSG > 1
|
||||
__local float2 part[IQ3_XXS_MV_NSG][64];
|
||||
part[sgi][lid] = (float2)(sumf, sumf1);
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
if (sgi != 0) {
|
||||
return;
|
||||
}
|
||||
for (uint s = 1; s < IQ3_XXS_MV_NSG; ++s) {
|
||||
const float2 p = part[s][lid];
|
||||
sumf += p.s0;
|
||||
sumf1 += p.s1;
|
||||
}
|
||||
#endif
|
||||
|
||||
if (j < mh) {
|
||||
vstore2((float2)(sumf, sumf1), 0, dst + (ulong)col * (uint)ne0 + row);
|
||||
}
|
||||
#undef IQ3_XXS_GRID
|
||||
}
|
||||
|
||||
kernel void kernel_iq3_xxs_grid_export(global uint * out) {
|
||||
const uint i = get_global_id(0);
|
||||
if (i < 256u) {
|
||||
out[i] = iq3_xxs_grid[i];
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,66 @@
|
||||
constant uint iq3_s_grid[512] = {
|
||||
0x01010101, 0x01010103, 0x01010105, 0x0101010b, 0x0101010f, 0x01010301, 0x01010303, 0x01010305,
|
||||
0x01010309, 0x0101030d, 0x01010501, 0x01010503, 0x0101050b, 0x01010707, 0x01010901, 0x01010905,
|
||||
0x0101090b, 0x0101090f, 0x01010b03, 0x01010b07, 0x01010d01, 0x01010d05, 0x01010f03, 0x01010f09,
|
||||
0x01010f0f, 0x01030101, 0x01030103, 0x01030105, 0x01030109, 0x01030301, 0x01030303, 0x0103030b,
|
||||
0x01030501, 0x01030507, 0x0103050f, 0x01030703, 0x0103070b, 0x01030909, 0x01030d03, 0x01030d0b,
|
||||
0x01030f05, 0x01050101, 0x01050103, 0x0105010b, 0x0105010f, 0x01050301, 0x01050307, 0x0105030d,
|
||||
0x01050503, 0x0105050b, 0x01050701, 0x01050709, 0x01050905, 0x0105090b, 0x0105090f, 0x01050b03,
|
||||
0x01050b07, 0x01050f01, 0x01050f07, 0x01070107, 0x01070303, 0x0107030b, 0x01070501, 0x01070505,
|
||||
0x01070703, 0x01070707, 0x0107070d, 0x01070909, 0x01070b01, 0x01070b05, 0x01070d0f, 0x01070f03,
|
||||
0x01070f0b, 0x01090101, 0x01090307, 0x0109030f, 0x01090503, 0x01090509, 0x01090705, 0x01090901,
|
||||
0x01090907, 0x01090b03, 0x01090f01, 0x010b0105, 0x010b0109, 0x010b0501, 0x010b0505, 0x010b050d,
|
||||
0x010b0707, 0x010b0903, 0x010b090b, 0x010b090f, 0x010b0d0d, 0x010b0f07, 0x010d010d, 0x010d0303,
|
||||
0x010d0307, 0x010d0703, 0x010d0b05, 0x010d0f03, 0x010f0101, 0x010f0105, 0x010f0109, 0x010f0501,
|
||||
0x010f0505, 0x010f050d, 0x010f0707, 0x010f0b01, 0x010f0b09, 0x03010101, 0x03010103, 0x03010105,
|
||||
0x03010109, 0x03010301, 0x03010303, 0x03010307, 0x0301030b, 0x0301030f, 0x03010501, 0x03010505,
|
||||
0x03010703, 0x03010709, 0x0301070d, 0x03010b09, 0x03010b0d, 0x03010d03, 0x03010f05, 0x03030101,
|
||||
0x03030103, 0x03030107, 0x0303010d, 0x03030301, 0x03030309, 0x03030503, 0x03030701, 0x03030707,
|
||||
0x03030903, 0x03030b01, 0x03030b05, 0x03030f01, 0x03030f0d, 0x03050101, 0x03050305, 0x0305030b,
|
||||
0x0305030f, 0x03050501, 0x03050509, 0x03050705, 0x03050901, 0x03050907, 0x03050b0b, 0x03050d01,
|
||||
0x03050f05, 0x03070103, 0x03070109, 0x0307010f, 0x03070301, 0x03070307, 0x03070503, 0x0307050f,
|
||||
0x03070701, 0x03070709, 0x03070903, 0x03070d05, 0x03070f01, 0x03090107, 0x0309010b, 0x03090305,
|
||||
0x03090309, 0x03090703, 0x03090707, 0x03090905, 0x0309090d, 0x03090b01, 0x03090b09, 0x030b0103,
|
||||
0x030b0301, 0x030b0307, 0x030b0503, 0x030b0701, 0x030b0705, 0x030b0b03, 0x030d0501, 0x030d0509,
|
||||
0x030d050f, 0x030d0909, 0x030d090d, 0x030f0103, 0x030f0107, 0x030f0301, 0x030f0305, 0x030f0503,
|
||||
0x030f070b, 0x030f0903, 0x030f0d05, 0x030f0f01, 0x05010101, 0x05010103, 0x05010107, 0x0501010b,
|
||||
0x0501010f, 0x05010301, 0x05010305, 0x05010309, 0x0501030d, 0x05010503, 0x05010507, 0x0501050f,
|
||||
0x05010701, 0x05010705, 0x05010903, 0x05010907, 0x0501090b, 0x05010b01, 0x05010b05, 0x05010d0f,
|
||||
0x05010f01, 0x05010f07, 0x05010f0b, 0x05030101, 0x05030105, 0x05030301, 0x05030307, 0x0503030f,
|
||||
0x05030505, 0x0503050b, 0x05030703, 0x05030709, 0x05030905, 0x05030b03, 0x05050103, 0x05050109,
|
||||
0x0505010f, 0x05050503, 0x05050507, 0x05050701, 0x0505070f, 0x05050903, 0x05050b07, 0x05050b0f,
|
||||
0x05050f03, 0x05050f09, 0x05070101, 0x05070105, 0x0507010b, 0x05070303, 0x05070505, 0x05070509,
|
||||
0x05070703, 0x05070707, 0x05070905, 0x05070b01, 0x05070d0d, 0x05090103, 0x0509010f, 0x05090501,
|
||||
0x05090507, 0x05090705, 0x0509070b, 0x05090903, 0x05090f05, 0x05090f0b, 0x050b0109, 0x050b0303,
|
||||
0x050b0505, 0x050b070f, 0x050b0901, 0x050b0b07, 0x050b0f01, 0x050d0101, 0x050d0105, 0x050d010f,
|
||||
0x050d0503, 0x050d0b0b, 0x050d0d03, 0x050f010b, 0x050f0303, 0x050f050d, 0x050f0701, 0x050f0907,
|
||||
0x050f0b01, 0x07010105, 0x07010303, 0x07010307, 0x0701030b, 0x0701030f, 0x07010505, 0x07010703,
|
||||
0x07010707, 0x0701070b, 0x07010905, 0x07010909, 0x0701090f, 0x07010b03, 0x07010d07, 0x07010f03,
|
||||
0x07030103, 0x07030107, 0x0703010b, 0x07030309, 0x07030503, 0x07030507, 0x07030901, 0x07030d01,
|
||||
0x07030f05, 0x07030f0d, 0x07050101, 0x07050305, 0x07050501, 0x07050705, 0x07050709, 0x07050b01,
|
||||
0x07070103, 0x07070301, 0x07070309, 0x07070503, 0x07070507, 0x0707050f, 0x07070701, 0x07070903,
|
||||
0x07070907, 0x0707090f, 0x07070b0b, 0x07070f07, 0x07090107, 0x07090303, 0x0709030d, 0x07090505,
|
||||
0x07090703, 0x07090b05, 0x07090d01, 0x07090d09, 0x070b0103, 0x070b0301, 0x070b0305, 0x070b050b,
|
||||
0x070b0705, 0x070b0909, 0x070b0b0d, 0x070b0f07, 0x070d030d, 0x070d0903, 0x070f0103, 0x070f0107,
|
||||
0x070f0501, 0x070f0505, 0x070f070b, 0x09010101, 0x09010109, 0x09010305, 0x09010501, 0x09010509,
|
||||
0x0901050f, 0x09010705, 0x09010903, 0x09010b01, 0x09010f01, 0x09030105, 0x0903010f, 0x09030303,
|
||||
0x09030307, 0x09030505, 0x09030701, 0x0903070b, 0x09030907, 0x09030b03, 0x09030b0b, 0x09050103,
|
||||
0x09050107, 0x09050301, 0x0905030b, 0x09050503, 0x09050707, 0x09050901, 0x09050b0f, 0x09050d05,
|
||||
0x09050f01, 0x09070109, 0x09070303, 0x09070307, 0x09070501, 0x09070505, 0x09070703, 0x0907070b,
|
||||
0x09090101, 0x09090105, 0x09090509, 0x0909070f, 0x09090901, 0x09090f03, 0x090b010b, 0x090b010f,
|
||||
0x090b0503, 0x090b0d05, 0x090d0307, 0x090d0709, 0x090d0d01, 0x090f0301, 0x090f030b, 0x090f0701,
|
||||
0x090f0907, 0x090f0b03, 0x0b010105, 0x0b010301, 0x0b010309, 0x0b010505, 0x0b010901, 0x0b010909,
|
||||
0x0b01090f, 0x0b010b05, 0x0b010d0d, 0x0b010f09, 0x0b030103, 0x0b030107, 0x0b03010b, 0x0b030305,
|
||||
0x0b030503, 0x0b030705, 0x0b030f05, 0x0b050101, 0x0b050303, 0x0b050507, 0x0b050701, 0x0b05070d,
|
||||
0x0b050b07, 0x0b070105, 0x0b07010f, 0x0b070301, 0x0b07050f, 0x0b070909, 0x0b070b03, 0x0b070d0b,
|
||||
0x0b070f07, 0x0b090103, 0x0b090109, 0x0b090501, 0x0b090705, 0x0b09090d, 0x0b0b0305, 0x0b0b050d,
|
||||
0x0b0b0b03, 0x0b0b0b07, 0x0b0d0905, 0x0b0f0105, 0x0b0f0109, 0x0b0f0505, 0x0d010303, 0x0d010307,
|
||||
0x0d01030b, 0x0d010703, 0x0d010707, 0x0d010d01, 0x0d030101, 0x0d030501, 0x0d03050f, 0x0d030d09,
|
||||
0x0d050305, 0x0d050709, 0x0d050905, 0x0d050b0b, 0x0d050d05, 0x0d050f01, 0x0d070101, 0x0d070309,
|
||||
0x0d070503, 0x0d070901, 0x0d09050b, 0x0d090907, 0x0d090d05, 0x0d0b0101, 0x0d0b0107, 0x0d0b0709,
|
||||
0x0d0b0d01, 0x0d0d010b, 0x0d0d0901, 0x0d0f0303, 0x0d0f0307, 0x0f010101, 0x0f010109, 0x0f01010f,
|
||||
0x0f010501, 0x0f010505, 0x0f01070d, 0x0f010901, 0x0f010b09, 0x0f010d05, 0x0f030105, 0x0f030303,
|
||||
0x0f030509, 0x0f030907, 0x0f03090b, 0x0f050103, 0x0f050109, 0x0f050301, 0x0f05030d, 0x0f050503,
|
||||
0x0f050701, 0x0f050b03, 0x0f070105, 0x0f070705, 0x0f07070b, 0x0f070b07, 0x0f090103, 0x0f09010b,
|
||||
0x0f090307, 0x0f090501, 0x0f090b01, 0x0f0b0505, 0x0f0b0905, 0x0f0d0105, 0x0f0d0703, 0x0f0f0101
|
||||
};
|
||||
@@ -0,0 +1,38 @@
|
||||
constant uint iq3_xxs_grid[256] = {
|
||||
0x04040404, 0x04040414, 0x04040424, 0x04040c0c, 0x04040c1c, 0x04040c3e, 0x04041404, 0x04041414,
|
||||
0x04041c0c, 0x04042414, 0x04043e1c, 0x04043e2c, 0x040c040c, 0x040c041c, 0x040c0c04, 0x040c0c14,
|
||||
0x040c140c, 0x040c142c, 0x040c1c04, 0x040c1c14, 0x040c240c, 0x040c2c24, 0x040c3e04, 0x04140404,
|
||||
0x04140414, 0x04140424, 0x04140c0c, 0x04141404, 0x04141414, 0x04141c0c, 0x04141c1c, 0x04141c3e,
|
||||
0x04142c0c, 0x04142c3e, 0x04143e2c, 0x041c040c, 0x041c043e, 0x041c0c04, 0x041c0c14, 0x041c142c,
|
||||
0x041c3e04, 0x04240c1c, 0x04241c3e, 0x04242424, 0x04242c3e, 0x04243e1c, 0x04243e2c, 0x042c040c,
|
||||
0x042c043e, 0x042c1c14, 0x042c2c14, 0x04341c2c, 0x04343424, 0x043e0c04, 0x043e0c24, 0x043e0c34,
|
||||
0x043e241c, 0x043e340c, 0x0c04040c, 0x0c04041c, 0x0c040c04, 0x0c040c14, 0x0c04140c, 0x0c04141c,
|
||||
0x0c041c04, 0x0c041c14, 0x0c041c24, 0x0c04243e, 0x0c042c04, 0x0c0c0404, 0x0c0c0414, 0x0c0c0c0c,
|
||||
0x0c0c1404, 0x0c0c1414, 0x0c14040c, 0x0c14041c, 0x0c140c04, 0x0c140c14, 0x0c14140c, 0x0c141c04,
|
||||
0x0c143e14, 0x0c1c0404, 0x0c1c0414, 0x0c1c1404, 0x0c1c1c0c, 0x0c1c2434, 0x0c1c3434, 0x0c24040c,
|
||||
0x0c24042c, 0x0c242c04, 0x0c2c1404, 0x0c2c1424, 0x0c2c2434, 0x0c2c3e0c, 0x0c34042c, 0x0c3e1414,
|
||||
0x0c3e2404, 0x14040404, 0x14040414, 0x14040c0c, 0x14040c1c, 0x14041404, 0x14041414, 0x14041434,
|
||||
0x14041c0c, 0x14042414, 0x140c040c, 0x140c041c, 0x140c042c, 0x140c0c04, 0x140c0c14, 0x140c140c,
|
||||
0x140c1c04, 0x140c341c, 0x140c343e, 0x140c3e04, 0x14140404, 0x14140414, 0x14140c0c, 0x14140c3e,
|
||||
0x14141404, 0x14141414, 0x14141c3e, 0x14142404, 0x14142c2c, 0x141c040c, 0x141c0c04, 0x141c0c24,
|
||||
0x141c3e04, 0x141c3e24, 0x14241c2c, 0x14242c1c, 0x142c041c, 0x142c143e, 0x142c240c, 0x142c3e24,
|
||||
0x143e040c, 0x143e041c, 0x143e0c34, 0x143e242c, 0x1c04040c, 0x1c040c04, 0x1c040c14, 0x1c04140c,
|
||||
0x1c04141c, 0x1c042c04, 0x1c04342c, 0x1c043e14, 0x1c0c0404, 0x1c0c0414, 0x1c0c1404, 0x1c0c1c0c,
|
||||
0x1c0c2424, 0x1c0c2434, 0x1c14040c, 0x1c14041c, 0x1c140c04, 0x1c14142c, 0x1c142c14, 0x1c143e14,
|
||||
0x1c1c0c0c, 0x1c1c1c1c, 0x1c241c04, 0x1c24243e, 0x1c243e14, 0x1c2c0404, 0x1c2c0434, 0x1c2c1414,
|
||||
0x1c2c2c2c, 0x1c340c24, 0x1c341c34, 0x1c34341c, 0x1c3e1c1c, 0x1c3e3404, 0x24040424, 0x24040c3e,
|
||||
0x24041c2c, 0x24041c3e, 0x24042c1c, 0x24042c3e, 0x240c3e24, 0x24141404, 0x24141c3e, 0x24142404,
|
||||
0x24143404, 0x24143434, 0x241c043e, 0x241c242c, 0x24240424, 0x24242c0c, 0x24243424, 0x242c142c,
|
||||
0x242c241c, 0x242c3e04, 0x243e042c, 0x243e0c04, 0x243e0c14, 0x243e1c04, 0x2c040c14, 0x2c04240c,
|
||||
0x2c043e04, 0x2c0c0404, 0x2c0c0434, 0x2c0c1434, 0x2c0c2c2c, 0x2c140c24, 0x2c141c14, 0x2c143e14,
|
||||
0x2c1c0414, 0x2c1c2c1c, 0x2c240c04, 0x2c24141c, 0x2c24143e, 0x2c243e14, 0x2c2c0414, 0x2c2c1c0c,
|
||||
0x2c342c04, 0x2c3e1424, 0x2c3e2414, 0x34041424, 0x34042424, 0x34042434, 0x34043424, 0x340c140c,
|
||||
0x340c340c, 0x34140c3e, 0x34143424, 0x341c1c04, 0x341c1c34, 0x34242424, 0x342c042c, 0x342c2c14,
|
||||
0x34341c1c, 0x343e041c, 0x343e140c, 0x3e04041c, 0x3e04042c, 0x3e04043e, 0x3e040c04, 0x3e041c14,
|
||||
0x3e042c14, 0x3e0c1434, 0x3e0c2404, 0x3e140c14, 0x3e14242c, 0x3e142c14, 0x3e1c0404, 0x3e1c0c2c,
|
||||
0x3e1c1c1c, 0x3e1c3404, 0x3e24140c, 0x3e24240c, 0x3e2c0404, 0x3e2c0414, 0x3e2c1424, 0x3e341c04
|
||||
};
|
||||
|
||||
inline uint iq3_xxs_signs(uint code) {
|
||||
return code | ((popcount(code) & 1u) << 7);
|
||||
}
|
||||
@@ -0,0 +1,193 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#ifdef cl_intel_required_subgroup_size
|
||||
#define INTEL_GPU 1
|
||||
#endif
|
||||
|
||||
#define QK_K 256
|
||||
|
||||
struct block_iq3_s {
|
||||
half d;
|
||||
uchar qs[QK_K/4];
|
||||
uchar qh[QK_K/32];
|
||||
uchar signs[QK_K/8];
|
||||
uchar scales[QK_K/64];
|
||||
};
|
||||
|
||||
#define LOAD_VEC_A 4
|
||||
#define LOAD_VEC_B 4
|
||||
|
||||
#define BM 64
|
||||
#define BN 64
|
||||
#define BK 32
|
||||
#ifdef INTEL_GPU
|
||||
#define TM 8
|
||||
#define TN 8
|
||||
#else
|
||||
#define TM 4
|
||||
#define TN 8
|
||||
#endif
|
||||
|
||||
kernel void kernel_mul_mm_iq3_s_f32_l4_lm(
|
||||
global uchar * src0_qs,
|
||||
global uchar * src0_sg,
|
||||
global uchar * src0_qh,
|
||||
global uchar * src0_sc,
|
||||
global half * src0_d,
|
||||
ulong offset0,
|
||||
global float4 * src1,
|
||||
ulong offset1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
|
||||
int ne00,
|
||||
int ne01,
|
||||
int ne02,
|
||||
int ne11,
|
||||
int ne12,
|
||||
|
||||
int stride_a,
|
||||
int stride_b,
|
||||
int stride_d,
|
||||
|
||||
int batch_stride_a,
|
||||
int batch_stride_b,
|
||||
int batch_stride_d,
|
||||
|
||||
int r2,
|
||||
int r3
|
||||
) {
|
||||
src1 = (global float4*)((global char*)src1 + offset1);
|
||||
dst = (global float *)((global char*)dst + offsetd);
|
||||
|
||||
local float buf_a[BM * BK];
|
||||
local float buf_b[BN * BK];
|
||||
|
||||
const int batch_idx = get_global_id(2);
|
||||
|
||||
const int i13 = batch_idx / ne12;
|
||||
const int i12 = batch_idx % ne12;
|
||||
|
||||
const int i03 = i13 / r3;
|
||||
const int i02 = i12 / r2;
|
||||
|
||||
const int batch_idx_a = i03 * ne02 + i02;
|
||||
|
||||
const int ir = get_group_id(0);
|
||||
const int ic = get_group_id(1);
|
||||
|
||||
const int tid = get_local_id(0);
|
||||
const int th_r = tid % (BM / TM);
|
||||
const int th_c = tid / (BM / TM);
|
||||
|
||||
const int loadr_a = get_local_id(0) % (BK / LOAD_VEC_A);
|
||||
const int loadc_a = get_local_id(0) / (BK / LOAD_VEC_A);
|
||||
const int loadr_b = get_local_id(0) % (BK / LOAD_VEC_B);
|
||||
const int loadc_b = get_local_id(0) / (BK / LOAD_VEC_B);
|
||||
|
||||
const int loadstride_a = get_local_size(0) * LOAD_VEC_A / BK;
|
||||
const int loadstride_b = get_local_size(0) * LOAD_VEC_B / BK;
|
||||
|
||||
// pos_a counts elements, not blocks - one iq3_s super block holds QK_K of them
|
||||
int pos_a = batch_idx_a * batch_stride_a + ir * BM * stride_a;
|
||||
int pos_b = (batch_idx * batch_stride_b + ic * BN * stride_b) / LOAD_VEC_B;
|
||||
|
||||
float sums[TM * TN];
|
||||
float cache_a[TM];
|
||||
float cache_b[TN];
|
||||
|
||||
for (int i = 0; i < TM * TN; i++) {
|
||||
sums[i] = 0.0f;
|
||||
}
|
||||
|
||||
for (int block = 0; block < ne00; block += BK) {
|
||||
for (int l = 0; l < BM; l += loadstride_a) {
|
||||
if (ir*BM + loadc_a + l < ne01) {
|
||||
int idx = pos_a + (loadc_a + l) * stride_a + loadr_a * LOAD_VEC_A;
|
||||
int ib = idx / QK_K;
|
||||
ib += offset0/sizeof(struct block_iq3_s);
|
||||
int e = idx % QK_K;
|
||||
|
||||
int ib32 = e >> 5;
|
||||
int rem = e & 31;
|
||||
int lg = rem >> 3;
|
||||
int j0 = rem & 7;
|
||||
|
||||
uchar scb = src0_sc[4*ib + ib32/2];
|
||||
float db = (float)src0_d[ib] * (float)(1 + 2*((scb >> (4*(ib32 & 1))) & 15));
|
||||
int which = j0 >> 2;
|
||||
uint gi = (uint)src0_qs[64*ib + 8*ib32 + 2*lg + which]
|
||||
| (((uint)src0_qh[8*ib + ib32] << (8 - 2*lg - which)) & 256);
|
||||
uchar sg = src0_sg[32*ib + 4*ib32 + lg];
|
||||
uint g = iq3_s_grid[gi];
|
||||
|
||||
float4 v1;
|
||||
v1.s0 = db * (float)((g >> 0) & 0xFF) * ((sg & (1 << (j0+0))) ? -1.f : 1.f);
|
||||
v1.s1 = db * (float)((g >> 8) & 0xFF) * ((sg & (1 << (j0+1))) ? -1.f : 1.f);
|
||||
v1.s2 = db * (float)((g >> 16) & 0xFF) * ((sg & (1 << (j0+2))) ? -1.f : 1.f);
|
||||
v1.s3 = db * (float)((g >> 24) & 0xFF) * ((sg & (1 << (j0+3))) ? -1.f : 1.f);
|
||||
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = v1.s0;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = v1.s1;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = v1.s2;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = v1.s3;
|
||||
} else {
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = 0.0f;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = 0.0f;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = 0.0f;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = 0.0f;
|
||||
}
|
||||
}
|
||||
|
||||
for (int l = 0; l < BN; l += loadstride_b) {
|
||||
if (ic*BN + loadc_b + l < ne11) {
|
||||
int idx = pos_b + (loadc_b + l) * stride_b / LOAD_VEC_B + loadr_b;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = src1[idx].s0;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = src1[idx].s1;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = src1[idx].s2;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = src1[idx].s3;
|
||||
} else {
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = 0.0f;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = 0.0f;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = 0.0f;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = 0.0f;
|
||||
}
|
||||
}
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
pos_a += BK;
|
||||
pos_b += BK / LOAD_VEC_B;
|
||||
|
||||
for (int i = 0; i < BK; i++) {
|
||||
for (int j = 0; j < TM; j++) {
|
||||
cache_a[j] = buf_a[(i) * BM + th_r * TM + j];
|
||||
}
|
||||
|
||||
for (int j = 0; j < TN; j++) {
|
||||
cache_b[j] = buf_b[(i) * BN + th_c * TN + j];
|
||||
}
|
||||
|
||||
for (int cc = 0; cc < TN; cc++) {
|
||||
for (int cr = 0; cr < TM; cr++) {
|
||||
const int sums_idx = cc*TM + cr;
|
||||
sums[sums_idx] = mad(cache_a[cr], cache_b[cc], sums[sums_idx]);
|
||||
}
|
||||
}
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const int dr = ir * BM + th_r * TM;
|
||||
const int dc = ic * BN + th_c * TN;
|
||||
|
||||
const ulong offsets = batch_idx * batch_stride_d;
|
||||
|
||||
for (int cc = 0; cc < TN; cc++) {
|
||||
for (int cr = 0; cr < TM; cr++) {
|
||||
if (dr + cr < ne01 && dc + cc < ne11) {
|
||||
dst[offsets + (dc + cc) * stride_d + dr + cr] = sums[cc * TM + cr];
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,187 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#ifdef cl_intel_required_subgroup_size
|
||||
#define INTEL_GPU 1
|
||||
#endif
|
||||
|
||||
#define QK_K 256
|
||||
|
||||
struct block_iq3_xxs {
|
||||
half d;
|
||||
uchar qs[3*QK_K/8];
|
||||
};
|
||||
|
||||
#define LOAD_VEC_A 4
|
||||
#define LOAD_VEC_B 4
|
||||
|
||||
#define BM 64
|
||||
#define BN 64
|
||||
#define BK 32
|
||||
#ifdef INTEL_GPU
|
||||
#define TM 8
|
||||
#define TN 8
|
||||
#else
|
||||
#define TM 4
|
||||
#define TN 8
|
||||
#endif
|
||||
|
||||
kernel void kernel_mul_mm_iq3_xxs_f32_l4_lm(
|
||||
global uchar * src0_qs,
|
||||
global uint * src0_sas,
|
||||
global half * src0_d,
|
||||
ulong offset0,
|
||||
global float4 * src1,
|
||||
ulong offset1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
|
||||
int ne00,
|
||||
int ne01,
|
||||
int ne02,
|
||||
int ne11,
|
||||
int ne12,
|
||||
|
||||
int stride_a,
|
||||
int stride_b,
|
||||
int stride_d,
|
||||
|
||||
int batch_stride_a,
|
||||
int batch_stride_b,
|
||||
int batch_stride_d,
|
||||
|
||||
int r2,
|
||||
int r3
|
||||
) {
|
||||
src1 = (global float4*)((global char*)src1 + offset1);
|
||||
dst = (global float *)((global char*)dst + offsetd);
|
||||
|
||||
local float buf_a[BM * BK];
|
||||
local float buf_b[BN * BK];
|
||||
|
||||
const int batch_idx = get_global_id(2);
|
||||
|
||||
const int i13 = batch_idx / ne12;
|
||||
const int i12 = batch_idx % ne12;
|
||||
|
||||
const int i03 = i13 / r3;
|
||||
const int i02 = i12 / r2;
|
||||
|
||||
const int batch_idx_a = i03 * ne02 + i02;
|
||||
|
||||
const int ir = get_group_id(0);
|
||||
const int ic = get_group_id(1);
|
||||
|
||||
const int tid = get_local_id(0);
|
||||
const int th_r = tid % (BM / TM);
|
||||
const int th_c = tid / (BM / TM);
|
||||
|
||||
const int loadr_a = get_local_id(0) % (BK / LOAD_VEC_A);
|
||||
const int loadc_a = get_local_id(0) / (BK / LOAD_VEC_A);
|
||||
const int loadr_b = get_local_id(0) % (BK / LOAD_VEC_B);
|
||||
const int loadc_b = get_local_id(0) / (BK / LOAD_VEC_B);
|
||||
|
||||
const int loadstride_a = get_local_size(0) * LOAD_VEC_A / BK;
|
||||
const int loadstride_b = get_local_size(0) * LOAD_VEC_B / BK;
|
||||
|
||||
// pos_a counts elements, not blocks - one iq3_xxs super block holds QK_K of them
|
||||
int pos_a = batch_idx_a * batch_stride_a + ir * BM * stride_a;
|
||||
int pos_b = (batch_idx * batch_stride_b + ic * BN * stride_b) / LOAD_VEC_B;
|
||||
|
||||
float sums[TM * TN];
|
||||
float cache_a[TM];
|
||||
float cache_b[TN];
|
||||
|
||||
for (int i = 0; i < TM * TN; i++) {
|
||||
sums[i] = 0.0f;
|
||||
}
|
||||
|
||||
for (int block = 0; block < ne00; block += BK) {
|
||||
for (int l = 0; l < BM; l += loadstride_a) {
|
||||
if (ir*BM + loadc_a + l < ne01) {
|
||||
int idx = pos_a + (loadc_a + l) * stride_a + loadr_a * LOAD_VEC_A;
|
||||
int ib = idx / QK_K;
|
||||
ib += offset0/sizeof(struct block_iq3_xxs);
|
||||
int e = idx % QK_K;
|
||||
|
||||
int ib32 = e >> 5;
|
||||
int rem = e & 31;
|
||||
int lg = rem >> 3;
|
||||
int j0 = rem & 7;
|
||||
|
||||
uint aux = src0_sas[8*ib + ib32];
|
||||
float db = (float)src0_d[ib] * (0.5f + (float)(aux >> 28)) * 0.5f;
|
||||
|
||||
uint gi = src0_qs[64*ib + 8*ib32 + 2*lg + (j0 >> 2)];
|
||||
uint sg = iq3_xxs_signs((aux >> (7*lg)) & 127u);
|
||||
uint g = iq3_xxs_grid[gi];
|
||||
|
||||
float4 v1;
|
||||
v1.s0 = db * (float)((g >> 0) & 0xFF) * ((sg & (1 << (j0+0))) ? -1.f : 1.f);
|
||||
v1.s1 = db * (float)((g >> 8) & 0xFF) * ((sg & (1 << (j0+1))) ? -1.f : 1.f);
|
||||
v1.s2 = db * (float)((g >> 16) & 0xFF) * ((sg & (1 << (j0+2))) ? -1.f : 1.f);
|
||||
v1.s3 = db * (float)((g >> 24) & 0xFF) * ((sg & (1 << (j0+3))) ? -1.f : 1.f);
|
||||
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = v1.s0;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = v1.s1;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = v1.s2;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = v1.s3;
|
||||
} else {
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = 0.0f;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = 0.0f;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = 0.0f;
|
||||
buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = 0.0f;
|
||||
}
|
||||
}
|
||||
|
||||
for (int l = 0; l < BN; l += loadstride_b) {
|
||||
if (ic*BN + loadc_b + l < ne11) {
|
||||
int idx = pos_b + (loadc_b + l) * stride_b / LOAD_VEC_B + loadr_b;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = src1[idx].s0;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = src1[idx].s1;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = src1[idx].s2;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = src1[idx].s3;
|
||||
} else {
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = 0.0f;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = 0.0f;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = 0.0f;
|
||||
buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = 0.0f;
|
||||
}
|
||||
}
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
pos_a += BK;
|
||||
pos_b += BK / LOAD_VEC_B;
|
||||
|
||||
for (int i = 0; i < BK; i++) {
|
||||
for (int j = 0; j < TM; j++) {
|
||||
cache_a[j] = buf_a[(i) * BM + th_r * TM + j];
|
||||
}
|
||||
|
||||
for (int j = 0; j < TN; j++) {
|
||||
cache_b[j] = buf_b[(i) * BN + th_c * TN + j];
|
||||
}
|
||||
|
||||
for (int cc = 0; cc < TN; cc++) {
|
||||
for (int cr = 0; cr < TM; cr++) {
|
||||
const int sums_idx = cc*TM + cr;
|
||||
sums[sums_idx] = mad(cache_a[cr], cache_b[cc], sums[sums_idx]);
|
||||
}
|
||||
}
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const int dr = ir * BM + th_r * TM;
|
||||
const int dc = ic * BN + th_c * TN;
|
||||
|
||||
const ulong offsets = batch_idx * batch_stride_d;
|
||||
|
||||
for (int cc = 0; cc < TN; cc++) {
|
||||
for (int cr = 0; cr < TM; cr++) {
|
||||
if (dr + cr < ne01 && dc + cc < ne11) {
|
||||
dst[offsets + (dc + cc) * stride_d + dr + cr] = sums[cc * TM + cr];
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,146 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#ifdef cl_intel_required_subgroup_size
|
||||
#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
|
||||
#define INTEL_GPU 1
|
||||
#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
|
||||
#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
|
||||
#elif defined(cl_qcom_reqd_sub_group_size)
|
||||
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
|
||||
#define ADRENO_GPU 1
|
||||
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
|
||||
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
|
||||
#endif
|
||||
|
||||
#define QK_K 256
|
||||
#define IQ3_S_N_SCALE (QK_K/64)
|
||||
|
||||
typedef struct {
|
||||
half d;
|
||||
uchar qs[QK_K/4];
|
||||
uchar qh[QK_K/32];
|
||||
uchar signs[QK_K/8];
|
||||
uchar scales[IQ3_S_N_SCALE];
|
||||
} block_iq3_s;
|
||||
|
||||
#undef N_DST
|
||||
#undef N_SIMDGROUP
|
||||
#undef N_SIMDWIDTH
|
||||
|
||||
#ifdef INTEL_GPU
|
||||
#define N_DST 4
|
||||
#define N_SIMDGROUP 1
|
||||
#define N_SIMDWIDTH 16
|
||||
#elif defined (ADRENO_GPU)
|
||||
#define N_DST 4
|
||||
#define N_SIMDGROUP 1
|
||||
#define N_SIMDWIDTH 64
|
||||
#endif
|
||||
|
||||
#undef BLOCK_STRIDE
|
||||
// 8 threads cover one super block, one 32 element sub block each
|
||||
#define BLOCK_STRIDE (N_SIMDWIDTH/8)
|
||||
|
||||
#ifdef INTEL_GPU
|
||||
REQD_SUBGROUP_SIZE_16
|
||||
#elif defined (ADRENO_GPU)
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_mul_mv_iq3_s_f32(
|
||||
global char * src0,
|
||||
ulong offset0,
|
||||
global char * src1,
|
||||
ulong offset1,
|
||||
global char * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
ulong nb01,
|
||||
ulong nb02,
|
||||
ulong nb03,
|
||||
int ne12,
|
||||
ulong nb11,
|
||||
ulong nb12,
|
||||
ulong nb13,
|
||||
int ne0,
|
||||
int ne1,
|
||||
int r2,
|
||||
int r3
|
||||
) {
|
||||
src0 = src0 + offset0;
|
||||
src1 = src1 + offset1;
|
||||
dst = dst + offsetd;
|
||||
|
||||
int ix = get_sub_group_local_id()/8; // super block index
|
||||
int it = get_sub_group_local_id()%8; // sub block inside the super block
|
||||
|
||||
int nb = ne00/QK_K;
|
||||
|
||||
int r0 = get_group_id(0);
|
||||
int r1 = get_group_id(1);
|
||||
int im = get_group_id(2);
|
||||
int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST;
|
||||
|
||||
int i12 = im%ne12;
|
||||
int i13 = im/ne12;
|
||||
|
||||
ulong offset_src0 = first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03;
|
||||
ulong offset_src1 = r1*nb11 + (i12 )*nb12 + (i13 )*nb13;
|
||||
|
||||
global block_iq3_s * x = (global block_iq3_s *) (src0 + offset_src0);
|
||||
global float * y = (global float *) (src1 + offset_src1);
|
||||
|
||||
float yl[32];
|
||||
float sumf[N_DST] = {0.f};
|
||||
float all_sum;
|
||||
|
||||
global float * y4 = y + ix * QK_K + 32 * it;
|
||||
|
||||
for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) {
|
||||
for (int i = 0; i < 32; ++i) {
|
||||
yl[i] = y4[i];
|
||||
}
|
||||
|
||||
global char * xrow = (global char *)(x + ib);
|
||||
|
||||
// only ne01 output rows exist; reading past them can pull in an inf scale
|
||||
for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
|
||||
global block_iq3_s * xb = (global block_iq3_s *)(xrow + row*nb01);
|
||||
|
||||
uchar nib = (it & 1) ? (xb->scales[it >> 1] >> 4) : (xb->scales[it >> 1] & 0xf);
|
||||
float db = (float)xb->d * (float)(1 + 2*nib);
|
||||
|
||||
uchar qhb = xb->qh[it];
|
||||
global uchar * qsb = xb->qs + 8*it;
|
||||
global uchar * sgb = xb->signs + 4*it;
|
||||
|
||||
float acc = 0.f;
|
||||
for (int l = 0; l < 4; ++l) {
|
||||
uint i1 = (uint)qsb[2*l+0] | (((uint)qhb << (8-2*l)) & 256);
|
||||
uint i2 = (uint)qsb[2*l+1] | (((uint)qhb << (7-2*l)) & 256);
|
||||
uint g1 = iq3_s_grid[i1];
|
||||
uint g2 = iq3_s_grid[i2];
|
||||
uchar sg = sgb[l];
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
acc += yl[8*l+j+0] * (float)((g1 >> (8*j)) & 0xFF) * ((sg & (1 << j)) ? -1.f : 1.f);
|
||||
acc += yl[8*l+j+4] * (float)((g2 >> (8*j)) & 0xFF) * ((sg & (1 << (j+4))) ? -1.f : 1.f);
|
||||
}
|
||||
}
|
||||
|
||||
sumf[row] += db * acc;
|
||||
}
|
||||
|
||||
y4 += BLOCK_STRIDE * QK_K;
|
||||
}
|
||||
|
||||
global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0;
|
||||
|
||||
for (int row = 0; row < N_DST; ++row) {
|
||||
all_sum = sub_group_reduce_add(sumf[row]);
|
||||
if (first_row + row < ne01) {
|
||||
if (get_sub_group_local_id() == 0) {
|
||||
dst_f32[first_row + row] = all_sum;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,147 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#ifdef cl_intel_required_subgroup_size
|
||||
#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
|
||||
#define INTEL_GPU 1
|
||||
#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
|
||||
#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
|
||||
#elif defined(cl_qcom_reqd_sub_group_size)
|
||||
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
|
||||
#define ADRENO_GPU 1
|
||||
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
|
||||
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
|
||||
#endif
|
||||
|
||||
#define QK_K 256
|
||||
#define IQ3_S_N_SCALE (QK_K/64)
|
||||
|
||||
typedef struct {
|
||||
half d;
|
||||
uchar qs[QK_K/4];
|
||||
uchar qh[QK_K/32];
|
||||
uchar signs[QK_K/8];
|
||||
uchar scales[IQ3_S_N_SCALE];
|
||||
} block_iq3_s;
|
||||
|
||||
#undef N_DST
|
||||
#undef N_SIMDGROUP
|
||||
#undef N_SIMDWIDTH
|
||||
|
||||
#ifdef INTEL_GPU
|
||||
#define N_DST 4
|
||||
#define N_SIMDGROUP 1
|
||||
#define N_SIMDWIDTH 16
|
||||
#elif defined (ADRENO_GPU)
|
||||
#define N_DST 4
|
||||
#define N_SIMDGROUP 1
|
||||
#define N_SIMDWIDTH 64
|
||||
#endif
|
||||
|
||||
#undef BLOCK_STRIDE
|
||||
// 8 threads cover one super block, one 32 element sub block each
|
||||
#define BLOCK_STRIDE (N_SIMDWIDTH/8)
|
||||
|
||||
#ifdef INTEL_GPU
|
||||
REQD_SUBGROUP_SIZE_16
|
||||
#elif defined (ADRENO_GPU)
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_mul_mv_iq3_s_f32_flat(
|
||||
global uchar * src0_qs,
|
||||
global uchar * src0_sg,
|
||||
global uchar * src0_qh,
|
||||
global uchar * src0_sc,
|
||||
global half * src0_d,
|
||||
ulong offset0,
|
||||
global char * src1,
|
||||
ulong offset1,
|
||||
global char * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
ulong nb01,
|
||||
ulong nb02,
|
||||
ulong nb03,
|
||||
int ne12,
|
||||
ulong nb11,
|
||||
ulong nb12,
|
||||
ulong nb13,
|
||||
int ne0,
|
||||
int ne1,
|
||||
int r2,
|
||||
int r3
|
||||
) {
|
||||
src1 = src1 + offset1;
|
||||
dst = dst + offsetd;
|
||||
|
||||
int ix = get_sub_group_local_id()/8; // super block index
|
||||
int it = get_sub_group_local_id()%8; // sub block inside the super block
|
||||
|
||||
int nb = ne00/QK_K;
|
||||
|
||||
int r0 = get_group_id(0);
|
||||
int r1 = get_group_id(1);
|
||||
int im = get_group_id(2);
|
||||
int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST;
|
||||
|
||||
int i12 = im%ne12;
|
||||
int i13 = im/ne12;
|
||||
|
||||
ulong offset_src0 = (offset0 + first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03)/sizeof(block_iq3_s);
|
||||
ulong row_stride = nb01/sizeof(block_iq3_s);
|
||||
ulong offset_src1 = r1*nb11 + (i12 )*nb12 + (i13 )*nb13;
|
||||
|
||||
global float * y = (global float *) (src1 + offset_src1);
|
||||
|
||||
float yl[32];
|
||||
float sumf[N_DST] = {0.f};
|
||||
float all_sum;
|
||||
|
||||
global float * y4 = y + ix * QK_K + 32 * it;
|
||||
|
||||
for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) {
|
||||
for (int i = 0; i < 32; ++i) {
|
||||
yl[i] = y4[i];
|
||||
}
|
||||
|
||||
// only ne01 output rows exist; reading past them can pull in an inf scale
|
||||
for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
|
||||
const ulong b = offset_src0 + ib + row*row_stride;
|
||||
|
||||
uchar nib = (it & 1) ? (src0_sc[4*b + (it >> 1)] >> 4) : (src0_sc[4*b + (it >> 1)] & 0xf);
|
||||
float db = (float)src0_d[b] * (float)(1 + 2*nib);
|
||||
|
||||
uchar qhb = src0_qh[8*b + it];
|
||||
global uchar * qsb = src0_qs + 64*b + 8*it;
|
||||
global uchar * sgb = src0_sg + 32*b + 4*it;
|
||||
|
||||
float acc = 0.f;
|
||||
for (int l = 0; l < 4; ++l) {
|
||||
uint i1 = (uint)qsb[2*l+0] | (((uint)qhb << (8-2*l)) & 256);
|
||||
uint i2 = (uint)qsb[2*l+1] | (((uint)qhb << (7-2*l)) & 256);
|
||||
uint g1 = iq3_s_grid[i1];
|
||||
uint g2 = iq3_s_grid[i2];
|
||||
uchar sg = sgb[l];
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
acc += yl[8*l+j+0] * (float)((g1 >> (8*j)) & 0xFF) * ((sg & (1 << j)) ? -1.f : 1.f);
|
||||
acc += yl[8*l+j+4] * (float)((g2 >> (8*j)) & 0xFF) * ((sg & (1 << (j+4))) ? -1.f : 1.f);
|
||||
}
|
||||
}
|
||||
|
||||
sumf[row] += db * acc;
|
||||
}
|
||||
|
||||
y4 += BLOCK_STRIDE * QK_K;
|
||||
}
|
||||
|
||||
global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0;
|
||||
|
||||
for (int row = 0; row < N_DST; ++row) {
|
||||
all_sum = sub_group_reduce_add(sumf[row]);
|
||||
if (first_row + row < ne01) {
|
||||
if (get_sub_group_local_id() == 0) {
|
||||
dst_f32[first_row + row] = all_sum;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,145 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#ifdef cl_intel_required_subgroup_size
|
||||
#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
|
||||
#define INTEL_GPU 1
|
||||
#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
|
||||
#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
|
||||
#elif defined(cl_qcom_reqd_sub_group_size)
|
||||
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
|
||||
#define ADRENO_GPU 1
|
||||
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
|
||||
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
|
||||
#endif
|
||||
|
||||
#define QK_K 256
|
||||
|
||||
// d, then 64 grid index bytes, then 8 uint32 of packed scale+signs
|
||||
typedef struct {
|
||||
half d;
|
||||
uchar qs[3*QK_K/8];
|
||||
} block_iq3_xxs;
|
||||
|
||||
#undef N_DST
|
||||
#undef N_SIMDGROUP
|
||||
#undef N_SIMDWIDTH
|
||||
|
||||
#ifdef INTEL_GPU
|
||||
#define N_DST 4
|
||||
#define N_SIMDGROUP 1
|
||||
#define N_SIMDWIDTH 16
|
||||
#elif defined (ADRENO_GPU)
|
||||
#define N_DST 4
|
||||
#define N_SIMDGROUP 1
|
||||
#define N_SIMDWIDTH 64
|
||||
#endif
|
||||
|
||||
#undef BLOCK_STRIDE
|
||||
// 8 threads cover one super block, one 32 element sub block each
|
||||
#define BLOCK_STRIDE (N_SIMDWIDTH/8)
|
||||
|
||||
#ifdef INTEL_GPU
|
||||
REQD_SUBGROUP_SIZE_16
|
||||
#elif defined (ADRENO_GPU)
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_mul_mv_iq3_xxs_f32(
|
||||
global char * src0,
|
||||
ulong offset0,
|
||||
global char * src1,
|
||||
ulong offset1,
|
||||
global char * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
ulong nb01,
|
||||
ulong nb02,
|
||||
ulong nb03,
|
||||
int ne12,
|
||||
ulong nb11,
|
||||
ulong nb12,
|
||||
ulong nb13,
|
||||
int ne0,
|
||||
int ne1,
|
||||
int r2,
|
||||
int r3
|
||||
) {
|
||||
src0 = src0 + offset0;
|
||||
src1 = src1 + offset1;
|
||||
dst = dst + offsetd;
|
||||
|
||||
int ix = get_sub_group_local_id()/8; // super block index
|
||||
int it = get_sub_group_local_id()%8; // sub block inside the super block
|
||||
|
||||
int nb = ne00/QK_K;
|
||||
|
||||
int r0 = get_group_id(0);
|
||||
int r1 = get_group_id(1);
|
||||
int im = get_group_id(2);
|
||||
int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST;
|
||||
|
||||
int i12 = im%ne12;
|
||||
int i13 = im/ne12;
|
||||
|
||||
ulong offset_src0 = first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03;
|
||||
ulong offset_src1 = r1*nb11 + (i12 )*nb12 + (i13 )*nb13;
|
||||
|
||||
global block_iq3_xxs * x = (global block_iq3_xxs *) (src0 + offset_src0);
|
||||
global float * y = (global float *) (src1 + offset_src1);
|
||||
|
||||
float yl[32];
|
||||
float sumf[N_DST] = {0.f};
|
||||
float all_sum;
|
||||
|
||||
global float * y4 = y + ix * QK_K + 32 * it;
|
||||
|
||||
for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) {
|
||||
for (int i = 0; i < 32; ++i) {
|
||||
yl[i] = y4[i];
|
||||
}
|
||||
|
||||
global char * xrow = (global char *)(x + ib);
|
||||
|
||||
// only ne01 output rows exist; reading past them can pull in an inf scale
|
||||
for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
|
||||
global block_iq3_xxs * xb = (global block_iq3_xxs *)(xrow + row*nb01);
|
||||
|
||||
global uchar * qs = xb->qs + 8*it;
|
||||
|
||||
// Blocks are 98 bytes, so packed signs are only 2-byte aligned.
|
||||
global ushort * sw = (global ushort *)(xb->qs + 64 + 4*it);
|
||||
uint aux32 = (uint)sw[0] | ((uint)sw[1] << 16);
|
||||
|
||||
float db = (float)xb->d * (0.5f + (float)(aux32 >> 28)) * 0.5f;
|
||||
|
||||
float acc = 0.f;
|
||||
for (int l = 0; l < 4; ++l) {
|
||||
uchar signs = iq3_xxs_signs((aux32 >> (7*l)) & 127u);
|
||||
uint g1 = iq3_xxs_grid[qs[2*l+0]];
|
||||
uint g2 = iq3_xxs_grid[qs[2*l+1]];
|
||||
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
float s1 = (signs & (1 << j)) ? -1.f : 1.f;
|
||||
float s2 = (signs & (1 << (j+4))) ? -1.f : 1.f;
|
||||
acc += yl[8*l+j+0] * (float)((g1 >> (8*j)) & 0xFF) * s1;
|
||||
acc += yl[8*l+j+4] * (float)((g2 >> (8*j)) & 0xFF) * s2;
|
||||
}
|
||||
}
|
||||
|
||||
sumf[row] += db * acc;
|
||||
}
|
||||
|
||||
y4 += BLOCK_STRIDE * QK_K;
|
||||
}
|
||||
|
||||
global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0;
|
||||
|
||||
for (int row = 0; row < N_DST; ++row) {
|
||||
all_sum = sub_group_reduce_add(sumf[row]);
|
||||
if (first_row + row < ne01) {
|
||||
if (get_sub_group_local_id() == 0) {
|
||||
dst_f32[first_row + row] = all_sum;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,142 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#ifdef cl_intel_required_subgroup_size
|
||||
#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
|
||||
#define INTEL_GPU 1
|
||||
#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
|
||||
#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
|
||||
#elif defined(cl_qcom_reqd_sub_group_size)
|
||||
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
|
||||
#define ADRENO_GPU 1
|
||||
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
|
||||
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
|
||||
#endif
|
||||
|
||||
#define QK_K 256
|
||||
|
||||
// d, then 64 grid index bytes, then 8 uint32 of packed scale+signs
|
||||
typedef struct {
|
||||
half d;
|
||||
uchar qs[3*QK_K/8];
|
||||
} block_iq3_xxs;
|
||||
|
||||
#undef N_DST
|
||||
#undef N_SIMDGROUP
|
||||
#undef N_SIMDWIDTH
|
||||
|
||||
#ifdef INTEL_GPU
|
||||
#define N_DST 4
|
||||
#define N_SIMDGROUP 1
|
||||
#define N_SIMDWIDTH 16
|
||||
#elif defined (ADRENO_GPU)
|
||||
#define N_DST 4
|
||||
#define N_SIMDGROUP 1
|
||||
#define N_SIMDWIDTH 64
|
||||
#endif
|
||||
|
||||
#undef BLOCK_STRIDE
|
||||
// 8 threads cover one super block, one 32 element sub block each
|
||||
#define BLOCK_STRIDE (N_SIMDWIDTH/8)
|
||||
|
||||
#ifdef INTEL_GPU
|
||||
REQD_SUBGROUP_SIZE_16
|
||||
#elif defined (ADRENO_GPU)
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_mul_mv_iq3_xxs_f32_flat(
|
||||
global uchar * src0_qs,
|
||||
global uint * src0_sas,
|
||||
global half * src0_d,
|
||||
ulong offset0,
|
||||
global char * src1,
|
||||
ulong offset1,
|
||||
global char * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
ulong nb01,
|
||||
ulong nb02,
|
||||
ulong nb03,
|
||||
int ne12,
|
||||
ulong nb11,
|
||||
ulong nb12,
|
||||
ulong nb13,
|
||||
int ne0,
|
||||
int ne1,
|
||||
int r2,
|
||||
int r3
|
||||
) {
|
||||
src1 = src1 + offset1;
|
||||
dst = dst + offsetd;
|
||||
|
||||
int ix = get_sub_group_local_id()/8; // super block index
|
||||
int it = get_sub_group_local_id()%8; // sub block inside the super block
|
||||
|
||||
int nb = ne00/QK_K;
|
||||
|
||||
int r0 = get_group_id(0);
|
||||
int r1 = get_group_id(1);
|
||||
int im = get_group_id(2);
|
||||
int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST;
|
||||
|
||||
int i12 = im%ne12;
|
||||
int i13 = im/ne12;
|
||||
|
||||
ulong offset_src0 = (offset0 + first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03)/sizeof(block_iq3_xxs);
|
||||
ulong row_stride = nb01/sizeof(block_iq3_xxs);
|
||||
ulong offset_src1 = r1*nb11 + (i12 )*nb12 + (i13 )*nb13;
|
||||
|
||||
global float * y = (global float *) (src1 + offset_src1);
|
||||
|
||||
float yl[32];
|
||||
float sumf[N_DST] = {0.f};
|
||||
float all_sum;
|
||||
|
||||
global float * y4 = y + ix * QK_K + 32 * it;
|
||||
|
||||
for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) {
|
||||
for (int i = 0; i < 32; ++i) {
|
||||
yl[i] = y4[i];
|
||||
}
|
||||
|
||||
// only ne01 output rows exist; reading past them can pull in an inf scale
|
||||
for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
|
||||
const ulong b = offset_src0 + ib + row*row_stride;
|
||||
|
||||
global uchar * qs = src0_qs + 64*b + 8*it;
|
||||
|
||||
uint aux32 = src0_sas[8*b + it];
|
||||
|
||||
float db = (float)src0_d[b] * (0.5f + (float)(aux32 >> 28)) * 0.5f;
|
||||
|
||||
float acc = 0.f;
|
||||
for (int l = 0; l < 4; ++l) {
|
||||
uchar signs = iq3_xxs_signs((aux32 >> (7*l)) & 127u);
|
||||
uint g1 = iq3_xxs_grid[qs[2*l+0]];
|
||||
uint g2 = iq3_xxs_grid[qs[2*l+1]];
|
||||
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
float s1 = (signs & (1 << j)) ? -1.f : 1.f;
|
||||
float s2 = (signs & (1 << (j+4))) ? -1.f : 1.f;
|
||||
acc += yl[8*l+j+0] * (float)((g1 >> (8*j)) & 0xFF) * s1;
|
||||
acc += yl[8*l+j+4] * (float)((g2 >> (8*j)) & 0xFF) * s2;
|
||||
}
|
||||
}
|
||||
|
||||
sumf[row] += db * acc;
|
||||
}
|
||||
|
||||
y4 += BLOCK_STRIDE * QK_K;
|
||||
}
|
||||
|
||||
global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0;
|
||||
|
||||
for (int row = 0; row < N_DST; ++row) {
|
||||
all_sum = sub_group_reduce_add(sumf[row]);
|
||||
if (first_row + row < ne01) {
|
||||
if (get_sub_group_local_id() == 0) {
|
||||
dst_f32[first_row + row] = all_sum;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
Reference in new issue
Block a user