2023-06-04 22:34:30 +02:00
|
|
|
#import "ggml-metal.h"
|
|
|
|
|
|
|
|
#import "ggml.h"
|
|
|
|
|
|
|
|
#import <Foundation/Foundation.h>
|
|
|
|
|
|
|
|
#import <Metal/Metal.h>
|
|
|
|
|
2023-08-07 09:52:57 +02:00
|
|
|
#undef MIN
|
|
|
|
#undef MAX
|
|
|
|
#define MIN(a, b) ((a) < (b) ? (a) : (b))
|
|
|
|
#define MAX(a, b) ((a) > (b) ? (a) : (b))
|
|
|
|
|
2023-06-04 22:34:30 +02:00
|
|
|
#ifdef GGML_METAL_NDEBUG
|
|
|
|
#define metal_printf(...)
|
|
|
|
#else
|
|
|
|
#define metal_printf(...) fprintf(stderr, __VA_ARGS__)
|
|
|
|
#endif
|
|
|
|
|
|
|
|
#define UNUSED(x) (void)(x)
|
|
|
|
|
2023-08-07 09:52:57 +02:00
|
|
|
#define GGML_MAX_CONCUR (2*GGML_MAX_NODES)
|
|
|
|
|
2023-06-04 22:34:30 +02:00
|
|
|
struct ggml_metal_buffer {
|
|
|
|
const char * name;
|
|
|
|
|
|
|
|
void * data;
|
|
|
|
size_t size;
|
|
|
|
|
|
|
|
id<MTLBuffer> metal;
|
|
|
|
};
|
|
|
|
|
|
|
|
struct ggml_metal_context {
|
2023-07-07 18:24:01 +02:00
|
|
|
int n_cb;
|
|
|
|
|
2023-06-04 22:34:30 +02:00
|
|
|
float * logits;
|
|
|
|
|
|
|
|
id<MTLDevice> device;
|
|
|
|
id<MTLCommandQueue> queue;
|
|
|
|
id<MTLLibrary> library;
|
|
|
|
|
|
|
|
int n_buffers;
|
|
|
|
struct ggml_metal_buffer buffers[GGML_METAL_MAX_BUFFERS];
|
|
|
|
|
2023-08-07 09:52:57 +02:00
|
|
|
int concur_list[GGML_MAX_CONCUR];
|
2023-07-25 14:00:19 +02:00
|
|
|
int concur_list_len;
|
|
|
|
|
2023-06-04 22:34:30 +02:00
|
|
|
// custom kernels
|
|
|
|
#define GGML_METAL_DECL_KERNEL(name) \
|
|
|
|
id<MTLFunction> function_##name; \
|
|
|
|
id<MTLComputePipelineState> pipeline_##name
|
|
|
|
|
|
|
|
GGML_METAL_DECL_KERNEL(add);
|
2023-07-23 13:00:37 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(add_row); // TODO: avoid this extra kernel, instead extend the "add" kernel to support broadcast
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(mul);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_row); // TODO: avoid this extra kernel, instead extend the "mul" kernel to support broadcast
|
|
|
|
GGML_METAL_DECL_KERNEL(scale);
|
|
|
|
GGML_METAL_DECL_KERNEL(silu);
|
|
|
|
GGML_METAL_DECL_KERNEL(relu);
|
2023-06-09 10:00:51 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(gelu);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(soft_max);
|
|
|
|
GGML_METAL_DECL_KERNEL(diag_mask_inf);
|
2023-06-06 19:16:57 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(get_rows_f16);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(get_rows_q4_0);
|
2023-06-10 10:28:11 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(get_rows_q4_1);
|
2023-08-24 15:19:57 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(get_rows_q8_0);
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(get_rows_q2_K);
|
|
|
|
GGML_METAL_DECL_KERNEL(get_rows_q3_K);
|
|
|
|
GGML_METAL_DECL_KERNEL(get_rows_q4_K);
|
|
|
|
GGML_METAL_DECL_KERNEL(get_rows_q5_K);
|
|
|
|
GGML_METAL_DECL_KERNEL(get_rows_q6_K);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(rms_norm);
|
2023-06-17 16:37:49 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(norm);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(mul_mat_f16_f32);
|
2023-06-06 19:16:57 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(mul_mat_q4_0_f32);
|
2023-06-10 10:28:11 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(mul_mat_q4_1_f32);
|
2023-08-24 15:19:57 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(mul_mat_q8_0_f32);
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(mul_mat_q2_K_f32);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_mat_q3_K_f32);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_mat_q4_K_f32);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_mat_q5_K_f32);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_mat_q6_K_f32);
|
2023-08-16 22:07:04 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(mul_mm_f16_f32);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_mm_q4_0_f32);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_mm_q4_1_f32);
|
2023-08-24 15:19:57 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(mul_mm_q8_0_f32);
|
2023-08-16 22:07:04 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(mul_mm_q2_K_f32);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_mm_q3_K_f32);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_mm_q4_K_f32);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_mm_q5_K_f32);
|
|
|
|
GGML_METAL_DECL_KERNEL(mul_mm_q6_K_f32);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(rope);
|
2023-06-17 16:37:49 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(alibi_f32);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(cpy_f32_f16);
|
|
|
|
GGML_METAL_DECL_KERNEL(cpy_f32_f32);
|
2023-06-17 16:37:49 +02:00
|
|
|
GGML_METAL_DECL_KERNEL(cpy_f16_f16);
|
2023-06-04 22:34:30 +02:00
|
|
|
|
|
|
|
#undef GGML_METAL_DECL_KERNEL
|
|
|
|
};
|
|
|
|
|
|
|
|
// MSL code
|
|
|
|
// TODO: move the contents here when ready
|
|
|
|
// for now it is easier to work in a separate file
|
|
|
|
static NSString * const msl_library_source = @"see metal.metal";
|
|
|
|
|
2023-06-10 16:47:34 +02:00
|
|
|
// Here to assist with NSBundle Path Hack
|
|
|
|
@interface GGMLMetalClass : NSObject
|
|
|
|
@end
|
|
|
|
@implementation GGMLMetalClass
|
|
|
|
@end
|
|
|
|
|
2023-07-07 18:24:01 +02:00
|
|
|
struct ggml_metal_context * ggml_metal_init(int n_cb) {
|
2023-06-04 22:34:30 +02:00
|
|
|
fprintf(stderr, "%s: allocating\n", __func__);
|
|
|
|
|
|
|
|
struct ggml_metal_context * ctx = malloc(sizeof(struct ggml_metal_context));
|
|
|
|
|
2023-07-07 18:24:01 +02:00
|
|
|
ctx->n_cb = n_cb;
|
2023-06-04 22:34:30 +02:00
|
|
|
ctx->device = MTLCreateSystemDefaultDevice();
|
|
|
|
ctx->queue = [ctx->device newCommandQueue];
|
2023-06-12 13:31:36 +02:00
|
|
|
ctx->n_buffers = 0;
|
2023-07-25 14:00:19 +02:00
|
|
|
ctx->concur_list_len = 0;
|
2023-06-04 22:34:30 +02:00
|
|
|
|
|
|
|
|
|
|
|
#if 0
|
|
|
|
// compile from source string and show compile log
|
|
|
|
{
|
|
|
|
NSError * error = nil;
|
|
|
|
|
|
|
|
ctx->library = [ctx->device newLibraryWithSource:msl_library_source options:nil error:&error];
|
|
|
|
if (error) {
|
|
|
|
fprintf(stderr, "%s: error: %s\n", __func__, [[error description] UTF8String]);
|
2023-08-14 15:37:39 +02:00
|
|
|
return NULL;
|
2023-06-04 22:34:30 +02:00
|
|
|
}
|
|
|
|
}
|
|
|
|
#else
|
|
|
|
UNUSED(msl_library_source);
|
|
|
|
|
|
|
|
// read the source from "ggml-metal.metal" into a string and use newLibraryWithSource
|
|
|
|
{
|
|
|
|
NSError * error = nil;
|
|
|
|
|
|
|
|
//NSString * path = [[NSBundle mainBundle] pathForResource:@"../../examples/metal/metal" ofType:@"metal"];
|
2023-06-10 16:47:34 +02:00
|
|
|
NSBundle * bundle = [NSBundle bundleForClass:[GGMLMetalClass class]];
|
|
|
|
NSString * path = [bundle pathForResource:@"ggml-metal" ofType:@"metal"];
|
2023-06-04 22:34:30 +02:00
|
|
|
fprintf(stderr, "%s: loading '%s'\n", __func__, [path UTF8String]);
|
|
|
|
|
|
|
|
NSString * src = [NSString stringWithContentsOfFile:path encoding:NSUTF8StringEncoding error:&error];
|
|
|
|
if (error) {
|
|
|
|
fprintf(stderr, "%s: error: %s\n", __func__, [[error description] UTF8String]);
|
2023-08-14 15:37:39 +02:00
|
|
|
return NULL;
|
2023-06-04 22:34:30 +02:00
|
|
|
}
|
|
|
|
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
#ifdef GGML_QKK_64
|
|
|
|
MTLCompileOptions* options = [MTLCompileOptions new];
|
|
|
|
options.preprocessorMacros = @{ @"QK_K" : @(64) };
|
|
|
|
ctx->library = [ctx->device newLibraryWithSource:src options:options error:&error];
|
|
|
|
#else
|
2023-06-04 22:34:30 +02:00
|
|
|
ctx->library = [ctx->device newLibraryWithSource:src options:nil error:&error];
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
#endif
|
2023-06-04 22:34:30 +02:00
|
|
|
if (error) {
|
|
|
|
fprintf(stderr, "%s: error: %s\n", __func__, [[error description] UTF8String]);
|
2023-08-14 15:37:39 +02:00
|
|
|
return NULL;
|
2023-06-04 22:34:30 +02:00
|
|
|
}
|
|
|
|
}
|
|
|
|
#endif
|
|
|
|
|
|
|
|
// load kernels
|
|
|
|
{
|
2023-08-16 22:09:03 +02:00
|
|
|
NSError * error = nil;
|
2023-06-04 22:34:30 +02:00
|
|
|
#define GGML_METAL_ADD_KERNEL(name) \
|
|
|
|
ctx->function_##name = [ctx->library newFunctionWithName:@"kernel_"#name]; \
|
2023-08-16 22:09:03 +02:00
|
|
|
ctx->pipeline_##name = [ctx->device newComputePipelineStateWithFunction:ctx->function_##name error:&error]; \
|
2023-08-23 22:08:04 +02:00
|
|
|
fprintf(stderr, "%s: loaded %-32s %16p | th_max = %4d | th_width = %4d\n", __func__, "kernel_"#name, (void *) ctx->pipeline_##name, \
|
|
|
|
(int) ctx->pipeline_##name.maxTotalThreadsPerThreadgroup, \
|
|
|
|
(int) ctx->pipeline_##name.threadExecutionWidth); \
|
2023-08-16 22:09:03 +02:00
|
|
|
if (error) { \
|
|
|
|
fprintf(stderr, "%s: load pipeline error: %s\n", __func__, [[error description] UTF8String]); \
|
|
|
|
return NULL; \
|
|
|
|
}
|
2023-06-04 22:34:30 +02:00
|
|
|
|
|
|
|
GGML_METAL_ADD_KERNEL(add);
|
2023-07-23 13:00:37 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(add_row);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(mul);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_row);
|
|
|
|
GGML_METAL_ADD_KERNEL(scale);
|
|
|
|
GGML_METAL_ADD_KERNEL(silu);
|
|
|
|
GGML_METAL_ADD_KERNEL(relu);
|
2023-06-09 10:00:51 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(gelu);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(soft_max);
|
|
|
|
GGML_METAL_ADD_KERNEL(diag_mask_inf);
|
2023-06-06 19:16:57 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(get_rows_f16);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(get_rows_q4_0);
|
2023-06-10 10:28:11 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(get_rows_q4_1);
|
2023-08-24 15:19:57 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(get_rows_q8_0);
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(get_rows_q2_K);
|
|
|
|
GGML_METAL_ADD_KERNEL(get_rows_q3_K);
|
|
|
|
GGML_METAL_ADD_KERNEL(get_rows_q4_K);
|
|
|
|
GGML_METAL_ADD_KERNEL(get_rows_q5_K);
|
|
|
|
GGML_METAL_ADD_KERNEL(get_rows_q6_K);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(rms_norm);
|
2023-06-17 16:37:49 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(norm);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(mul_mat_f16_f32);
|
2023-06-06 19:16:57 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(mul_mat_q4_0_f32);
|
2023-06-10 10:28:11 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(mul_mat_q4_1_f32);
|
2023-08-24 15:19:57 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(mul_mat_q8_0_f32);
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(mul_mat_q2_K_f32);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_mat_q3_K_f32);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_mat_q4_K_f32);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_mat_q5_K_f32);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_mat_q6_K_f32);
|
2023-08-16 22:07:04 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(mul_mm_f16_f32);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_mm_q4_0_f32);
|
2023-08-24 15:19:57 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(mul_mm_q8_0_f32);
|
2023-08-16 22:07:04 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(mul_mm_q4_1_f32);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_mm_q2_K_f32);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_mm_q3_K_f32);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_mm_q4_K_f32);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_mm_q5_K_f32);
|
|
|
|
GGML_METAL_ADD_KERNEL(mul_mm_q6_K_f32);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(rope);
|
2023-06-17 16:37:49 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(alibi_f32);
|
2023-06-04 22:34:30 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(cpy_f32_f16);
|
|
|
|
GGML_METAL_ADD_KERNEL(cpy_f32_f32);
|
2023-06-17 16:37:49 +02:00
|
|
|
GGML_METAL_ADD_KERNEL(cpy_f16_f16);
|
2023-06-04 22:34:30 +02:00
|
|
|
|
|
|
|
#undef GGML_METAL_ADD_KERNEL
|
|
|
|
}
|
|
|
|
|
2023-08-23 22:08:04 +02:00
|
|
|
fprintf(stderr, "%s: recommendedMaxWorkingSetSize = %8.2f MB\n", __func__, ctx->device.recommendedMaxWorkingSetSize / 1024.0 / 1024.0);
|
|
|
|
fprintf(stderr, "%s: hasUnifiedMemory = %s\n", __func__, ctx->device.hasUnifiedMemory ? "true" : "false");
|
2023-06-18 08:09:47 +02:00
|
|
|
if (ctx->device.maxTransferRate != 0) {
|
2023-08-23 22:08:04 +02:00
|
|
|
fprintf(stderr, "%s: maxTransferRate = %8.2f MB/s\n", __func__, ctx->device.maxTransferRate / 1024.0 / 1024.0);
|
2023-06-18 08:09:47 +02:00
|
|
|
} else {
|
2023-08-23 22:08:04 +02:00
|
|
|
fprintf(stderr, "%s: maxTransferRate = built-in GPU\n", __func__);
|
2023-06-18 08:09:47 +02:00
|
|
|
}
|
|
|
|
|
2023-06-04 22:34:30 +02:00
|
|
|
return ctx;
|
|
|
|
}
|
|
|
|
|
|
|
|
void ggml_metal_free(struct ggml_metal_context * ctx) {
|
|
|
|
fprintf(stderr, "%s: deallocating\n", __func__);
|
2023-07-01 20:14:59 +02:00
|
|
|
for (int i = 0; i < ctx->n_buffers; ++i) {
|
|
|
|
[ctx->buffers[i].metal release];
|
|
|
|
}
|
2023-06-04 22:34:30 +02:00
|
|
|
free(ctx);
|
|
|
|
}
|
|
|
|
|
2023-08-21 22:07:43 +02:00
|
|
|
void * ggml_metal_host_malloc(size_t n) {
|
|
|
|
void * data = NULL;
|
|
|
|
const int result = posix_memalign((void **) &data, getpagesize(), n);
|
|
|
|
if (result != 0) {
|
|
|
|
fprintf(stderr, "%s: error: posix_memalign failed\n", __func__);
|
|
|
|
return NULL;
|
|
|
|
}
|
|
|
|
|
|
|
|
return data;
|
|
|
|
}
|
|
|
|
|
|
|
|
void ggml_metal_host_free(void * data) {
|
|
|
|
free(data);
|
|
|
|
}
|
|
|
|
|
2023-07-07 18:24:01 +02:00
|
|
|
void ggml_metal_set_n_cb(struct ggml_metal_context * ctx, int n_cb) {
|
|
|
|
ctx->n_cb = n_cb;
|
|
|
|
}
|
|
|
|
|
2023-08-16 22:08:28 +02:00
|
|
|
int ggml_metal_if_optimized(struct ggml_metal_context * ctx) {
|
|
|
|
return ctx->concur_list_len;
|
|
|
|
}
|
|
|
|
|
|
|
|
int * ggml_metal_get_concur_list(struct ggml_metal_context * ctx) {
|
|
|
|
return ctx->concur_list;
|
2023-07-25 14:00:19 +02:00
|
|
|
}
|
|
|
|
|
2023-06-04 22:34:30 +02:00
|
|
|
// finds the Metal buffer that contains the tensor data on the GPU device
|
|
|
|
// the assumption is that there is 1-to-1 mapping between the host and device memory buffers, so we can find the
|
|
|
|
// Metal buffer based on the host memory pointer
|
|
|
|
//
|
|
|
|
static id<MTLBuffer> ggml_metal_get_buffer(struct ggml_metal_context * ctx, struct ggml_tensor * t, size_t * offs) {
|
|
|
|
//fprintf(stderr, "%s: data tensor '%16s', offs_data = %8ld, offs_eval = %8ld, offs_cach = %8ld\n", __func__, t->name, offs_data, offs_eval, offs_cach);
|
|
|
|
|
2023-06-18 08:09:47 +02:00
|
|
|
const int64_t tsize = ggml_nbytes(t);
|
|
|
|
|
|
|
|
// find the view that contains the tensor fully
|
2023-06-04 22:34:30 +02:00
|
|
|
for (int i = 0; i < ctx->n_buffers; ++i) {
|
|
|
|
const int64_t ioffs = (int64_t) t->data - (int64_t) ctx->buffers[i].data;
|
|
|
|
|
2023-06-18 08:09:47 +02:00
|
|
|
if (ioffs >= 0 && ioffs + tsize <= (int64_t) ctx->buffers[i].size) {
|
2023-06-04 22:34:30 +02:00
|
|
|
*offs = (size_t) ioffs;
|
|
|
|
|
|
|
|
//fprintf(stderr, "%s: '%s' tensor '%16s', offs = %8ld\n", __func__, ctx->buffers[i].name, t->name, *offs);
|
|
|
|
|
|
|
|
return ctx->buffers[i].metal;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
fprintf(stderr, "%s: error: buffer is nil\n", __func__);
|
|
|
|
|
|
|
|
return nil;
|
|
|
|
}
|
|
|
|
|
|
|
|
bool ggml_metal_add_buffer(
|
|
|
|
struct ggml_metal_context * ctx,
|
|
|
|
const char * name,
|
|
|
|
void * data,
|
2023-06-18 08:09:47 +02:00
|
|
|
size_t size,
|
|
|
|
size_t max_size) {
|
2023-06-04 22:34:30 +02:00
|
|
|
if (ctx->n_buffers >= GGML_METAL_MAX_BUFFERS) {
|
|
|
|
fprintf(stderr, "%s: too many buffers\n", __func__);
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
|
|
|
if (data) {
|
|
|
|
// verify that the buffer does not overlap with any of the existing buffers
|
|
|
|
for (int i = 0; i < ctx->n_buffers; ++i) {
|
|
|
|
const int64_t ioffs = (int64_t) data - (int64_t) ctx->buffers[i].data;
|
|
|
|
|
|
|
|
if (ioffs >= 0 && ioffs < (int64_t) ctx->buffers[i].size) {
|
|
|
|
fprintf(stderr, "%s: error: buffer '%s' overlaps with '%s'\n", __func__, name, ctx->buffers[i].name);
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2023-06-18 08:09:47 +02:00
|
|
|
const size_t size_page = getpagesize();
|
|
|
|
|
|
|
|
size_t size_aligned = size;
|
|
|
|
if ((size_aligned % size_page) != 0) {
|
|
|
|
size_aligned += (size_page - (size_aligned % size_page));
|
2023-06-05 22:24:04 +02:00
|
|
|
}
|
|
|
|
|
2023-06-18 08:09:47 +02:00
|
|
|
// the buffer fits into the max buffer size allowed by the device
|
|
|
|
if (size_aligned <= ctx->device.maxBufferLength) {
|
|
|
|
ctx->buffers[ctx->n_buffers].name = name;
|
|
|
|
ctx->buffers[ctx->n_buffers].data = data;
|
|
|
|
ctx->buffers[ctx->n_buffers].size = size;
|
2023-06-06 05:28:17 +02:00
|
|
|
|
2023-06-18 08:09:47 +02:00
|
|
|
ctx->buffers[ctx->n_buffers].metal = [ctx->device newBufferWithBytesNoCopy:data length:size_aligned options:MTLResourceStorageModeShared deallocator:nil];
|
|
|
|
|
|
|
|
if (ctx->buffers[ctx->n_buffers].metal == nil) {
|
|
|
|
fprintf(stderr, "%s: failed to allocate '%-16s' buffer, size = %8.2f MB\n", __func__, name, size_aligned / 1024.0 / 1024.0);
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
|
|
|
fprintf(stderr, "%s: allocated '%-16s' buffer, size = %8.2f MB", __func__, name, size_aligned / 1024.0 / 1024.0);
|
|
|
|
|
|
|
|
++ctx->n_buffers;
|
|
|
|
} else {
|
|
|
|
// this overlap between the views will guarantee that the tensor with the maximum size will fully fit into
|
|
|
|
// one of the views
|
|
|
|
const size_t size_ovlp = ((max_size + size_page - 1) / size_page + 1) * size_page; // round-up 2 pages just in case
|
|
|
|
const size_t size_step = ctx->device.maxBufferLength - size_ovlp;
|
|
|
|
const size_t size_view = ctx->device.maxBufferLength;
|
|
|
|
|
|
|
|
for (size_t i = 0; i < size; i += size_step) {
|
|
|
|
const size_t size_step_aligned = (i + size_view <= size) ? size_view : (size_aligned - i);
|
2023-06-04 22:34:30 +02:00
|
|
|
|
2023-06-18 08:09:47 +02:00
|
|
|
ctx->buffers[ctx->n_buffers].name = name;
|
|
|
|
ctx->buffers[ctx->n_buffers].data = (void *) ((uint8_t *) data + i);
|
|
|
|
ctx->buffers[ctx->n_buffers].size = size_step_aligned;
|
|
|
|
|
|
|
|
ctx->buffers[ctx->n_buffers].metal = [ctx->device newBufferWithBytesNoCopy:(void *) ((uint8_t *) data + i) length:size_step_aligned options:MTLResourceStorageModeShared deallocator:nil];
|
|
|
|
|
|
|
|
if (ctx->buffers[ctx->n_buffers].metal == nil) {
|
|
|
|
fprintf(stderr, "%s: failed to allocate '%-16s' buffer, size = %8.2f MB\n", __func__, name, size_step_aligned / 1024.0 / 1024.0);
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
|
|
|
fprintf(stderr, "%s: allocated '%-16s' buffer, size = %8.2f MB, offs = %12ld", __func__, name, size_step_aligned / 1024.0 / 1024.0, i);
|
|
|
|
if (i + size_step < size) {
|
|
|
|
fprintf(stderr, "\n");
|
|
|
|
}
|
|
|
|
|
|
|
|
++ctx->n_buffers;
|
|
|
|
}
|
2023-06-05 22:24:04 +02:00
|
|
|
}
|
2023-06-04 22:34:30 +02:00
|
|
|
|
2023-06-18 08:09:47 +02:00
|
|
|
fprintf(stderr, ", (%8.2f / %8.2f)",
|
|
|
|
ctx->device.currentAllocatedSize / 1024.0 / 1024.0,
|
|
|
|
ctx->device.recommendedMaxWorkingSetSize / 1024.0 / 1024.0);
|
2023-06-17 19:24:11 +02:00
|
|
|
|
2023-06-18 08:09:47 +02:00
|
|
|
if (ctx->device.currentAllocatedSize > ctx->device.recommendedMaxWorkingSetSize) {
|
|
|
|
fprintf(stderr, ", warning: current allocated size is greater than the recommended max working set size\n");
|
|
|
|
} else {
|
|
|
|
fprintf(stderr, "\n");
|
|
|
|
}
|
2023-06-04 22:34:30 +02:00
|
|
|
}
|
|
|
|
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
void ggml_metal_set_tensor(
|
|
|
|
struct ggml_metal_context * ctx,
|
|
|
|
struct ggml_tensor * t) {
|
|
|
|
metal_printf("%s: set input for tensor '%s'\n", __func__, t->name);
|
|
|
|
|
|
|
|
size_t offs;
|
|
|
|
id<MTLBuffer> id_dst = ggml_metal_get_buffer(ctx, t, &offs);
|
|
|
|
|
|
|
|
memcpy((void *) ((uint8_t *) id_dst.contents + offs), t->data, ggml_nbytes(t));
|
|
|
|
}
|
|
|
|
|
|
|
|
void ggml_metal_get_tensor(
|
|
|
|
struct ggml_metal_context * ctx,
|
|
|
|
struct ggml_tensor * t) {
|
|
|
|
metal_printf("%s: extract results for tensor '%s'\n", __func__, t->name);
|
|
|
|
|
|
|
|
size_t offs;
|
|
|
|
id<MTLBuffer> id_src = ggml_metal_get_buffer(ctx, t, &offs);
|
|
|
|
|
|
|
|
memcpy(t->data, (void *) ((uint8_t *) id_src.contents + offs), ggml_nbytes(t));
|
|
|
|
}
|
|
|
|
|
2023-07-25 14:00:19 +02:00
|
|
|
void ggml_metal_graph_find_concurrency(
|
|
|
|
struct ggml_metal_context * ctx,
|
2023-08-16 22:08:28 +02:00
|
|
|
struct ggml_cgraph * gf, bool check_mem) {
|
2023-07-25 14:00:19 +02:00
|
|
|
int search_depth = gf->n_nodes; //we only find concurrency in this range to avoid wasting too much time
|
2023-08-07 09:52:57 +02:00
|
|
|
int nodes_unused[GGML_MAX_CONCUR];
|
2023-07-25 14:00:19 +02:00
|
|
|
|
2023-08-07 09:52:57 +02:00
|
|
|
for (int i = 0; i < GGML_MAX_CONCUR; i++) { ctx->concur_list[i] = 0; }
|
|
|
|
for (int i = 0; i < gf->n_nodes; i++) { nodes_unused[i] = 1; }
|
2023-07-25 14:00:19 +02:00
|
|
|
ctx->concur_list_len = 0;
|
|
|
|
|
2023-08-07 09:52:57 +02:00
|
|
|
int n_left = gf->n_nodes;
|
|
|
|
int n_start = 0; // all nodes before n_start at nodes_unused array have been sorted and store back to ctx->concur_list
|
|
|
|
int level_pos = 0; // at ctx->concur_list, the last layer (level) ends at level_pos
|
2023-07-25 14:00:19 +02:00
|
|
|
|
|
|
|
while (n_left > 0) {
|
|
|
|
// number of nodes at a layer (that can be issued concurrently)
|
|
|
|
int concurrency = 0;
|
|
|
|
for (int i = n_start; i < ((n_start + search_depth > gf->n_nodes) ? gf->n_nodes : n_start + search_depth); i++) {
|
|
|
|
if (nodes_unused[i]) {
|
|
|
|
// if the requirements for gf->nodes[i] are satisfied
|
2023-08-07 09:52:57 +02:00
|
|
|
int exe_flag = 1;
|
|
|
|
|
2023-07-25 14:00:19 +02:00
|
|
|
// scan all srcs
|
|
|
|
for (int src_ind = 0; src_ind < GGML_MAX_SRC; src_ind++) {
|
|
|
|
struct ggml_tensor * src_cur = gf->nodes[i]->src[src_ind];
|
|
|
|
if (src_cur) {
|
|
|
|
// if is leaf nodes it's satisfied.
|
2023-08-07 09:52:57 +02:00
|
|
|
// TODO: ggml_is_leaf()
|
|
|
|
if (src_cur->op == GGML_OP_NONE && src_cur->grad == NULL) {
|
|
|
|
continue;
|
|
|
|
}
|
2023-07-25 14:00:19 +02:00
|
|
|
|
|
|
|
// otherwise this src should be the output from previous nodes.
|
|
|
|
int is_found = 0;
|
2023-08-07 09:52:57 +02:00
|
|
|
|
2023-07-25 14:00:19 +02:00
|
|
|
// scan 2*search_depth back because we inserted barrier.
|
2023-08-07 09:52:57 +02:00
|
|
|
//for (int j = ((level_pos - 2*search_depth) < 0 ? 0 : (level_pos - 2*search_depth)); j < level_pos; j++) {
|
|
|
|
for (int j = MAX(0, level_pos - 2*search_depth); j < level_pos; j++) {
|
|
|
|
if (ctx->concur_list[j] >= 0 && gf->nodes[ctx->concur_list[j]] == src_cur) {
|
|
|
|
is_found = 1;
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
if (is_found == 0) {
|
|
|
|
exe_flag = 0;
|
|
|
|
break;
|
2023-07-25 14:00:19 +02:00
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
2023-08-16 22:08:28 +02:00
|
|
|
if (exe_flag && check_mem) {
|
2023-07-25 14:00:19 +02:00
|
|
|
// check if nodes[i]'s data will be overwritten by a node before nodes[i].
|
|
|
|
// if node[5] and node[3] write to the same memory region, then we can't issue node[5] before node[3]
|
|
|
|
int64_t data_start = (int64_t) gf->nodes[i]->data;
|
2023-08-07 09:52:57 +02:00
|
|
|
int64_t length = (int64_t) ggml_nbytes(gf->nodes[i]);
|
2023-07-25 14:00:19 +02:00
|
|
|
for (int j = n_start; j < i; j++) {
|
|
|
|
if (nodes_unused[j] && gf->nodes[j]->op != GGML_OP_RESHAPE \
|
|
|
|
&& gf->nodes[j]->op != GGML_OP_VIEW \
|
|
|
|
&& gf->nodes[j]->op != GGML_OP_TRANSPOSE \
|
|
|
|
&& gf->nodes[j]->op != GGML_OP_PERMUTE) {
|
|
|
|
if (((int64_t)gf->nodes[j]->data) >= data_start + length || \
|
|
|
|
((int64_t)gf->nodes[j]->data) + (int64_t) ggml_nbytes(gf->nodes[j]) <= data_start) {
|
|
|
|
continue;
|
|
|
|
}
|
2023-08-07 09:52:57 +02:00
|
|
|
|
|
|
|
exe_flag = 0;
|
2023-07-25 14:00:19 +02:00
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
if (exe_flag) {
|
|
|
|
ctx->concur_list[level_pos + concurrency] = i;
|
|
|
|
nodes_unused[i] = 0;
|
|
|
|
concurrency++;
|
|
|
|
ctx->concur_list_len++;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
n_left -= concurrency;
|
|
|
|
// adding a barrier different layer
|
|
|
|
ctx->concur_list[level_pos + concurrency] = -1;
|
|
|
|
ctx->concur_list_len++;
|
|
|
|
// jump all sorted nodes at nodes_bak
|
2023-08-07 09:52:57 +02:00
|
|
|
while (!nodes_unused[n_start]) {
|
|
|
|
n_start++;
|
|
|
|
}
|
2023-07-25 14:00:19 +02:00
|
|
|
level_pos += concurrency + 1;
|
|
|
|
}
|
|
|
|
|
2023-08-07 09:52:57 +02:00
|
|
|
if (ctx->concur_list_len > GGML_MAX_CONCUR) {
|
2023-07-25 14:00:19 +02:00
|
|
|
fprintf(stderr, "%s: too many elements for metal ctx->concur_list!\n", __func__);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2023-06-04 22:34:30 +02:00
|
|
|
void ggml_metal_graph_compute(
|
|
|
|
struct ggml_metal_context * ctx,
|
2023-06-15 19:29:48 +02:00
|
|
|
struct ggml_cgraph * gf) {
|
2023-06-04 22:34:30 +02:00
|
|
|
metal_printf("%s: evaluating graph\n", __func__);
|
|
|
|
|
2023-07-25 14:00:19 +02:00
|
|
|
// if there is ctx->concur_list, dispatch concurrently
|
|
|
|
// else fallback to serial dispatch
|
|
|
|
MTLComputePassDescriptor * edesc = MTLComputePassDescriptor.computePassDescriptor;
|
|
|
|
|
2023-08-07 09:52:57 +02:00
|
|
|
const bool has_concur = ctx->concur_list_len && ctx->concur_list_len <= GGML_MAX_CONCUR;
|
2023-07-25 14:00:19 +02:00
|
|
|
|
|
|
|
const int n_nodes = has_concur ? ctx->concur_list_len : gf->n_nodes;
|
|
|
|
edesc.dispatchType = has_concur ? MTLDispatchTypeConcurrent : MTLDispatchTypeSerial;
|
|
|
|
|
2023-06-15 19:29:48 +02:00
|
|
|
// create multiple command buffers and enqueue them
|
|
|
|
// then, we encode the graph into the command buffers in parallel
|
|
|
|
|
2023-07-07 18:24:01 +02:00
|
|
|
const int n_cb = ctx->n_cb;
|
2023-06-15 19:29:48 +02:00
|
|
|
|
|
|
|
NSMutableArray * command_buffers = [NSMutableArray arrayWithCapacity:n_cb];
|
|
|
|
|
|
|
|
for (int i = 0; i < n_cb; ++i) {
|
|
|
|
command_buffers[i] = [ctx->queue commandBuffer];
|
2023-06-04 22:34:30 +02:00
|
|
|
|
2023-06-15 19:29:48 +02:00
|
|
|
// enqueue the command buffers in order to specify their execution order
|
|
|
|
[command_buffers[i] enqueue];
|
2023-06-04 22:34:30 +02:00
|
|
|
}
|
|
|
|
|
2023-06-15 19:29:48 +02:00
|
|
|
// TODO: is this the best way to start threads?
|
|
|
|
dispatch_queue_t queue = dispatch_queue_create("llama.cpp", DISPATCH_QUEUE_CONCURRENT);
|
|
|
|
|
|
|
|
for (int cb_idx = 0; cb_idx < n_cb; ++cb_idx) {
|
2023-07-25 14:00:19 +02:00
|
|
|
const int n_nodes_per_cb = (n_nodes + n_cb - 1) / n_cb;
|
2023-06-15 19:29:48 +02:00
|
|
|
|
|
|
|
dispatch_async(queue, ^{
|
|
|
|
size_t offs_src0 = 0;
|
|
|
|
size_t offs_src1 = 0;
|
|
|
|
size_t offs_dst = 0;
|
|
|
|
|
|
|
|
id<MTLCommandBuffer> command_buffer = command_buffers[cb_idx];
|
|
|
|
|
2023-08-16 22:07:04 +02:00
|
|
|
id<MTLComputeCommandEncoder> encoder = [command_buffer computeCommandEncoderWithDescriptor: edesc];
|
2023-06-15 19:29:48 +02:00
|
|
|
|
2023-08-23 22:08:04 +02:00
|
|
|
const int node_start = (cb_idx + 0) * n_nodes_per_cb;
|
|
|
|
const int node_end = MIN((cb_idx == n_cb - 1) ? n_nodes : (cb_idx + 1) * n_nodes_per_cb, n_nodes);
|
2023-07-25 14:00:19 +02:00
|
|
|
|
|
|
|
for (int ind = node_start; ind < node_end; ++ind) {
|
|
|
|
const int i = has_concur ? ctx->concur_list[ind] : ind;
|
|
|
|
|
|
|
|
if (i == -1) {
|
|
|
|
[encoder memoryBarrierWithScope:MTLBarrierScopeBuffers];
|
|
|
|
continue;
|
|
|
|
}
|
2023-06-15 19:29:48 +02:00
|
|
|
|
|
|
|
metal_printf("%s: encoding node %3d, op = %8s\n", __func__, i, ggml_op_name(gf->nodes[i]->op));
|
|
|
|
|
2023-07-11 18:31:10 +02:00
|
|
|
struct ggml_tensor * src0 = gf->nodes[i]->src[0];
|
|
|
|
struct ggml_tensor * src1 = gf->nodes[i]->src[1];
|
2023-06-15 19:29:48 +02:00
|
|
|
struct ggml_tensor * dst = gf->nodes[i];
|
|
|
|
|
|
|
|
const int64_t ne00 = src0 ? src0->ne[0] : 0;
|
|
|
|
const int64_t ne01 = src0 ? src0->ne[1] : 0;
|
|
|
|
const int64_t ne02 = src0 ? src0->ne[2] : 0;
|
|
|
|
const int64_t ne03 = src0 ? src0->ne[3] : 0;
|
|
|
|
|
|
|
|
const uint64_t nb00 = src0 ? src0->nb[0] : 0;
|
|
|
|
const uint64_t nb01 = src0 ? src0->nb[1] : 0;
|
|
|
|
const uint64_t nb02 = src0 ? src0->nb[2] : 0;
|
|
|
|
const uint64_t nb03 = src0 ? src0->nb[3] : 0;
|
|
|
|
|
|
|
|
const int64_t ne10 = src1 ? src1->ne[0] : 0;
|
|
|
|
const int64_t ne11 = src1 ? src1->ne[1] : 0;
|
|
|
|
const int64_t ne12 = src1 ? src1->ne[2] : 0;
|
|
|
|
const int64_t ne13 = src1 ? src1->ne[3] : 0; UNUSED(ne13);
|
|
|
|
|
|
|
|
const uint64_t nb10 = src1 ? src1->nb[0] : 0;
|
|
|
|
const uint64_t nb11 = src1 ? src1->nb[1] : 0;
|
|
|
|
const uint64_t nb12 = src1 ? src1->nb[2] : 0;
|
|
|
|
const uint64_t nb13 = src1 ? src1->nb[3] : 0; UNUSED(nb13);
|
|
|
|
|
|
|
|
const int64_t ne0 = dst ? dst->ne[0] : 0;
|
|
|
|
const int64_t ne1 = dst ? dst->ne[1] : 0;
|
|
|
|
const int64_t ne2 = dst ? dst->ne[2] : 0;
|
|
|
|
const int64_t ne3 = dst ? dst->ne[3] : 0;
|
|
|
|
|
|
|
|
const uint64_t nb0 = dst ? dst->nb[0] : 0;
|
|
|
|
const uint64_t nb1 = dst ? dst->nb[1] : 0;
|
|
|
|
const uint64_t nb2 = dst ? dst->nb[2] : 0;
|
|
|
|
const uint64_t nb3 = dst ? dst->nb[3] : 0;
|
|
|
|
|
|
|
|
const enum ggml_type src0t = src0 ? src0->type : GGML_TYPE_COUNT;
|
|
|
|
const enum ggml_type src1t = src1 ? src1->type : GGML_TYPE_COUNT;
|
|
|
|
const enum ggml_type dstt = dst ? dst->type : GGML_TYPE_COUNT;
|
|
|
|
|
|
|
|
id<MTLBuffer> id_src0 = src0 ? ggml_metal_get_buffer(ctx, src0, &offs_src0) : nil;
|
|
|
|
id<MTLBuffer> id_src1 = src1 ? ggml_metal_get_buffer(ctx, src1, &offs_src1) : nil;
|
|
|
|
id<MTLBuffer> id_dst = dst ? ggml_metal_get_buffer(ctx, dst, &offs_dst) : nil;
|
|
|
|
|
|
|
|
//metal_printf("%s: op - %s\n", __func__, ggml_op_name(dst->op));
|
|
|
|
//if (src0) {
|
|
|
|
// metal_printf("%s: src0 - %4s [%5lld, %5lld, %5lld], %d, %s\n", __func__, ggml_type_name(src0t), ne00, ne01, ne02,
|
|
|
|
// ggml_is_contiguous(src0), src0->name);
|
|
|
|
//}
|
|
|
|
//if (src1) {
|
|
|
|
// metal_printf("%s: src1 - %4s [%5lld, %5lld, %5lld], %d, %s\n", __func__, ggml_type_name(src1t), ne10, ne11, ne12,
|
|
|
|
// ggml_is_contiguous(src1), src1->name);
|
|
|
|
//}
|
|
|
|
//if (dst) {
|
|
|
|
// metal_printf("%s: dst - %4s [%5lld, %5lld, %5lld], 1, %s\n", __func__, ggml_type_name(dstt), ne0, ne1, ne2,
|
|
|
|
// dst->name);
|
|
|
|
//}
|
|
|
|
|
|
|
|
switch (dst->op) {
|
2023-07-10 17:49:56 +02:00
|
|
|
case GGML_OP_NONE:
|
2023-06-15 19:29:48 +02:00
|
|
|
case GGML_OP_RESHAPE:
|
|
|
|
case GGML_OP_VIEW:
|
|
|
|
case GGML_OP_TRANSPOSE:
|
|
|
|
case GGML_OP_PERMUTE:
|
|
|
|
{
|
|
|
|
// noop
|
|
|
|
} break;
|
|
|
|
case GGML_OP_ADD:
|
|
|
|
{
|
2023-07-23 13:00:37 +02:00
|
|
|
if (ggml_nelements(src1) == ne10) {
|
|
|
|
// src1 is a row
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_add_row];
|
|
|
|
} else {
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_add];
|
|
|
|
}
|
2023-06-15 19:29:48 +02:00
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_src1 offset:offs_src1 atIndex:1];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:2];
|
2023-07-23 13:00:37 +02:00
|
|
|
[encoder setBytes:&ne00 length:sizeof(ne00) atIndex:3];
|
2023-06-15 19:29:48 +02:00
|
|
|
|
|
|
|
const int64_t n = ggml_nelements(dst);
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
|
|
|
|
} break;
|
|
|
|
case GGML_OP_MUL:
|
|
|
|
{
|
|
|
|
if (ggml_nelements(src1) == ne10) {
|
|
|
|
// src1 is a row
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul_row];
|
|
|
|
} else {
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul];
|
|
|
|
}
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_src1 offset:offs_src1 atIndex:1];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:2];
|
|
|
|
[encoder setBytes:&ne00 length:sizeof(ne00) atIndex:3];
|
|
|
|
|
|
|
|
const int64_t n = ggml_nelements(dst);
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
|
|
|
|
} break;
|
|
|
|
case GGML_OP_SCALE:
|
|
|
|
{
|
|
|
|
const float scale = *(const float *) src1->data;
|
|
|
|
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_scale];
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
|
|
|
[encoder setBytes:&scale length:sizeof(scale) atIndex:2];
|
|
|
|
|
|
|
|
const int64_t n = ggml_nelements(dst);
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
|
|
|
|
} break;
|
2023-07-24 13:46:21 +02:00
|
|
|
case GGML_OP_UNARY:
|
|
|
|
switch (ggml_get_unary_op(gf->nodes[i])) {
|
|
|
|
case GGML_UNARY_OP_SILU:
|
|
|
|
{
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_silu];
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
|
|
|
|
|
|
|
const int64_t n = ggml_nelements(dst);
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
|
|
|
|
} break;
|
|
|
|
case GGML_UNARY_OP_RELU:
|
|
|
|
{
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_relu];
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
|
|
|
|
|
|
|
const int64_t n = ggml_nelements(dst);
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
|
|
|
|
} break;
|
|
|
|
case GGML_UNARY_OP_GELU:
|
|
|
|
{
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_gelu];
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
|
|
|
|
|
|
|
const int64_t n = ggml_nelements(dst);
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
|
|
|
|
} break;
|
|
|
|
default:
|
|
|
|
{
|
|
|
|
fprintf(stderr, "%s: node %3d, op = %8s not implemented\n", __func__, i, ggml_op_name(dst->op));
|
|
|
|
GGML_ASSERT(false);
|
|
|
|
}
|
2023-06-15 19:29:48 +02:00
|
|
|
} break;
|
|
|
|
case GGML_OP_SOFT_MAX:
|
|
|
|
{
|
|
|
|
const int nth = 32;
|
|
|
|
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_soft_max];
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
|
|
|
[encoder setBytes:&ne00 length:sizeof(ne00) atIndex:2];
|
|
|
|
[encoder setBytes:&ne01 length:sizeof(ne01) atIndex:3];
|
|
|
|
[encoder setBytes:&ne02 length:sizeof(ne02) atIndex:4];
|
|
|
|
[encoder setThreadgroupMemoryLength:nth*sizeof(float) atIndex:0];
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(ne01, ne02, ne03) threadsPerThreadgroup:MTLSizeMake(nth, 1, 1)];
|
|
|
|
} break;
|
|
|
|
case GGML_OP_DIAG_MASK_INF:
|
|
|
|
{
|
2023-07-23 14:36:02 +02:00
|
|
|
const int n_past = ((int32_t *)(dst->op_params))[0];
|
2023-06-15 19:29:48 +02:00
|
|
|
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_diag_mask_inf];
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
|
|
|
[encoder setBytes:&ne00 length:sizeof(ne00) atIndex:2];
|
|
|
|
[encoder setBytes:&ne01 length:sizeof(ne01) atIndex:3];
|
|
|
|
[encoder setBytes:&n_past length:sizeof(int) atIndex:4];
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(ne00, ne01, ne02) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
|
|
|
|
} break;
|
|
|
|
case GGML_OP_MUL_MAT:
|
|
|
|
{
|
|
|
|
// TODO: needs to be updated after PR: https://github.com/ggerganov/ggml/pull/224
|
|
|
|
|
|
|
|
GGML_ASSERT(ne00 == ne10);
|
2023-08-01 09:43:12 +02:00
|
|
|
// GGML_ASSERT(ne02 == ne12); // Should be checked on individual data types until broadcast is implemented everywhere
|
2023-08-16 22:07:04 +02:00
|
|
|
uint gqa = ne12/ne02;
|
2023-08-01 09:43:12 +02:00
|
|
|
GGML_ASSERT(ne03 == ne13);
|
2023-06-15 19:29:48 +02:00
|
|
|
|
2023-08-16 22:07:04 +02:00
|
|
|
// for now the matrix-matrix multiplication kernel only works on A14+/M1+ SoCs
|
|
|
|
// AMD GPU and older A-chips will reuse matrix-vector multiplication kernel
|
2023-06-15 19:29:48 +02:00
|
|
|
if (ggml_is_contiguous(src0) &&
|
|
|
|
ggml_is_contiguous(src1) &&
|
2023-08-16 22:07:04 +02:00
|
|
|
src1t == GGML_TYPE_F32 &&
|
|
|
|
[ctx->device supportsFamily:MTLGPUFamilyApple7] &&
|
|
|
|
ne00%32 == 0 &&
|
|
|
|
ne11 > 1) {
|
2023-08-23 22:08:04 +02:00
|
|
|
switch (src0->type) {
|
2023-08-24 15:19:57 +02:00
|
|
|
case GGML_TYPE_F16: [encoder setComputePipelineState:ctx->pipeline_mul_mm_f16_f32]; break;
|
2023-08-23 22:08:04 +02:00
|
|
|
case GGML_TYPE_Q4_0: [encoder setComputePipelineState:ctx->pipeline_mul_mm_q4_0_f32]; break;
|
|
|
|
case GGML_TYPE_Q4_1: [encoder setComputePipelineState:ctx->pipeline_mul_mm_q4_1_f32]; break;
|
2023-08-24 15:19:57 +02:00
|
|
|
case GGML_TYPE_Q8_0: [encoder setComputePipelineState:ctx->pipeline_mul_mm_q8_0_f32]; break;
|
2023-08-23 22:08:04 +02:00
|
|
|
case GGML_TYPE_Q2_K: [encoder setComputePipelineState:ctx->pipeline_mul_mm_q2_K_f32]; break;
|
|
|
|
case GGML_TYPE_Q3_K: [encoder setComputePipelineState:ctx->pipeline_mul_mm_q3_K_f32]; break;
|
|
|
|
case GGML_TYPE_Q4_K: [encoder setComputePipelineState:ctx->pipeline_mul_mm_q4_K_f32]; break;
|
|
|
|
case GGML_TYPE_Q5_K: [encoder setComputePipelineState:ctx->pipeline_mul_mm_q5_K_f32]; break;
|
|
|
|
case GGML_TYPE_Q6_K: [encoder setComputePipelineState:ctx->pipeline_mul_mm_q6_K_f32]; break;
|
|
|
|
default: GGML_ASSERT(false && "MUL MAT-MAT not implemented");
|
2023-06-15 19:29:48 +02:00
|
|
|
}
|
2023-08-23 22:08:04 +02:00
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_src1 offset:offs_src1 atIndex:1];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:2];
|
|
|
|
[encoder setBytes:&ne00 length:sizeof(ne00) atIndex:3];
|
|
|
|
[encoder setBytes:&ne02 length:sizeof(ne02) atIndex:4];
|
|
|
|
[encoder setBytes:&nb01 length:sizeof(nb01) atIndex:5];
|
|
|
|
[encoder setBytes:&nb02 length:sizeof(nb02) atIndex:6];
|
|
|
|
[encoder setBytes:&ne12 length:sizeof(ne12) atIndex:7];
|
|
|
|
[encoder setBytes:&ne0 length:sizeof(ne0) atIndex:8];
|
|
|
|
[encoder setBytes:&ne1 length:sizeof(ne1) atIndex:9];
|
|
|
|
[encoder setBytes:&gqa length:sizeof(gqa) atIndex:10];
|
|
|
|
[encoder setThreadgroupMemoryLength:8192 atIndex:0];
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake( (ne11+31)/32, (ne01+63) / 64, ne12) threadsPerThreadgroup:MTLSizeMake(128, 1, 1)];
|
|
|
|
} else {
|
2023-06-15 19:29:48 +02:00
|
|
|
int nth0 = 32;
|
|
|
|
int nth1 = 1;
|
|
|
|
|
|
|
|
// use custom matrix x vector kernel
|
|
|
|
switch (src0t) {
|
|
|
|
case GGML_TYPE_F16:
|
|
|
|
{
|
|
|
|
nth0 = 64;
|
|
|
|
nth1 = 1;
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul_mat_f16_f32];
|
|
|
|
} break;
|
|
|
|
case GGML_TYPE_Q4_0:
|
|
|
|
{
|
|
|
|
GGML_ASSERT(ne02 == 1);
|
|
|
|
GGML_ASSERT(ne12 == 1);
|
|
|
|
|
|
|
|
nth0 = 8;
|
|
|
|
nth1 = 8;
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul_mat_q4_0_f32];
|
|
|
|
} break;
|
|
|
|
case GGML_TYPE_Q4_1:
|
|
|
|
{
|
|
|
|
GGML_ASSERT(ne02 == 1);
|
|
|
|
GGML_ASSERT(ne12 == 1);
|
|
|
|
|
|
|
|
nth0 = 8;
|
|
|
|
nth1 = 8;
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul_mat_q4_1_f32];
|
|
|
|
} break;
|
2023-08-24 15:19:57 +02:00
|
|
|
case GGML_TYPE_Q8_0:
|
|
|
|
{
|
|
|
|
GGML_ASSERT(ne02 == 1);
|
|
|
|
GGML_ASSERT(ne12 == 1);
|
|
|
|
|
|
|
|
nth0 = 8;
|
|
|
|
nth1 = 8;
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul_mat_q8_0_f32];
|
|
|
|
} break;
|
2023-06-15 19:29:48 +02:00
|
|
|
case GGML_TYPE_Q2_K:
|
|
|
|
{
|
|
|
|
GGML_ASSERT(ne02 == 1);
|
|
|
|
GGML_ASSERT(ne12 == 1);
|
|
|
|
|
2023-07-21 09:44:40 +02:00
|
|
|
nth0 = 2;
|
|
|
|
nth1 = 32;
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul_mat_q2_K_f32];
|
2023-06-15 19:29:48 +02:00
|
|
|
} break;
|
|
|
|
case GGML_TYPE_Q3_K:
|
|
|
|
{
|
|
|
|
GGML_ASSERT(ne02 == 1);
|
|
|
|
GGML_ASSERT(ne12 == 1);
|
|
|
|
|
2023-07-21 16:05:30 +02:00
|
|
|
nth0 = 2;
|
|
|
|
nth1 = 32;
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul_mat_q3_K_f32];
|
2023-06-15 19:29:48 +02:00
|
|
|
} break;
|
|
|
|
case GGML_TYPE_Q4_K:
|
|
|
|
{
|
|
|
|
GGML_ASSERT(ne02 == 1);
|
|
|
|
GGML_ASSERT(ne12 == 1);
|
|
|
|
|
2023-07-20 14:18:43 +02:00
|
|
|
nth0 = 2;
|
|
|
|
nth1 = 32;
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul_mat_q4_K_f32];
|
2023-06-15 19:29:48 +02:00
|
|
|
} break;
|
|
|
|
case GGML_TYPE_Q5_K:
|
|
|
|
{
|
|
|
|
GGML_ASSERT(ne02 == 1);
|
|
|
|
GGML_ASSERT(ne12 == 1);
|
|
|
|
|
2023-07-20 17:19:45 +02:00
|
|
|
nth0 = 2;
|
|
|
|
nth1 = 32;
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul_mat_q5_K_f32];
|
2023-06-15 19:29:48 +02:00
|
|
|
} break;
|
|
|
|
case GGML_TYPE_Q6_K:
|
|
|
|
{
|
|
|
|
GGML_ASSERT(ne02 == 1);
|
|
|
|
GGML_ASSERT(ne12 == 1);
|
|
|
|
|
2023-07-20 17:19:45 +02:00
|
|
|
nth0 = 2;
|
|
|
|
nth1 = 32;
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
[encoder setComputePipelineState:ctx->pipeline_mul_mat_q6_K_f32];
|
2023-06-15 19:29:48 +02:00
|
|
|
} break;
|
|
|
|
default:
|
|
|
|
{
|
|
|
|
fprintf(stderr, "Asserting on type %d\n",(int)src0t);
|
|
|
|
GGML_ASSERT(false && "not implemented");
|
|
|
|
}
|
|
|
|
};
|
|
|
|
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_src1 offset:offs_src1 atIndex:1];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:2];
|
|
|
|
[encoder setBytes:&ne00 length:sizeof(ne00) atIndex:3];
|
|
|
|
[encoder setBytes:&ne01 length:sizeof(ne01) atIndex:4];
|
2023-08-01 09:43:12 +02:00
|
|
|
[encoder setBytes:&ne02 length:sizeof(ne02) atIndex:5];
|
|
|
|
[encoder setBytes:&nb00 length:sizeof(nb00) atIndex:6];
|
|
|
|
[encoder setBytes:&nb01 length:sizeof(nb01) atIndex:7];
|
|
|
|
[encoder setBytes:&nb02 length:sizeof(nb02) atIndex:8];
|
|
|
|
[encoder setBytes:&ne10 length:sizeof(ne10) atIndex:9];
|
|
|
|
[encoder setBytes:&ne11 length:sizeof(ne11) atIndex:10];
|
|
|
|
[encoder setBytes:&ne12 length:sizeof(ne12) atIndex:11];
|
|
|
|
[encoder setBytes:&nb10 length:sizeof(nb10) atIndex:12];
|
|
|
|
[encoder setBytes:&nb11 length:sizeof(nb11) atIndex:13];
|
|
|
|
[encoder setBytes:&nb12 length:sizeof(nb12) atIndex:14];
|
|
|
|
[encoder setBytes:&ne0 length:sizeof(ne0) atIndex:15];
|
|
|
|
[encoder setBytes:&ne1 length:sizeof(ne1) atIndex:16];
|
2023-08-23 22:08:04 +02:00
|
|
|
[encoder setBytes:&gqa length:sizeof(gqa) atIndex:17];
|
2023-06-15 19:29:48 +02:00
|
|
|
|
2023-08-24 15:19:57 +02:00
|
|
|
if (src0t == GGML_TYPE_Q4_0 || src0t == GGML_TYPE_Q4_1 || src0t == GGML_TYPE_Q8_0 ||
|
2023-07-21 09:44:40 +02:00
|
|
|
src0t == GGML_TYPE_Q2_K || src0t == GGML_TYPE_Q4_K) {
|
2023-08-23 22:08:04 +02:00
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake((ne01 + 7)/8, ne11, ne12) threadsPerThreadgroup:MTLSizeMake(nth0, nth1, 1)];
|
2023-06-15 19:29:48 +02:00
|
|
|
}
|
2023-07-21 16:05:30 +02:00
|
|
|
else if (src0t == GGML_TYPE_Q3_K) {
|
|
|
|
#ifdef GGML_QKK_64
|
2023-08-23 22:08:04 +02:00
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake((ne01 + 1)/2, ne11, ne12) threadsPerThreadgroup:MTLSizeMake(nth0, nth1, 1)];
|
2023-07-21 16:05:30 +02:00
|
|
|
#else
|
2023-08-23 22:08:04 +02:00
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake((ne01 + 3)/4, ne11, ne12) threadsPerThreadgroup:MTLSizeMake(nth0, nth1, 1)];
|
2023-07-21 16:05:30 +02:00
|
|
|
#endif
|
|
|
|
}
|
2023-07-20 17:19:45 +02:00
|
|
|
else if (src0t == GGML_TYPE_Q5_K) {
|
2023-08-23 22:08:04 +02:00
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake((ne01 + 3)/4, ne11, ne12) threadsPerThreadgroup:MTLSizeMake(nth0, nth1, 1)];
|
2023-07-20 17:19:45 +02:00
|
|
|
}
|
|
|
|
else if (src0t == GGML_TYPE_Q6_K) {
|
2023-08-23 22:08:04 +02:00
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake((ne01 + 1)/2, ne11, ne12) threadsPerThreadgroup:MTLSizeMake(nth0, nth1, 1)];
|
2023-06-15 19:29:48 +02:00
|
|
|
} else {
|
|
|
|
[encoder setThreadgroupMemoryLength:nth0*sizeof(float) atIndex:0];
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(ne01, ne11, ne12) threadsPerThreadgroup:MTLSizeMake(nth0, nth1, 1)];
|
|
|
|
}
|
|
|
|
}
|
|
|
|
} break;
|
|
|
|
case GGML_OP_GET_ROWS:
|
|
|
|
{
|
|
|
|
switch (src0->type) {
|
2023-08-24 15:19:57 +02:00
|
|
|
case GGML_TYPE_F16: [encoder setComputePipelineState:ctx->pipeline_get_rows_f16]; break;
|
2023-06-15 19:29:48 +02:00
|
|
|
case GGML_TYPE_Q4_0: [encoder setComputePipelineState:ctx->pipeline_get_rows_q4_0]; break;
|
|
|
|
case GGML_TYPE_Q4_1: [encoder setComputePipelineState:ctx->pipeline_get_rows_q4_1]; break;
|
2023-08-24 15:19:57 +02:00
|
|
|
case GGML_TYPE_Q8_0: [encoder setComputePipelineState:ctx->pipeline_get_rows_q8_0]; break;
|
k-quants : support for super-block size of 64 (#2001)
* k_quants: WIP super-blocks with 64 weights
* k_quants: WIP super-blocks with 64 weights
Q6_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q4_K scalar and AVX2 works
* k_quants: WIP super-blocks with 64 weights
Q2_K scalar and AVX2 works. Q2_K is way too slow (it is actually slower
than the scalar implementation)
* k_quants: WIP super-blocks with 64 weights
Q3_K scalar and AVX2 works.
* k_quants: WIP super-blocks with 64 weights
Q5_K scalar and AVX2 works, and with that all
k_quants are done on AVX2 and scalar
* k_quants: WIP super-blocks with 64 weights
Q6_K working on CUDA. Cannot make it run quite as gast as
with super-blocks with 256 weigths: 8% slower on 4080,
20% slower on the 1660 (but there we fit 1 less layer on the
GPU because pf the larger model size), so some fraction of
these 20% is due to that,
* k_quants: WIP super-blocks with 64 weights
Q4_K working on CUDA. ~10% slower on GTX-1660,
16% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q2_K working on CUDA. ~3% slower on GTX-1660,
10% slower on 4080.
* k_quants: WIP super-blocks with 64 weights
Q3_K working on CUDA.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on CUDA, and with this CUDA is done.
* k_quants: WIP super-blocks with 64 weights
Q6_K working on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Q4_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q2_K working on ARM_NEON, but quite a bit slower than 256 weights
* k_quants: WIP super-blocks with 64 weights
Q3_K working on ARM_NEON, but quite a bit slower than 256 weights.
* k_quants: WIP super-blocks with 64 weights
Q5_K working on ARM_NEON, but quite a bit slower than 256 weights.
With that, we have full support for ARM_NEON, although
performance is not quite there.
* k_quants: WIP super-blocks with 64 weights
Slightly more efficient Q3_K and Q5_K
* k_quants: WIP super-blocks with 64 weights
Another small improvement for Q3_K and Q5_K on ARM_NEON
* k_quants: WIP super-blocks with 64 weights
Yet another speedup for Q5_K on ARM_NEON.
We are now within 10% of the QK_K = 256 version.
* k_quants: WIP super-blocks with 64 weights
* We are able to pass preprocessor macros to the Metal
compiler
* Q6_K works and is actually slightly more efficient than
the QK_K = 256 version (25.2 ms vs 25.8 ms)
* k_quants: WIP super-blocks with 64 weights
Q4_K works on Metal and is actually slightly faster
than QK_K = 256 (21.95 ms vs 24.0 ms).
* k_quants: WIP super-blocks with 64 weights
Q2_K works on Metal and is very slightly faster
than QK_K = 256 (23.8 ms vs 24.2 ms).
* k_quants: WIP super-blocks with 64 weights
Q3_K works on Metal and is slightly faster
than QK_K = 256 (26.6 ms vs 28.3 ms).
* k_quants: WIP super-blocks with 64 weights
Q5_K works on Metal and is slightly faster
than QK_K = 256 (23.7 ms vs 26.3 ms).
* k_quants: call them _K, not _k, also on Metal
* k_quants: correctly define QK_K in llama.cpp
* Fixed bug in q4_K quantization added with the 64-block addition
* Simplify via lambda
* k_quants: swicth Q3_K to 4-bit scales when QK_K = 64
Otherwise there isn't much benefit from this
quantization type. There is some very slight loss
in accuracy, but we reduce size by ~7%.
E.g., for OpenLLaMA-3B, Q3_K_S perplexity is
8.6131 with 8-bit scales and 8.6352 with 4-bit,
while file size decreases from 1.53G to 1.44G.
* k_quants: switch Q4_K to 4-bit scales when QK_K = 64
Here the loss in accuracy is greater than for Q3_K,
but the Q4_K points still move further to the left on
the perplexity vs size curve.
* k_quants: forgot to add the Metal changes in last commit
* k_quants: change Q5_K to be type 0 when QK_K = 64
Still needs AVX2 implementation
* k_quants: AVX2 implementation for new 64-weight Q5_K
* k_quants: 10% faster ARM_NEON Q5_K dot product
* k_quants: fixed issue caused by merging with master
---------
Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
2023-06-26 18:43:07 +02:00
|
|
|
case GGML_TYPE_Q2_K: [encoder setComputePipelineState:ctx->pipeline_get_rows_q2_K]; break;
|
|
|
|
case GGML_TYPE_Q3_K: [encoder setComputePipelineState:ctx->pipeline_get_rows_q3_K]; break;
|
|
|
|
case GGML_TYPE_Q4_K: [encoder setComputePipelineState:ctx->pipeline_get_rows_q4_K]; break;
|
|
|
|
case GGML_TYPE_Q5_K: [encoder setComputePipelineState:ctx->pipeline_get_rows_q5_K]; break;
|
|
|
|
case GGML_TYPE_Q6_K: [encoder setComputePipelineState:ctx->pipeline_get_rows_q6_K]; break;
|
2023-06-15 19:29:48 +02:00
|
|
|
default: GGML_ASSERT(false && "not implemented");
|
|
|
|
}
|
|
|
|
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_src1 offset:offs_src1 atIndex:1];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:2];
|
|
|
|
[encoder setBytes:&(src0->ne[0]) length:sizeof( int64_t) atIndex:3];
|
|
|
|
[encoder setBytes:&(src0->nb[1]) length:sizeof(uint64_t) atIndex:4];
|
|
|
|
[encoder setBytes:&(dst->nb[1]) length:sizeof(uint64_t) atIndex:5];
|
|
|
|
|
|
|
|
const int64_t n = ggml_nelements(src1);
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
|
|
|
|
} break;
|
|
|
|
case GGML_OP_RMS_NORM:
|
|
|
|
{
|
2023-07-24 17:57:12 +02:00
|
|
|
float eps;
|
|
|
|
memcpy(&eps, dst->op_params, sizeof(float));
|
2023-06-15 19:29:48 +02:00
|
|
|
|
2023-07-20 12:32:22 +02:00
|
|
|
const int nth = 512;
|
2023-06-15 19:29:48 +02:00
|
|
|
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_rms_norm];
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
|
|
|
[encoder setBytes:&ne00 length:sizeof( int64_t) atIndex:2];
|
|
|
|
[encoder setBytes:&nb01 length:sizeof(uint64_t) atIndex:3];
|
|
|
|
[encoder setBytes:&eps length:sizeof( float) atIndex:4];
|
2023-07-20 12:32:22 +02:00
|
|
|
[encoder setThreadgroupMemoryLength:nth/32*sizeof(float) atIndex:0];
|
2023-06-15 19:29:48 +02:00
|
|
|
|
|
|
|
const int64_t nrows = ggml_nrows(src0);
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(nrows, 1, 1) threadsPerThreadgroup:MTLSizeMake(nth, 1, 1)];
|
|
|
|
} break;
|
2023-06-17 16:37:49 +02:00
|
|
|
case GGML_OP_NORM:
|
|
|
|
{
|
2023-08-23 22:08:04 +02:00
|
|
|
float eps;
|
|
|
|
memcpy(&eps, dst->op_params, sizeof(float));
|
2023-06-17 16:37:49 +02:00
|
|
|
|
|
|
|
const int nth = 256;
|
|
|
|
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_norm];
|
2023-08-23 22:08:04 +02:00
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
|
|
|
[encoder setBytes:&ne00 length:sizeof( int64_t) atIndex:2];
|
|
|
|
[encoder setBytes:&nb01 length:sizeof(uint64_t) atIndex:3];
|
|
|
|
[encoder setBytes:&eps length:sizeof( float) atIndex:4];
|
2023-06-17 16:37:49 +02:00
|
|
|
[encoder setThreadgroupMemoryLength:nth*sizeof(float) atIndex:0];
|
|
|
|
|
|
|
|
const int64_t nrows = ggml_nrows(src0);
|
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(nrows, 1, 1) threadsPerThreadgroup:MTLSizeMake(nth, 1, 1)];
|
|
|
|
} break;
|
|
|
|
case GGML_OP_ALIBI:
|
|
|
|
{
|
|
|
|
GGML_ASSERT((src0t == GGML_TYPE_F32));
|
2023-06-17 19:24:11 +02:00
|
|
|
|
2023-07-23 14:36:02 +02:00
|
|
|
const int n_past = ((int32_t *) dst->op_params)[0]; UNUSED(n_past);
|
|
|
|
const int n_head = ((int32_t *) dst->op_params)[1];
|
|
|
|
float max_bias;
|
|
|
|
memcpy(&max_bias, (int32_t *) dst->op_params + 2, sizeof(float));
|
2023-06-17 19:24:11 +02:00
|
|
|
|
2023-06-17 16:37:49 +02:00
|
|
|
if (__builtin_popcount(n_head) != 1) {
|
|
|
|
GGML_ASSERT(false && "only power-of-two n_head implemented");
|
|
|
|
}
|
2023-06-17 19:24:11 +02:00
|
|
|
|
2023-06-17 16:37:49 +02:00
|
|
|
const int n_heads_log2_floor = 1 << (int) floor(log2(n_head));
|
|
|
|
const float m0 = powf(2.0f, -(max_bias) / n_heads_log2_floor);
|
2023-06-17 19:24:11 +02:00
|
|
|
|
2023-06-17 16:37:49 +02:00
|
|
|
[encoder setComputePipelineState:ctx->pipeline_alibi_f32];
|
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
|
|
|
[encoder setBytes:&ne00 length:sizeof( int64_t) atIndex:2];
|
|
|
|
[encoder setBytes:&ne01 length:sizeof( int64_t) atIndex:3];
|
|
|
|
[encoder setBytes:&ne02 length:sizeof( int64_t) atIndex:4];
|
|
|
|
[encoder setBytes:&ne03 length:sizeof( int64_t) atIndex:5];
|
|
|
|
[encoder setBytes:&nb00 length:sizeof(uint64_t) atIndex:6];
|
|
|
|
[encoder setBytes:&nb01 length:sizeof(uint64_t) atIndex:7];
|
|
|
|
[encoder setBytes:&nb02 length:sizeof(uint64_t) atIndex:8];
|
|
|
|
[encoder setBytes:&nb03 length:sizeof(uint64_t) atIndex:9];
|
|
|
|
[encoder setBytes:&ne0 length:sizeof( int64_t) atIndex:10];
|
|
|
|
[encoder setBytes:&ne1 length:sizeof( int64_t) atIndex:11];
|
|
|
|
[encoder setBytes:&ne2 length:sizeof( int64_t) atIndex:12];
|
|
|
|
[encoder setBytes:&ne3 length:sizeof( int64_t) atIndex:13];
|
|
|
|
[encoder setBytes:&nb0 length:sizeof(uint64_t) atIndex:14];
|
|
|
|
[encoder setBytes:&nb1 length:sizeof(uint64_t) atIndex:15];
|
|
|
|
[encoder setBytes:&nb2 length:sizeof(uint64_t) atIndex:16];
|
|
|
|
[encoder setBytes:&nb3 length:sizeof(uint64_t) atIndex:17];
|
|
|
|
[encoder setBytes:&m0 length:sizeof( float) atIndex:18];
|
2023-08-23 22:08:04 +02:00
|
|
|
|
2023-06-17 16:37:49 +02:00
|
|
|
const int nth = 32;
|
2023-08-23 22:08:04 +02:00
|
|
|
|
2023-06-17 16:37:49 +02:00
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(ne01, ne02, ne03) threadsPerThreadgroup:MTLSizeMake(nth, 1, 1)];
|
|
|
|
} break;
|
2023-06-15 19:29:48 +02:00
|
|
|
case GGML_OP_ROPE:
|
|
|
|
{
|
2023-07-23 14:36:02 +02:00
|
|
|
const int n_past = ((int32_t *) dst->op_params)[0];
|
|
|
|
const int n_dims = ((int32_t *) dst->op_params)[1];
|
|
|
|
const int mode = ((int32_t *) dst->op_params)[2];
|
2023-06-15 19:29:48 +02:00
|
|
|
|
llama : add custom RoPE (#2054)
* Implement customizable RoPE
The original RoPE has pre-defined parameters
theta_i = 10000^(−2(i−1)/d), for i in [1, 2, ..., d/2]
Our customizable RoPE, ggml_rope_custom_inplace, uses
theta_i = scale * base^(−2(i−1)/d), for i in [1, 2, ..., d/2]
with the default matches the original
scale = 1.0
base = 10000
The new command line arguments
--rope-freq-base
--rope-freq-scale
set the two new RoPE parameter.
Recent researches show changing these two parameters extends the context limit with minimal loss.
1. Extending Context to 8K
kaiokendev
https://kaiokendev.github.io/til#extending-context-to-8k
2. Extending Context Window of Large Language Models via Positional Interpolation
Shouyuan Chen, Sherman Wong, Liangjian Chen, Yuandong Tian
https://arxiv.org/abs/2306.15595
3. NTK-Aware Scaled RoPE allows LLaMA models to have extended (8k+) context size without any fine-tuning and minimal perplexity degradation.
https://www.reddit.com/user/bloc97
https://www.reddit.com/r/LocalLLaMA/comments/14lz7j5/ntkaware_scaled_rope_allows_llama_models_to_have/
For the bold, try adding the following command line parameters to your favorite model:
-c 16384 --rope-freq-base 80000 --rope-freq-scale 0.5
* ggml-metal: fix custom rope
* common: fix argument names in help
* llama: increase MEM_REQ_EVAL for MODEL_3B
It avoids crashing for quantized weights on CPU.
Better ways to calculate the required buffer size would be better.
* llama: make MEM_REQ_EVAL depend on n_ctx
* server: use proper Content-Type in curl examples
Without the header Content-Type: application/json, curl will POST with
Content-Type: application/x-www-form-urlencoded
Though our simple server doesn't care, the httplib.h used has a limit
with CPPHTTPLIB_FORM_URL_ENCODED_PAYLOAD_MAX_LENGTH 8192
With Content-Type: application/json, we can send large json data.
* style : minor fixes, mostly indentations
* ggml : fix asserts
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2023-07-15 12:34:16 +02:00
|
|
|
float freq_base;
|
|
|
|
float freq_scale;
|
2023-07-23 14:36:02 +02:00
|
|
|
memcpy(&freq_base, (int32_t *) dst->op_params + 4, sizeof(float));
|
|
|
|
memcpy(&freq_scale, (int32_t *) dst->op_params + 5, sizeof(float));
|
llama : add custom RoPE (#2054)
* Implement customizable RoPE
The original RoPE has pre-defined parameters
theta_i = 10000^(−2(i−1)/d), for i in [1, 2, ..., d/2]
Our customizable RoPE, ggml_rope_custom_inplace, uses
theta_i = scale * base^(−2(i−1)/d), for i in [1, 2, ..., d/2]
with the default matches the original
scale = 1.0
base = 10000
The new command line arguments
--rope-freq-base
--rope-freq-scale
set the two new RoPE parameter.
Recent researches show changing these two parameters extends the context limit with minimal loss.
1. Extending Context to 8K
kaiokendev
https://kaiokendev.github.io/til#extending-context-to-8k
2. Extending Context Window of Large Language Models via Positional Interpolation
Shouyuan Chen, Sherman Wong, Liangjian Chen, Yuandong Tian
https://arxiv.org/abs/2306.15595
3. NTK-Aware Scaled RoPE allows LLaMA models to have extended (8k+) context size without any fine-tuning and minimal perplexity degradation.
https://www.reddit.com/user/bloc97
https://www.reddit.com/r/LocalLLaMA/comments/14lz7j5/ntkaware_scaled_rope_allows_llama_models_to_have/
For the bold, try adding the following command line parameters to your favorite model:
-c 16384 --rope-freq-base 80000 --rope-freq-scale 0.5
* ggml-metal: fix custom rope
* common: fix argument names in help
* llama: increase MEM_REQ_EVAL for MODEL_3B
It avoids crashing for quantized weights on CPU.
Better ways to calculate the required buffer size would be better.
* llama: make MEM_REQ_EVAL depend on n_ctx
* server: use proper Content-Type in curl examples
Without the header Content-Type: application/json, curl will POST with
Content-Type: application/x-www-form-urlencoded
Though our simple server doesn't care, the httplib.h used has a limit
with CPPHTTPLIB_FORM_URL_ENCODED_PAYLOAD_MAX_LENGTH 8192
With Content-Type: application/json, we can send large json data.
* style : minor fixes, mostly indentations
* ggml : fix asserts
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2023-07-15 12:34:16 +02:00
|
|
|
|
2023-06-15 19:29:48 +02:00
|
|
|
[encoder setComputePipelineState:ctx->pipeline_rope];
|
2023-08-23 22:08:04 +02:00
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
llama : add custom RoPE (#2054)
* Implement customizable RoPE
The original RoPE has pre-defined parameters
theta_i = 10000^(−2(i−1)/d), for i in [1, 2, ..., d/2]
Our customizable RoPE, ggml_rope_custom_inplace, uses
theta_i = scale * base^(−2(i−1)/d), for i in [1, 2, ..., d/2]
with the default matches the original
scale = 1.0
base = 10000
The new command line arguments
--rope-freq-base
--rope-freq-scale
set the two new RoPE parameter.
Recent researches show changing these two parameters extends the context limit with minimal loss.
1. Extending Context to 8K
kaiokendev
https://kaiokendev.github.io/til#extending-context-to-8k
2. Extending Context Window of Large Language Models via Positional Interpolation
Shouyuan Chen, Sherman Wong, Liangjian Chen, Yuandong Tian
https://arxiv.org/abs/2306.15595
3. NTK-Aware Scaled RoPE allows LLaMA models to have extended (8k+) context size without any fine-tuning and minimal perplexity degradation.
https://www.reddit.com/user/bloc97
https://www.reddit.com/r/LocalLLaMA/comments/14lz7j5/ntkaware_scaled_rope_allows_llama_models_to_have/
For the bold, try adding the following command line parameters to your favorite model:
-c 16384 --rope-freq-base 80000 --rope-freq-scale 0.5
* ggml-metal: fix custom rope
* common: fix argument names in help
* llama: increase MEM_REQ_EVAL for MODEL_3B
It avoids crashing for quantized weights on CPU.
Better ways to calculate the required buffer size would be better.
* llama: make MEM_REQ_EVAL depend on n_ctx
* server: use proper Content-Type in curl examples
Without the header Content-Type: application/json, curl will POST with
Content-Type: application/x-www-form-urlencoded
Though our simple server doesn't care, the httplib.h used has a limit
with CPPHTTPLIB_FORM_URL_ENCODED_PAYLOAD_MAX_LENGTH 8192
With Content-Type: application/json, we can send large json data.
* style : minor fixes, mostly indentations
* ggml : fix asserts
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2023-07-15 12:34:16 +02:00
|
|
|
[encoder setBytes:&ne00 length:sizeof( int64_t) atIndex:2];
|
|
|
|
[encoder setBytes:&ne01 length:sizeof( int64_t) atIndex:3];
|
|
|
|
[encoder setBytes:&ne02 length:sizeof( int64_t) atIndex:4];
|
|
|
|
[encoder setBytes:&ne03 length:sizeof( int64_t) atIndex:5];
|
|
|
|
[encoder setBytes:&nb00 length:sizeof(uint64_t) atIndex:6];
|
|
|
|
[encoder setBytes:&nb01 length:sizeof(uint64_t) atIndex:7];
|
|
|
|
[encoder setBytes:&nb02 length:sizeof(uint64_t) atIndex:8];
|
|
|
|
[encoder setBytes:&nb03 length:sizeof(uint64_t) atIndex:9];
|
|
|
|
[encoder setBytes:&ne0 length:sizeof( int64_t) atIndex:10];
|
|
|
|
[encoder setBytes:&ne1 length:sizeof( int64_t) atIndex:11];
|
|
|
|
[encoder setBytes:&ne2 length:sizeof( int64_t) atIndex:12];
|
|
|
|
[encoder setBytes:&ne3 length:sizeof( int64_t) atIndex:13];
|
|
|
|
[encoder setBytes:&nb0 length:sizeof(uint64_t) atIndex:14];
|
|
|
|
[encoder setBytes:&nb1 length:sizeof(uint64_t) atIndex:15];
|
|
|
|
[encoder setBytes:&nb2 length:sizeof(uint64_t) atIndex:16];
|
|
|
|
[encoder setBytes:&nb3 length:sizeof(uint64_t) atIndex:17];
|
|
|
|
[encoder setBytes:&n_past length:sizeof( int) atIndex:18];
|
|
|
|
[encoder setBytes:&n_dims length:sizeof( int) atIndex:19];
|
|
|
|
[encoder setBytes:&mode length:sizeof( int) atIndex:20];
|
|
|
|
[encoder setBytes:&freq_base length:sizeof(float) atIndex:21];
|
|
|
|
[encoder setBytes:&freq_scale length:sizeof(float) atIndex:22];
|
2023-06-15 19:29:48 +02:00
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(ne01, ne02, ne03) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
|
|
|
|
} break;
|
2023-07-23 13:00:37 +02:00
|
|
|
case GGML_OP_DUP:
|
2023-06-15 19:29:48 +02:00
|
|
|
case GGML_OP_CPY:
|
2023-07-23 13:00:37 +02:00
|
|
|
case GGML_OP_CONT:
|
2023-06-15 19:29:48 +02:00
|
|
|
{
|
|
|
|
const int nth = 32;
|
|
|
|
|
|
|
|
switch (src0t) {
|
|
|
|
case GGML_TYPE_F32:
|
|
|
|
{
|
|
|
|
switch (dstt) {
|
|
|
|
case GGML_TYPE_F16: [encoder setComputePipelineState:ctx->pipeline_cpy_f32_f16]; break;
|
|
|
|
case GGML_TYPE_F32: [encoder setComputePipelineState:ctx->pipeline_cpy_f32_f32]; break;
|
|
|
|
default: GGML_ASSERT(false && "not implemented");
|
|
|
|
};
|
|
|
|
} break;
|
2023-06-17 16:37:49 +02:00
|
|
|
case GGML_TYPE_F16:
|
|
|
|
{
|
|
|
|
switch (dstt) {
|
|
|
|
case GGML_TYPE_F16: [encoder setComputePipelineState:ctx->pipeline_cpy_f16_f16]; break;
|
|
|
|
case GGML_TYPE_F32: GGML_ASSERT(false && "cpy_f16_f32 not implemented"); break;
|
|
|
|
default: GGML_ASSERT(false && "not implemented");
|
|
|
|
};
|
|
|
|
} break;
|
2023-06-15 19:29:48 +02:00
|
|
|
default: GGML_ASSERT(false && "not implemented");
|
|
|
|
}
|
|
|
|
|
2023-08-23 22:08:04 +02:00
|
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
|
|
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
|
|
|
|
[encoder setBytes:&ne00 length:sizeof( int64_t) atIndex:2];
|
|
|
|
[encoder setBytes:&ne01 length:sizeof( int64_t) atIndex:3];
|
|
|
|
[encoder setBytes:&ne02 length:sizeof( int64_t) atIndex:4];
|
|
|
|
[encoder setBytes:&ne03 length:sizeof( int64_t) atIndex:5];
|
|
|
|
[encoder setBytes:&nb00 length:sizeof(uint64_t) atIndex:6];
|
|
|
|
[encoder setBytes:&nb01 length:sizeof(uint64_t) atIndex:7];
|
|
|
|
[encoder setBytes:&nb02 length:sizeof(uint64_t) atIndex:8];
|
|
|
|
[encoder setBytes:&nb03 length:sizeof(uint64_t) atIndex:9];
|
|
|
|
[encoder setBytes:&ne0 length:sizeof( int64_t) atIndex:10];
|
|
|
|
[encoder setBytes:&ne1 length:sizeof( int64_t) atIndex:11];
|
|
|
|
[encoder setBytes:&ne2 length:sizeof( int64_t) atIndex:12];
|
|
|
|
[encoder setBytes:&ne3 length:sizeof( int64_t) atIndex:13];
|
|
|
|
[encoder setBytes:&nb0 length:sizeof(uint64_t) atIndex:14];
|
|
|
|
[encoder setBytes:&nb1 length:sizeof(uint64_t) atIndex:15];
|
|
|
|
[encoder setBytes:&nb2 length:sizeof(uint64_t) atIndex:16];
|
|
|
|
[encoder setBytes:&nb3 length:sizeof(uint64_t) atIndex:17];
|
2023-06-15 19:29:48 +02:00
|
|
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(ne01, ne02, ne03) threadsPerThreadgroup:MTLSizeMake(nth, 1, 1)];
|
|
|
|
} break;
|
|
|
|
default:
|
2023-07-24 13:46:21 +02:00
|
|
|
{
|
|
|
|
fprintf(stderr, "%s: node %3d, op = %8s not implemented\n", __func__, i, ggml_op_name(dst->op));
|
|
|
|
GGML_ASSERT(false);
|
|
|
|
}
|
2023-06-15 19:29:48 +02:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
if (encoder != nil) {
|
|
|
|
[encoder endEncoding];
|
|
|
|
encoder = nil;
|
|
|
|
}
|
|
|
|
|
|
|
|
[command_buffer commit];
|
|
|
|
});
|
2023-06-04 22:34:30 +02:00
|
|
|
}
|
2023-06-15 19:29:48 +02:00
|
|
|
|
|
|
|
// wait for all threads to finish
|
|
|
|
dispatch_barrier_sync(queue, ^{});
|
|
|
|
|
|
|
|
[command_buffers[n_cb - 1] waitUntilCompleted];
|
2023-06-18 08:09:47 +02:00
|
|
|
|
|
|
|
// check status of command buffers
|
|
|
|
// needed to detect if the device ran out-of-memory for example (#1881)
|
|
|
|
for (int i = 0; i < n_cb; i++) {
|
|
|
|
MTLCommandBufferStatus status = (MTLCommandBufferStatus) [command_buffers[i] status];
|
|
|
|
if (status != MTLCommandBufferStatusCompleted) {
|
|
|
|
fprintf(stderr, "%s: command buffer %d failed with status %lu\n", __func__, i, status);
|
|
|
|
GGML_ASSERT(false);
|
|
|
|
}
|
|
|
|
}
|
2023-06-04 22:34:30 +02:00
|
|
|
}
|