Skip to content

Commit 59b33b9

Browse files
ggerganovggerganov
committed
metal : some mul_mv experiments
1 parent ab6a3b7 commit 59b33b9

File tree

3 files changed

+206
-1
lines changed

3 files changed

+206
-1
lines changed

‎ggml/src/ggml-metal-impl.h‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -192,6 +192,29 @@ typedef struct {
192192
int16_t r3;
193193
} ggml_metal_kargs_mul_mv;
194194

195+
typedef struct {
196+
int32_t ne00;
197+
int32_t ne01;
198+
int32_t ne02;
199+
uint64_t nb00;
200+
uint64_t nb01;
201+
uint64_t nb02;
202+
uint64_t nb03;
203+
int32_t ne10;
204+
int32_t ne11;
205+
int32_t ne12;
206+
uint64_t nb10;
207+
uint64_t nb11;
208+
uint64_t nb12;
209+
uint64_t nb13;
210+
int32_t ne0;
211+
int32_t ne1;
212+
int16_t r2;
213+
int16_t r3;
214+
int16_t nsg;
215+
int16_t nxpsg;
216+
} ggml_metal_kargs_mul_mv_ext;
217+
195218
typedef struct {
196219
int32_t nei0;
197220
int32_t nei1;

‎ggml/src/ggml-metal.m‎

Lines changed: 54 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,8 @@
11
#import "ggml-metal.h"
22

33
#import "ggml-impl.h"
4+
#define GGML_COMMON_DECL_C
5+
#import "ggml-common.h"
46
#import "ggml-backend-impl.h"
57
#import "ggml-metal-impl.h"
68

@@ -174,6 +176,7 @@ static void ggml_backend_metal_device_rel(struct ggml_backend_metal_device_conte
174176
GGML_METAL_KERNEL_TYPE_MUL_MV_Q5_0_F32,
175177
GGML_METAL_KERNEL_TYPE_MUL_MV_Q5_1_F32,
176178
GGML_METAL_KERNEL_TYPE_MUL_MV_Q8_0_F32,
179+
GGML_METAL_KERNEL_TYPE_MUL_MV_EXT_Q8_0_F32,
177180
GGML_METAL_KERNEL_TYPE_MUL_MV_Q2_K_F32,
178181
GGML_METAL_KERNEL_TYPE_MUL_MV_Q3_K_F32,
179182
GGML_METAL_KERNEL_TYPE_MUL_MV_Q4_K_F32,
@@ -693,6 +696,7 @@ @implementation GGMLMetalClass
693696
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MV_Q5_0_F32, mul_mv_q5_0_f32, has_simdgroup_reduction);
694697
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MV_Q5_1_F32, mul_mv_q5_1_f32, has_simdgroup_reduction);
695698
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MV_Q8_0_F32, mul_mv_q8_0_f32, has_simdgroup_reduction);
699+
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MV_EXT_Q8_0_F32, mul_mv_ext_q8_0_f32, has_simdgroup_reduction);
696700
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MV_Q2_K_F32, mul_mv_q2_K_f32, has_simdgroup_reduction);
697701
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MV_Q3_K_F32, mul_mv_q3_K_f32, has_simdgroup_reduction);
698702
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MV_Q4_K_F32, mul_mv_q4_K_f32, has_simdgroup_reduction);
@@ -1908,7 +1912,7 @@ static void ggml_metal_encode_node(
19081912

19091913
// find the break-even point where the matrix-matrix kernel becomes more efficient compared
19101914
// to the matrix-vector kernel
1911-
int ne11_mm_min = 1;
1915+
int ne11_mm_min = 16;
19121916

19131917
#if 0
19141918
// the numbers below are measured on M2 Ultra for 7B and 13B models
@@ -1932,6 +1936,55 @@ static void ggml_metal_encode_node(
19321936
}
19331937
#endif
19341938

1939+
if (src0t == GGML_TYPE_Q8_0 && (ne00%16 == 0)) {
1940+
//if (false) {
1941+
id<MTLComputePipelineState> pipeline = nil;
1942+
1943+
pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_MUL_MV_EXT_Q8_0_F32].pipeline;
1944+
1945+
const int nsg = 2;
1946+
const int r0pt = 2;
1947+
const int r1pt = 1;
1948+
const int nxpsg = ne11 > 1 ? 8 : 32;
1949+
const int nypsg = 32/nxpsg;
1950+
const int nr0ptg = nypsg*r0pt*nsg;
1951+
1952+
//GGML_ASSERT(ne00%1024 == 0);
1953+
//GGML_ASSERT(ne01%nr0ptg == 0);
1954+
//printf("ne01 = %lld, nr0ptg = %d, ne00 = %lld\n", ne01, nr0ptg, ne00);
1955+
1956+
ggml_metal_kargs_mul_mv_ext args = {
1957+
/*.ne00 =*/ ne00,
1958+
/*.ne01 =*/ ne01,
1959+
/*.ne02 =*/ ne02,
1960+
/*.nb00 =*/ nb00,
1961+
/*.nb01 =*/ nb01,
1962+
/*.nb02 =*/ nb02,
1963+
/*.nb03 =*/ nb03,
1964+
/*.ne10 =*/ ne10,
1965+
/*.ne11 =*/ ne11,
1966+
/*.ne12 =*/ ne12,
1967+
/*.nb10 =*/ nb10,
1968+
/*.nb11 =*/ nb11,
1969+
/*.nb12 =*/ nb12,
1970+
/*.nb13 =*/ nb13,
1971+
/*.ne0 =*/ ne0,
1972+
/*.ne1 =*/ ne1,
1973+
/*.r2 =*/ r2,
1974+
/*.r3 =*/ r3,
1975+
/*.nsg =*/ nsg,
1976+
/*.nxpsg =*/ nxpsg,
1977+
};
1978+
1979+
[encoder setComputePipelineState:pipeline];
1980+
[encoder setBytes:&args length:sizeof(args) atIndex:0];
1981+
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:1];
1982+
[encoder setBuffer:id_src1 offset:offs_src1 atIndex:2];
1983+
[encoder setBuffer:id_dst offset:offs_dst atIndex:3];
1984+
1985+
//printf("ne01 = %lld nr0ptg = %d\n", ne01, nr0ptg);
1986+
[encoder dispatchThreadgroups:MTLSizeMake((ne01 + nr0ptg - 1)/nr0ptg, (ne11 + r1pt - 1)/r1pt, ne12*ne13) threadsPerThreadgroup:MTLSizeMake(32, nsg, 1)];
1987+
} else
19351988
// for now the matrix-matrix multiplication kernel only works on A14+/M1+ SoCs
19361989
// AMD GPU and older A-chips will reuse matrix-vector multiplication kernel
19371990
if ([device supportsFamily:MTLGPUFamilyApple7] &&

‎ggml/src/ggml-metal.metal‎

Lines changed: 129 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1739,6 +1739,135 @@ kernel void kernel_mul_mv_q8_0_f32(
17391739
kernel_mul_mv_q8_0_f32_impl<constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg);
17401740
}
17411741

1742+
template<short nsg, short nxpsg>
1743+
void kernel_mul_mv_ext_q8_0_f32_impl(
1744+
constant ggml_metal_kargs_mul_mv_ext & args,
1745+
device const char * src0,
1746+
device const char * src1,
1747+
device char * dst,
1748+
uint3 tgpig[[threadgroup_position_in_grid]],
1749+
ushort3 ntg[[threads_per_threadgroup]],
1750+
ushort tiisg[[thread_index_in_simdgroup]],
1751+
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
1752+
const short chpt = 1;
1753+
const short r0pt = 2;
1754+
1755+
//const short nxpsg = (32);
1756+
const short nypsg = (32/nxpsg)*r0pt;
1757+
1758+
const short tx = tiisg%nxpsg;
1759+
const short ty = tiisg/nxpsg;
1760+
1761+
const int i01 = tgpig.x*(nypsg*nsg) + nypsg*sgitg + ty*r0pt;
1762+
const int i11 = tgpig.y;
1763+
const int i1m = tgpig.z;
1764+
1765+
const int i12 = i1m%args.ne12;
1766+
const int i13 = i1m/args.ne12;
1767+
1768+
const uint64_t offset0 = i01*args.nb01 + (i12/args.r2)*args.nb02 + (i13/args.r3)*args.nb03;
1769+
const uint64_t offset1 = i11*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
1770+
1771+
device const block_q8_0 * xq[r0pt];
1772+
1773+
for (short ir0 = 0; ir0 < r0pt; ++ir0) {
1774+
xq[ir0] = (i01 + ir0 < args.ne01) ? (device const block_q8_0 *) (src0 + offset0 + ir0*args.nb01) + (chpt*tx)/2 : (device const block_q8_0 *) src0;
1775+
}
1776+
1777+
device const float4x4 * y4x4 = (device const float4x4 *) (src1 + offset1) + chpt*tx;
1778+
1779+
float sumf[r0pt] = { [0 ... r0pt - 1] = 0.0f };
1780+
1781+
for (int iib = 0; (16*chpt)*(iib*nxpsg + tx) < args.ne00; ++iib) {
1782+
float4x4 lx;
1783+
1784+
#pragma unroll(2)
1785+
for (short ir0 = 0; ir0 < r0pt; ++ir0) {
1786+
#pragma unroll
1787+
for (short ch = 0; ch < chpt; ++ch) {
1788+
dequantize_q8_0(xq[ir0] + ch/2, (chpt*tx + ch)%2, lx);
1789+
1790+
const float4x4 ly = y4x4[ch];
1791+
1792+
sumf[ir0] +=
1793+
dot(lx[0], ly[0]) +
1794+
dot(lx[1], ly[1]) +
1795+
dot(lx[2], ly[2]) +
1796+
dot(lx[3], ly[3]);
1797+
}
1798+
}
1799+
1800+
y4x4 += ((16*chpt)*nxpsg)/16;
1801+
1802+
for (short ir0 = 0; ir0 < r0pt; ++ir0) {
1803+
xq[ir0] += ((16*chpt)*nxpsg)/32;
1804+
}
1805+
}
1806+
1807+
for (short ir0 = 0; ir0 < r0pt; ++ir0) {
1808+
if (nxpsg >= 32) {
1809+
sumf[ir0] += simd_shuffle_down(sumf[ir0], 16);
1810+
}
1811+
if (nxpsg >= 16) {
1812+
sumf[ir0] += simd_shuffle_down(sumf[ir0], 8);
1813+
}
1814+
if (nxpsg >= 8) {
1815+
sumf[ir0] += simd_shuffle_down(sumf[ir0], 4);
1816+
}
1817+
if (nxpsg >= 4) {
1818+
sumf[ir0] += simd_shuffle_down(sumf[ir0], 2);
1819+
}
1820+
if (nxpsg >= 2) {
1821+
sumf[ir0] += simd_shuffle_down(sumf[ir0], 1);
1822+
}
1823+
1824+
//sumf[ir0] = simd_sum(sumf[ir0]);
1825+
}
1826+
1827+
device float * dst_f32 = (device float *) dst + (uint64_t)i1m*args.ne0*args.ne1 + (uint64_t)i11*args.ne0;
1828+
1829+
if (tx == 0) {
1830+
for (short ir0 = 0; ir0 < r0pt && i01 + ir0 < args.ne01; ++ir0) {
1831+
dst_f32[i01 + ir0] = sumf[ir0];
1832+
}
1833+
}
1834+
}
1835+
1836+
[[host_name("kernel_mul_mv_ext_q8_0_f32")]]
1837+
kernel void kernel_mul_mv_ext_q8_0_f32(
1838+
constant ggml_metal_kargs_mul_mv_ext & args,
1839+
device const char * src0,
1840+
device const char * src1,
1841+
device char * dst,
1842+
uint3 tgpig[[threadgroup_position_in_grid]],
1843+
ushort3 ntg[[threads_per_threadgroup]],
1844+
ushort tiisg[[thread_index_in_simdgroup]],
1845+
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
1846+
switch (args.nsg) {
1847+
case 1:
1848+
switch (args.nxpsg) {
1849+
case 4: kernel_mul_mv_ext_q8_0_f32_impl<1, 4> (args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1850+
case 8: kernel_mul_mv_ext_q8_0_f32_impl<1, 8> (args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1851+
case 16: kernel_mul_mv_ext_q8_0_f32_impl<1, 16>(args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1852+
case 32: kernel_mul_mv_ext_q8_0_f32_impl<1, 32>(args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1853+
} break;
1854+
case 2:
1855+
switch (args.nxpsg) {
1856+
case 4: kernel_mul_mv_ext_q8_0_f32_impl<2, 4> (args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1857+
case 8: kernel_mul_mv_ext_q8_0_f32_impl<2, 8> (args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1858+
case 16: kernel_mul_mv_ext_q8_0_f32_impl<2, 16>(args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1859+
case 32: kernel_mul_mv_ext_q8_0_f32_impl<2, 32>(args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1860+
} break;
1861+
case 4:
1862+
switch (args.nxpsg) {
1863+
case 4: kernel_mul_mv_ext_q8_0_f32_impl<4, 4> (args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1864+
case 8: kernel_mul_mv_ext_q8_0_f32_impl<4, 8> (args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1865+
case 16: kernel_mul_mv_ext_q8_0_f32_impl<4, 16>(args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1866+
case 32: kernel_mul_mv_ext_q8_0_f32_impl<4, 32>(args, src0, src1, dst, tgpig, ntg, tiisg, sgitg); break;
1867+
} break;
1868+
}
1869+
}
1870+
17421871
#define N_MV_T_T 4
17431872

17441873
template<typename T0, typename T04, typename T1, typename T14, typename args_t>

0 commit comments

Comments
 (0)