feat: add Metal INT8 ConvRot support (#2297)

src/core/metal/CMakeLists.txt

@@ -0,0 +1,8 @@
+set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS int8.metal int8_mpp.metal)
+file(READ "${CMAKE_CURRENT_SOURCE_DIR}/int8.metal" SD_METAL_I8_SOURCE)
+file(READ "${CMAKE_CURRENT_SOURCE_DIR}/int8_mpp.metal" SD_METAL_I8_MPP_SOURCE)
+configure_file(int8_source.h.in "${CMAKE_CURRENT_BINARY_DIR}/int8_source.h" @ONLY)
+target_sources(${SD_LIB} PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/int8.mm")
+target_compile_options(${SD_LIB} PRIVATE "$<$<COMPILE_LANGUAGE:OBJCXX>:-fno-objc-arc>")
+target_include_directories(${SD_LIB} PRIVATE "${CMAKE_CURRENT_BINARY_DIR}")
+set_target_properties(${SD_LIB} PROPERTIES OBJCXX_STANDARD 17 OBJCXX_STANDARD_REQUIRED ON)

src/core/metal/int8.metal

@@ -0,0 +1,143 @@
+#include <metal_simdgroup_matrix>
+#include <metal_stdlib>
+using namespace metal;
+
+struct QuantParams {
+    uint k;
+    uint rows;
+    uint group;
+    uint unused;
+    ulong scale_offset;
+};
+
+kernel void sd_i8_rotate(device const float* input [[buffer(0)]],
+                         device float* output [[buffer(1)]],
+                         constant QuantParams& p [[buffer(2)]],
+                         uint block [[threadgroup_position_in_grid]],
+                         uint tid [[thread_index_in_threadgroup]]) {
+    threadgroup float values[256];
+    const ulong base = ulong(block) * p.group;
+    values[tid]      = input[base + tid] * rsqrt(float(p.group));
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    for (uint stride = 1; stride < p.group; stride *= 4) {
+        uint index = tid / (4 * stride) * (4 * stride) + tid % stride;
+        float a = values[index], b = values[index + stride];
+        float c = values[index + 2 * stride], d = values[index + 3 * stride];
+        float value;
+        switch ((tid / stride) % 4) {
+            case 0:
+                value = a + b + c - d;
+                break;
+            case 1:
+                value = a + b - c + d;
+                break;
+            case 2:
+                value = a - b + c + d;
+                break;
+            default:
+                value = -a + b + c + d;
+                break;
+        }
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        values[tid] = value;
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+    }
+    output[base + tid] = values[tid];
+}
+
+kernel void sd_i8_quantize(device const float* input [[buffer(0)]],
+                           device char* output [[buffer(1)]],
+                           constant QuantParams& p [[buffer(2)]],
+                           uint row [[threadgroup_position_in_grid]],
+                           uint tid [[thread_index_in_threadgroup]],
+                           uint lane [[thread_index_in_simdgroup]],
+                           uint simd [[simdgroup_index_in_threadgroup]]) {
+    threadgroup float maxima[8];
+    const ulong base = ulong(row) * p.k;
+    float amax       = 0;
+    for (uint i = tid; i < p.k; i += 256) {
+        amax = max(amax, abs(input[base + i]));
+    }
+    amax = simd_max(amax);
+    if (lane == 0)
+        maxima[simd] = amax;
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    amax          = simd_max(lane < 8 ? maxima[lane] : 0.f);
+    float scale   = amax / 127.f;
+    float inverse = scale == 0 ? 0 : 1.f / scale;
+    for (uint i = tid; i < p.k; i += 256) {
+        output[base + i] = char(clamp(rint(input[base + i] * inverse), -127.f, 127.f));
+    }
+    if (tid == 0)
+        ((device float*)(output + p.scale_offset))[row] = scale;
+}
+
+struct MatParams {
+    uint k;
+    uint n;
+    uint rows;
+    uint flags;
+};
+
+kernel void sd_i8_multiply(device const char* weight [[buffer(0)]],
+                           device const char* input [[buffer(1)]],
+                           device const float* xs [[buffer(2)]],
+                           device const float* ws [[buffer(3)]],
+                           device const float* bias [[buffer(4)]],
+                           device float* output [[buffer(5)]],
+                           constant MatParams& p [[buffer(6)]],
+                           uint2 tile [[threadgroup_position_in_grid]],
+                           uint tid [[thread_index_in_threadgroup]],
+                           uint simd [[simdgroup_index_in_threadgroup]]) {
+    threadgroup half a[32 * 32];
+    threadgroup half b[32 * 32];
+    threadgroup float partial[32 * 32];
+    uint sums[8]    = {};
+    const uint row0 = tile.y * 32, col0 = tile.x * 32;
+    const uint sr = (simd / 2) * 16, sc = (simd % 2) * 16;
+
+    // Half represents every INT8 value exactly. Each 256-term dot product fits
+    // in FP32's exact integer range; combine chunks as integers to avoid drift.
+    for (uint chunk = 0; chunk < p.k; chunk += 256) {
+        simdgroup_float8x8 acc[4];
+        for (uint i = 0; i < 4; ++i)
+            acc[i] = make_filled_simdgroup_matrix<float, 8>(0.f);
+        for (uint kk = chunk; kk < min(chunk + 256, p.k); kk += 32) {
+            for (uint index = tid; index < 1024; index += 128) {
+                uint r = index / 32, k = index % 32;
+                a[index] = row0 + r < p.rows && kk + k < p.k ? half(input[ulong(row0 + r) * p.k + kk + k]) : half(0);
+                b[index] = col0 + r < p.n && kk + k < p.k ? half(weight[ulong(col0 + r) * p.k + kk + k]) : half(0);
+            }
+            threadgroup_barrier(mem_flags::mem_threadgroup);
+            for (uint k = 0; k < 32; k += 8) {
+                simdgroup_half8x8 ma[2], mb[2];
+                for (uint i = 0; i < 2; ++i) {
+                    simdgroup_load(ma[i], a + (sr + i * 8) * 32 + k, 32);
+                    simdgroup_load(mb[i], b + (sc + i * 8) * 32 + k, 32, ulong2(0), true);
+                }
+                for (uint i = 0; i < 4; ++i) {
+                    simdgroup_multiply_accumulate(acc[i], ma[i / 2], mb[i % 2], acc[i]);
+                }
+            }
+            threadgroup_barrier(mem_flags::mem_threadgroup);
+        }
+        for (uint i = 0; i < 4; ++i) {
+            simdgroup_store(acc[i], partial + (sr + (i / 2) * 8) * 32 + sc + (i % 2) * 8, 32);
+        }
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        for (uint i = 0; i < 8; ++i)
+            sums[i] += uint(int(partial[tid + 128 * i]));
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+    }
+    for (uint i = 0; i < 8; ++i) {
+        uint index = tid + 128 * i, row = row0 + index / 32, col = col0 + index % 32;
+        if (row < p.rows && col < p.n) {
+            float value = float(as_type<int>(sums[i])) * xs[row];
+            if (p.flags & 1)
+                value *= ws[col];
+            if (p.flags & 2)
+                value += bias[col];
+            output[ulong(row) * p.n + col] = value;
+        }
+    }
+}

src/core/metal/int8.mm

@@ -0,0 +1,298 @@
+#include "core/ggml_extensions.h"
+#include "ggml-ext-metal.h"
+#include "int8_source.h"
+
+#import <Foundation/Foundation.h>
+#import <Metal/Metal.h>
+
+#include <climits>
+#include <memory>
+#include <vector>
+
+namespace sd_metal {
+
+    struct Kernels {
+        id<MTLComputePipelineState> rotate       = nil;
+        id<MTLComputePipelineState> quantize     = nil;
+        id<MTLComputePipelineState> multiply     = nil;
+        id<MTLComputePipelineState> multiply_mpp = nil;
+        ggml_ext_metal_buffer_fn buffer          = nullptr;
+        ggml_ext_metal_submit_fn submit          = nullptr;
+
+        ~Kernels() {
+            [rotate release];
+            [quantize release];
+            [multiply release];
+            [multiply_mpp release];
+        }
+    };
+
+    static id<MTLComputePipelineState> tensor_multiply(id<MTLDevice> device) {
+        if (@available(macOS 26.0, iOS 26.0, *)) {
+            MTLCompileOptions* options = [[MTLCompileOptions alloc] init];
+            // Keep building with older SDKs; the optional shader is compiled at runtime.
+            options.languageVersion = static_cast<MTLLanguageVersion>(4 << 16);
+            options.fastMathEnabled = NO;
+            NSError* error          = nil;
+            id<MTLLibrary> library  = [device newLibraryWithSource:[NSString stringWithUTF8String:sd_metal_i8_mpp_source] options:options error:&error];
+            [options release];
+            id<MTLFunction> function             = [library newFunctionWithName:@"sd_i8_multiply_mpp"];
+            id<MTLComputePipelineState> pipeline = function ? [device newComputePipelineStateWithFunction:function error:&error] : nil;
+            [function release];
+            [library release];
+            if (pipeline && pipeline.threadExecutionWidth == 32 && pipeline.maxTotalThreadsPerThreadgroup >= 128) {
+                return pipeline;
+            }
+            [pipeline release];
+            NSLog(@"sd.cpp Metal INT8 TensorOps unavailable; using SIMD-group fallback: %@", error);
+        }
+        return nil;
+    }
+
+    struct QuantParams {
+        uint32_t k, rows, group, unused;
+        uint64_t scale_offset;
+    };
+
+    struct MatParams {
+        uint32_t k, n, rows, flags;
+    };
+
+    struct Execution {
+        const ggml_ext_execution_context* context;
+        ggml_tensor* node;
+        Kernels* kernels;
+        SDI8Operation op = {};
+        int group        = 0;
+        bool linear      = false;
+    };
+
+    static bool dimensions(const ggml_tensor* node) {
+        return node->ne[0] <= INT_MAX && ggml_nrows(node) <= INT_MAX;
+    }
+
+    static bool supports_i8(ggml_backend_dev_t, const ggml_tensor* node, void*) {
+        SDI8Operation op;
+        return sd_ggml_i8_describe(node, &op) && dimensions(node) && dimensions(op.weight) &&
+               uint64_t(op.weight->ne[0]) * ggml_nrows(node) <= UINT32_MAX;
+    }
+
+    static bool supports_convrot(ggml_backend_dev_t, const ggml_tensor* node, void*) {
+        int group;
+        return sd_ggml_convrot_describe(node, &group) && dimensions(node->src[0]) &&
+               uint64_t(ggml_nelements(node->src[0])) <= UINT32_MAX;
+    }
+
+    static size_t aligned(size_t bytes) {
+        return (bytes + 255) / 256 * 256;
+    }
+
+    static size_t packed_size(int64_t k, int64_t rows) {
+        return aligned(size_t(k) * ((rows + 3) / 4 * 4) + rows * sizeof(float));
+    }
+
+    static ggml_status prepare(ggml_backend_t, ggml_tensor* const* nodes, int, void**, size_t* bytes, void*) {
+        auto* node = nodes[0];
+        SDI8Operation op;
+        const bool linear  = sd_ggml_i8_describe(node, &op);
+        const auto* input  = linear ? op.input : node->src[0];
+        const int64_t rows = linear ? ggml_nrows(node) : ggml_nrows(input);
+        int group;
+        if (linear) {
+            group = op.group_size;
+        } else if (!sd_ggml_convrot_describe(node, &group)) {
+            return GGML_STATUS_FAILED;
+        }
+        *bytes = aligned(sd_ggml_modulate_workspace_size(node));
+        if (input->type == GGML_TYPE_F32) {
+            if (group > 1)
+                *bytes += aligned(ggml_nbytes(input));
+            if (linear)
+                *bytes += packed_size(input->ne[0], rows);
+        }
+        return GGML_STATUS_SUCCESS;
+    }
+
+    static bool bind_tensor(Kernels* kernels, const ggml_tensor* tensor, ggml_ext_metal_buffer& result) {
+        auto buffer = tensor->view_src ? tensor->view_src->buffer : tensor->buffer;
+        if (!buffer || !tensor->data)
+            return false;
+        const size_t offset = uintptr_t(tensor->data) - uintptr_t(ggml_backend_buffer_get_base(buffer));
+        return kernels->buffer(buffer, offset, ggml_nbytes(tensor), &result);
+    }
+
+    static void set_buffer(id<MTLComputeCommandEncoder> encoder, const ggml_ext_metal_buffer& buffer, int index) {
+        [encoder setBuffer:(id<MTLBuffer>)buffer.buffer offset:buffer.offset atIndex:index];
+    }
+
+    static ggml_status encode(void* command_buffer, void* user_data) {
+        auto& e                      = *static_cast<Execution*>(user_data);
+        auto* kernels                = e.kernels;
+        const auto* input            = e.linear ? e.op.input : e.node->src[0];
+        const uint32_t k             = uint32_t(input->ne[0]);
+        const uint32_t rows          = uint32_t(e.linear ? ggml_nrows(e.node) : ggml_nrows(input));
+        id<MTLCommandBuffer> command = (id<MTLCommandBuffer>)command_buffer;
+        ggml_ext_metal_buffer x, dst, output, rotated, packed, xs, w, ws, bias;
+        if (!bind_tensor(kernels, input, x) || !bind_tensor(kernels, e.node, dst))
+            return GGML_STATUS_FAILED;
+        output        = dst;
+        size_t offset = 0;
+        auto scratch  = [&](size_t bytes, ggml_ext_metal_buffer& binding) {
+            bool ok = kernels->buffer(e.context->workspace, offset, bytes, &binding);
+            offset += aligned(bytes);
+            return ok;
+        };
+        const size_t output_scratch = sd_ggml_modulate_workspace_size(e.node);
+        if (output_scratch && !scratch(output_scratch, output))
+            return GGML_STATUS_FAILED;
+
+        if (input->type == GGML_TYPE_F32) {
+            rotated = x;
+            if (e.group > 1 && !scratch(ggml_nbytes(input), rotated))
+                return GGML_STATUS_FAILED;
+            packed = output;
+            if (e.linear && !scratch(packed_size(k, rows), packed))
+                return GGML_STATUS_FAILED;
+            QuantParams p                   = {k, rows, uint32_t(e.group), 0, uint64_t(k) * ((uint64_t(rows) + 3) / 4 * 4)};
+            id<MTLBlitCommandEncoder> clear = [command blitCommandEncoder];
+            [clear fillBuffer:(id<MTLBuffer>)packed.buffer range:NSMakeRange(packed.offset, e.linear ? packed_size(k, rows) : ggml_nbytes(e.node)) value:0];
+            [clear endEncoding];
+            if (e.group > 1) {
+                id<MTLComputeCommandEncoder> encoder = [command computeCommandEncoder];
+                [encoder setComputePipelineState:kernels->rotate];
+                set_buffer(encoder, x, 0);
+                set_buffer(encoder, rotated, 1);
+                [encoder setBytes:&p length:sizeof(p) atIndex:2];
+                [encoder dispatchThreadgroups:MTLSizeMake(uint64_t(k) * rows / e.group, 1, 1)
+                        threadsPerThreadgroup:MTLSizeMake(e.group, 1, 1)];
+                [encoder endEncoding];
+            }
+            id<MTLComputeCommandEncoder> encoder = [command computeCommandEncoder];
+            [encoder setComputePipelineState:kernels->quantize];
+            set_buffer(encoder, rotated, 0);
+            set_buffer(encoder, packed, 1);
+            [encoder setBytes:&p length:sizeof(p) atIndex:2];
+            [encoder dispatchThreadgroups:MTLSizeMake(rows, 1, 1) threadsPerThreadgroup:MTLSizeMake(256, 1, 1)];
+            [encoder endEncoding];
+            x  = packed;
+            xs = packed;
+            xs.offset += p.scale_offset;
+        } else if (e.op.packed) {
+            xs = x;
+            xs.offset += size_t(k) * ((uint64_t(rows) + 3) / 4 * 4);
+        } else if (!bind_tensor(kernels, e.op.input_scale, xs)) {
+            return GGML_STATUS_FAILED;
+        }
+
+        if (e.linear) {
+            if (!bind_tensor(kernels, e.op.weight, w))
+                return GGML_STATUS_FAILED;
+            ws = bias = xs;
+            if (e.op.weight_scale && !bind_tensor(kernels, e.op.weight_scale, ws))
+                return GGML_STATUS_FAILED;
+            if (e.op.bias && !bind_tensor(kernels, e.op.bias, bias))
+                return GGML_STATUS_FAILED;
+            MatParams p                          = {k, uint32_t(e.node->ne[0]), rows, uint32_t((e.op.weight_scale ? 1 : 0) | (e.op.bias ? 2 : 0))};
+            id<MTLComputeCommandEncoder> encoder = [command computeCommandEncoder];
+            [encoder setComputePipelineState:kernels->multiply_mpp ? kernels->multiply_mpp : kernels->multiply];
+            set_buffer(encoder, w, 0);
+            set_buffer(encoder, x, 1);
+            set_buffer(encoder, xs, 2);
+            set_buffer(encoder, ws, 3);
+            set_buffer(encoder, bias, 4);
+            set_buffer(encoder, output, 5);
+            [encoder setBytes:&p length:sizeof(p) atIndex:6];
+            [encoder dispatchThreadgroups:MTLSizeMake((p.n + 31) / 32, (rows + 31) / 32, 1)
+                    threadsPerThreadgroup:MTLSizeMake(128, 1, 1)];
+            [encoder endEncoding];
+        }
+        if (output_scratch) {
+            id<MTLBlitCommandEncoder> copy = [command blitCommandEncoder];
+            [copy copyFromBuffer:(id<MTLBuffer>)output.buffer
+                     sourceOffset:output.offset
+                         toBuffer:(id<MTLBuffer>)dst.buffer
+                destinationOffset:dst.offset
+                             size:ggml_nbytes(e.node)];
+            [copy endEncoding];
+        }
+        return GGML_STATUS_SUCCESS;
+    }
+
+    static ggml_status execute(const ggml_ext_execution_context* context, ggml_tensor* const* nodes, int, void*, void* user_data) {
+        Execution e = {context, nodes[0], static_cast<Kernels*>(user_data)};
+        e.linear    = sd_ggml_i8_describe(e.node, &e.op);
+        if (e.linear) {
+            e.group = e.op.group_size;
+        } else if (!sd_ggml_convrot_describe(e.node, &e.group)) {
+            return GGML_STATUS_FAILED;
+        }
+        return e.kernels->submit(context->backend, encode, &e);
+    }
+
+}
+
+void sd_ggml_register_metal_i8(ggml_backend_dev_t device) {
+    using namespace sd_metal;
+    auto reg        = ggml_backend_dev_backend_reg(device);
+    auto get_device = reinterpret_cast<ggml_ext_metal_device_fn>(ggml_backend_reg_get_proc_address(reg, "ggml_ext_metal_device_v1"));
+    auto get_buffer = reinterpret_cast<ggml_ext_metal_buffer_fn>(ggml_backend_reg_get_proc_address(reg, "ggml_ext_metal_buffer_v1"));
+    auto submit     = reinterpret_cast<ggml_ext_metal_submit_fn>(ggml_backend_reg_get_proc_address(reg, "ggml_ext_metal_submit_v1"));
+    if (!get_device || !get_buffer || !submit)
+        return;
+    @autoreleasepool {
+        id<MTLDevice> metal = (id<MTLDevice>)get_device(device);
+        if (!metal || ![metal supportsFamily:MTLGPUFamilyApple7])
+            return;
+        auto kernels               = std::make_unique<Kernels>();
+        kernels->buffer            = get_buffer;
+        kernels->submit            = submit;
+        MTLCompileOptions* options = [[MTLCompileOptions alloc] init];
+        options.fastMathEnabled    = NO;
+        NSError* error             = nil;
+        id<MTLLibrary> library     = [metal newLibraryWithSource:[NSString stringWithUTF8String:sd_metal_i8_source] options:options error:&error];
+        [options release];
+        if (!library) {
+            NSLog(@"sd.cpp Metal INT8 compilation failed: %@", error);
+            return;
+        }
+        auto pipeline = [&](NSString* name) {
+            id<MTLFunction> function           = [library newFunctionWithName:name];
+            id<MTLComputePipelineState> result = function ? [metal newComputePipelineStateWithFunction:function error:&error] : nil;
+            [function release];
+            return result;
+        };
+        kernels->rotate   = pipeline(@"sd_i8_rotate");
+        kernels->quantize = pipeline(@"sd_i8_quantize");
+        kernels->multiply = pipeline(@"sd_i8_multiply");
+        [library release];
+        if (!kernels->rotate || !kernels->quantize || !kernels->multiply ||
+            kernels->quantize.threadExecutionWidth != 32 || kernels->multiply.threadExecutionWidth != 32 ||
+            kernels->rotate.maxTotalThreadsPerThreadgroup < 256 ||
+            kernels->quantize.maxTotalThreadsPerThreadgroup < 256 || kernels->multiply.maxTotalThreadsPerThreadgroup < 128) {
+            NSLog(@"sd.cpp Metal INT8 pipelines unavailable: %@", error);
+            return;
+        }
+        kernels->multiply_mpp = tensor_multiply(metal);
+        if (kernels->multiply_mpp) {
+            NSLog(@"sd.cpp Metal INT8: using TensorOps INT8 x INT8 -> INT32");
+        }
+        ggml_ext_kernel_desc desc = {};
+        desc.abi_version          = GGML_EXT_ABI_VERSION;
+        desc.name                 = "sd.linear_i8.metal.v1";
+        desc.op                   = GGML_OP_CUSTOM;
+        desc.custom_op            = sd_ggml_i8_op();
+        desc.priority             = 100;
+        desc.supports             = supports_i8;
+        desc.prepare              = prepare;
+        desc.destroy              = [](void*, void*) {};
+        desc.execute              = execute;
+        desc.user_data            = kernels.get();
+        GGML_ASSERT(ggml_ext_register_kernel(device, &desc));
+        desc.name      = "sd.quantize_i8_convrot.metal.v1";
+        desc.custom_op = sd_ggml_convrot_op();
+        desc.supports  = supports_convrot;
+        GGML_ASSERT(ggml_ext_register_kernel(device, &desc));
+        static std::vector<std::unique_ptr<Kernels>> registrations;
+        registrations.push_back(std::move(kernels));
+    }
+}

src/core/metal/int8_mpp.metal

@@ -0,0 +1,65 @@
+#include <MetalPerformancePrimitives/MetalPerformancePrimitives.h>
+#include <metal_stdlib>
+#include <metal_tensor>
+using namespace metal;
+using namespace mpp::tensor_ops;
+
+struct MatParams {
+    uint k;
+    uint n;
+    uint rows;
+    uint flags;
+};
+
+// TensorOps requires explicit signed INT8 and non-const tensor element types.
+kernel void sd_i8_multiply_mpp(device int8_t* weight [[buffer(0)]],
+                               device int8_t* input [[buffer(1)]],
+                               device const float* xs [[buffer(2)]],
+                               device const float* ws [[buffer(3)]],
+                               device const float* bias [[buffer(4)]],
+                               device float* output [[buffer(5)]],
+                               constant MatParams& p [[buffer(6)]],
+                               uint2 tile [[threadgroup_position_in_grid]]) {
+    auto a                    = tensor<device int8_t, dextents<int, 2>, tensor_inline>(input, dextents<int, 2>(p.k, p.rows));
+    auto b                    = tensor<device int8_t, dextents<int, 2>, tensor_inline>(weight, dextents<int, 2>(p.k, p.n));
+    constexpr int chunk_size  = 1024;
+    constexpr auto descriptor = matmul2d_descriptor(32, 32, chunk_size, false, true, false);
+    matmul2d<descriptor, execution_simdgroups<4>> multiply;
+    auto left  = a.slice(0, tile.y * 32);
+    auto right = b.slice(0, tile.x * 32);
+    auto acc   = multiply.get_destination_cooperative_tensor<decltype(left), decltype(right), int>();
+    auto sums  = multiply.get_destination_cooperative_tensor<decltype(left), decltype(right), int>();
+#pragma unroll
+    for (uint i = 0; i < sums.get_capacity(); ++i) {
+        if (sums.is_valid_element(i))
+            sums[i] = 0;
+    }
+    // TensorOps can round long integer dot products on older GPUs. Bound each
+    // partial sum to FP32's exact integer range, then combine with wrapping
+    // integer addition outside TensorOps.
+    for (uint k = 0; k < p.k; k += chunk_size) {
+        left  = a.slice(k, tile.y * 32);
+        right = b.slice(k, tile.x * 32);
+        multiply.run(left, right, acc);
+#pragma unroll
+        for (uint i = 0; i < acc.get_capacity(); ++i) {
+            if (acc.is_valid_element(i))
+                sums[i] = as_type<int>(as_type<uint>(sums[i]) + as_type<uint>(acc[i]));
+        }
+    }
+#pragma unroll
+    for (uint i = 0; i < acc.get_capacity(); ++i) {
+        if (acc.is_valid_element(i)) {
+            auto position = acc.get_multidimensional_index(i);
+            uint col = tile.x * 32 + position[0], row = tile.y * 32 + position[1];
+            if (col < p.n && row < p.rows) {
+                float value = float(sums[i]) * xs[row];
+                if (p.flags & 1)
+                    value *= ws[col];
+                if (p.flags & 2)
+                    value += bias[col];
+                output[ulong(row) * p.n + col] = value;
+            }
+        }
+    }
+}

src/core/metal/int8_source.h.in

@@ -0,0 +1,7 @@
+#ifndef __SD_CORE_METAL_INT8_SOURCE_H__
+#define __SD_CORE_METAL_INT8_SOURCE_H__
+
+static const char sd_metal_i8_source[] = R"SD_METAL(@SD_METAL_I8_SOURCE@)SD_METAL";
+static const char sd_metal_i8_mpp_source[] = R"SD_METAL(@SD_METAL_I8_MPP_SOURCE@)SD_METAL";
+
+#endif

src/model_manager_prefetch.cpp

@@ -33,7 +33,10 @@ void ModelManager::synchronize_prefetch_block(PrefetchBlock& block) {
         ggml_backend_event_synchronize(block.event);
         ggml_backend_event_free(block.event);
         block.event = nullptr;
-    } else if (block.submitted && block.transfer_backend != nullptr) {
+    }
+    // Event completion does not release backend-owned transfer resources, such
+    // as Metal's temporary host mappings and command buffers.
+    if (block.submitted && block.transfer_backend != nullptr) {
         ggml_backend_synchronize(block.transfer_backend);
     }
     block.transfer_backend = nullptr;