Merge branch 'master' into avx512

.gitmodules

@@ -1,3 +0,0 @@
-[submodule "bindings/ios"]
-	path = bindings/ios
-	url = https://github.com/ggerganov/whisper.spm

CMakeLists.txt

@@ -9,11 +9,6 @@ if(CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
     set(WHISPER_STANDALONE ON)
     include(cmake/GitVars.cmake)
     include(cmake/BuildTypes.cmake)
-
-    # configure project version
-    if (EXISTS "${CMAKE_SOURCE_DIR}/bindings/ios/Makefile-tmpl")
-        configure_file(${CMAKE_SOURCE_DIR}/bindings/ios/Makefile-tmpl ${CMAKE_SOURCE_DIR}/bindings/ios/Makefile @ONLY)
-    endif()
 else()
     set(WHISPER_STANDALONE OFF)
 endif()
@@ -94,17 +89,6 @@ if (APPLE AND NOT WHISPER_NO_ACCELERATE)
     else()
         message(WARNING "Accelerate framework not found")
     endif()
-
-    find_library(FOUNDATION_LIBRARY Foundation REQUIRED)
-    find_library(METAL_FRAMEWORK    Metal      REQUIRED)
-    find_library(METALKIT_FRAMEWORK MetalKit   REQUIRED)
-    find_library(METALPERFORMANCE_FRAMEWORK MetalPerformanceShaders REQUIRED)
-
-    set(WHISPER_EXTRA_LIBS ${WHISPER_EXTRA_LIBS}
-        ${FOUNDATION_LIBRARY}
-        ${METAL_FRAMEWORK}
-        ${METALKIT_FRAMEWORK}
-        ${METALPERFORMANCE_FRAMEWORK})
 endif()
 
 if (WHISPER_SUPPORT_OPENBLAS)
@@ -179,7 +163,6 @@ set(TARGET whisper)
 
 add_library(${TARGET}
     ggml.c
-    ggml-mtl.m
     whisper.cpp
     )
 

Makefile

@@ -50,7 +50,11 @@ endif
 # TODO: probably these flags need to be tweaked on some architectures
 #       feel free to update the Makefile for your architecture and send a pull request or issue
 ifeq ($(UNAME_M),x86_64)
-	CFLAGS += -mavx -mavx2 -mfma -mf16c
+	# AVX 512
+	CFLAGS += -mavx512f -mfma -mf16c
+
+	# AVX 256
+	#CFLAGS += -mavx -mavx2 -mfma -mf16c
 endif
 ifeq ($(UNAME_M),amd64)
 	CFLAGS += -mavx -mavx2 -mfma -mf16c
@@ -58,8 +62,8 @@ endif
 ifndef WHISPER_NO_ACCELERATE
 	# Mac M1 - include Accelerate framework
 	ifeq ($(UNAME_S),Darwin)
-		CFLAGS  += -DGGML_USE_ACCELERATE -DGGML_PERF
+		CFLAGS  += -DGGML_USE_ACCELERATE
-		LDFLAGS += -framework Foundation -framework Accelerate -framework Metal -framework MetalKit -framework MetalPerformanceShaders
+		LDFLAGS += -framework Accelerate
 	endif
 endif
 ifneq ($(filter aarch64%,$(UNAME_M)),)
@@ -81,21 +85,18 @@ endif
 # Build library + main
 #
 
-main: examples/main/main.cpp ggml.o ggml-mtl.o whisper.o
+main: examples/main/main.cpp ggml.o whisper.o
-	$(CXX) $(CXXFLAGS) examples/main/main.cpp whisper.o ggml.o ggml-mtl.o -o main $(LDFLAGS)
+	$(CXX) $(CXXFLAGS) examples/main/main.cpp whisper.o ggml.o -o main $(LDFLAGS)
 	./main -h
 
 ggml.o: ggml.c ggml.h
 	$(CC)  $(CFLAGS)   -c ggml.c -o ggml.o
 
-ggml-mtl.o: ggml-mtl.m ggml-mtl.h
-	$(CC)  $(CFLAGS)   -c ggml-mtl.m -o ggml-mtl.o
-
 whisper.o: whisper.cpp whisper.h
 	$(CXX) $(CXXFLAGS) -c whisper.cpp -o whisper.o
 
-libwhisper.a: ggml.o ggml-mtl.o whisper.o
+libwhisper.a: ggml.o whisper.o
-	$(AR) rcs libwhisper.a ggml.o ggml-mtl.o whisper.o
+	$(AR) rcs libwhisper.a ggml.o whisper.o
 
 clean:
 	rm -f *.o main stream bench libwhisper.a

README.md

@@ -437,12 +437,9 @@ For more details, see the conversion script [models/convert-pt-to-ggml.py](model
 ## Bindings
 
 - [X] Rust: [tazz4843/whisper-rs](https://github.com/tazz4843/whisper-rs)
-- [X] Objective-C / Swift: [ggerganov/whisper.spm](https://github.com/ggerganov/whisper.spm)
 - [ ] Python:
 - [ ] Java:
 
 ## Examples
 
 There are various examples of using the library for different projects in the [examples](examples) folder. Check them out!
-
-## [Frequently asked questions (#126)](https://github.com/ggerganov/whisper.cpp/discussions/126)

examples/generate-karaoke.sh

@@ -1,49 +0,0 @@
-#!/bin/bash
-
-executable="./main"
-model="base.en"
-model_path="models/ggml-$model.bin"
-
-# require sox and ffmpeg to be installed
-if ! command -v sox &> /dev/null
-then
-    echo "sox could not be found"
-    exit 1
-fi
-
-if ! command -v ffmpeg &> /dev/null
-then
-    echo "ffmpeg could not be found"
-    exit 2
-fi
-
-if [ ! -f "$executable" ]; then
-    echo "'$executable' does not exist. Please build it first."
-    exit 3
-fi
-
-if [ ! -f "$model_path" ]; then
-    echo "'$model_path' does not exist. Please download it first."
-    exit 4
-fi
-
-# record some raw audio
-sox -d rec.wav
-
-# resample to 16kHz
-ffmpeg -y -i ./rec.wav -ar 16000 -ac 1 -c:a pcm_s16le ./rec16.wav > /dev/null 2>&1
-
-# run Whisper
-echo "Processing ..."
-./main -m models/ggml-base.en.bin rec16.wav -owts > /dev/null 2>&1
-
-# generate Karaoke video
-echo "Generating video ..."
-source rec16.wav.wts > /dev/null 2>&1
-
-# play the video
-echo "Playing ./rec16.wav.mp4 ..."
-ffplay -loglevel 0 -autoexit ./rec16.wav.mp4
-
-echo "Done"
-exit 0

examples/main/main.cpp

@@ -53,7 +53,6 @@ struct whisper_params {
     int32_t n_processors = 1;
     int32_t offset_t_ms  = 0;
     int32_t offset_n     = 0;
-    int32_t duration_ms  = 0;
     int32_t max_context  = -1;
     int32_t max_len      = 0;
 
@@ -96,8 +95,6 @@ bool whisper_params_parse(int argc, char ** argv, whisper_params & params) {
             params.offset_t_ms = std::stoi(argv[++i]);
         } else if (arg == "-on" || arg == "--offset-n") {
             params.offset_n = std::stoi(argv[++i]);
-        } else if (arg == "-d" || arg == "--duration") {
-            params.duration_ms = std::stoi(argv[++i]);
         } else if (arg == "-mc" || arg == "--max-context") {
             params.max_context = std::stoi(argv[++i]);
         } else if (arg == "-ml" || arg == "--max-len") {
@@ -157,7 +154,6 @@ void whisper_print_usage(int argc, char ** argv, const whisper_params & params)
     fprintf(stderr, "  -p N,     --processors N   number of processors to use during computation (default: %d)\n", params.n_processors);
     fprintf(stderr, "  -ot N,    --offset-t N     time offset in milliseconds (default: %d)\n", params.offset_t_ms);
     fprintf(stderr, "  -on N,    --offset-n N     segment index offset (default: %d)\n", params.offset_n);
-    fprintf(stderr, "  -d  N,    --duration N     duration of audio to process in milliseconds (default: %d)\n", params.duration_ms);
     fprintf(stderr, "  -mc N,    --max-context N  maximum number of text context tokens to store (default: max)\n");
     fprintf(stderr, "  -ml N,    --max-len N      maximum segment length in characters (default: %d)\n", params.max_len);
     fprintf(stderr, "  -wt N,    --word-thold N   word timestamp probability threshold (default: %f)\n", params.word_thold);
@@ -536,7 +532,6 @@ int main(int argc, char ** argv) {
             wparams.n_threads            = params.n_threads;
             wparams.n_max_text_ctx       = params.max_context >= 0 ? params.max_context : wparams.n_max_text_ctx;
             wparams.offset_ms            = params.offset_t_ms;
-            wparams.duration_ms          = params.duration_ms;
 
             wparams.token_timestamps     = params.output_wts || params.max_len > 0;
             wparams.thold_pt             = params.word_thold;

ggml-mtl.h

@@ -1,38 +0,0 @@
-#pragma once
-
-#include <stdint.h>
-#include <stddef.h>
-
-// TODO: this will hold dynamic context data in the future
-//       currently unused
-struct ggml_mtl_context {
-    void * dummy;
-};
-
-struct ggml_mtl_object {
-    int32_t id;
-    void * data;
-};
-
-struct ggml_mtl_context * ggml_mtl_init(void);
-
-struct ggml_mtl_object ggml_mtl_alloc(size_t size);
-
-// multiply matrix by vector
-void ggml_mtl_mul_mat_vec_f16(
-    struct ggml_mtl_context * ctx,
-    struct ggml_mtl_object    src0,  // matrix f16
-    const __fp16            * src1,  // vector f16
-    float                   * dst,   // vector f32
-    int                       nrows,
-    int                       ncols);
-
-// multiply matrix by matrix
-void ggml_mtl_mul_mat_f16(
-    struct ggml_mtl_context * ctx,
-    struct ggml_mtl_object    src0,  // matrix f16
-    const __fp16            * src1,  // matrix f16
-    float                   * dst,   // matrix f32
-    int                       nrows0,
-    int                       nrows1,
-    int                       ncols);

ggml-mtl.m

@@ -1,162 +0,0 @@
-#import "ggml-mtl.h"
-
-#import <Foundation/Foundation.h>
-#import <Metal/Metal.h>
-#import <MetalPerformanceShaders/MetalPerformanceShaders.h>
-
-#define GGML_MTL_MAX_BUFFERS 256
-
-// global static storage for Metal buffers
-// TODO: move this into a dynamic context
-static id<MTLBuffer> g_buffers[GGML_MTL_MAX_BUFFERS];
-
-// global MTL context
-// TODO: move this into a dynamic context
-static id<MTLDevice>       g_device;
-static id<MTLCommandQueue> g_command_queue;
-
-struct ggml_mtl_context * ggml_mtl_init() {
-    // TODO: implement properly
-    //       for now, init the global MTL context and MTL buffers
-    g_device = MTLCreateSystemDefaultDevice();
-
-    g_command_queue = [g_device newCommandQueue];
-    if (g_command_queue == nil)
-    {
-        NSLog(@"Failed to find the command queue.");
-        return nil;
-    }
-
-    return nil;
-}
-
-// search for unallocated buffer slot and use it
-struct ggml_mtl_object ggml_mtl_alloc(size_t size) {
-    // TODO: temporarily making sure that the buffers are nil at the start
-    static bool first = true;
-    if (first) {
-        for (int i = 0; i < GGML_MTL_MAX_BUFFERS; ++i) {
-            assert(g_buffers[i] == nil);
-        }
-        first = false;
-    }
-
-    struct ggml_mtl_object obj = { -1, nil };
-
-    for (int i = 0; i < GGML_MTL_MAX_BUFFERS; i++) {
-        if (g_buffers[i] == nil) {
-            g_buffers[i] = [g_device newBufferWithLength:size options:MTLResourceStorageModeManaged];
-
-            // lunk the MTL buffer to the ggml object
-            obj.id = i;
-            obj.data = [g_buffers[i] contents];
-
-            break;
-        }
-    }
-
-    return obj;
-}
-
-struct params_mul_mat_vec {
-    int  N; // rows
-    int  M; // cols
-};
-
-// multiply matrix with a vector using MPSMatrixVectorMultiplication
-void ggml_mtl_mul_mat_vec_f16(
-        struct ggml_mtl_context * ctx,
-        struct ggml_mtl_object    src0,
-        const __fp16            * src1,
-        float                   * dst,
-        int                       nrows,
-        int                       ncols) {
-    (void) ctx; // unused
-
-    // Create a command buffer to hold commands.
-    id<MTLCommandBuffer> commandBuffer = [g_command_queue commandBuffer];
-    assert(commandBuffer != nil);
-
-    // make managed device buffer to store src1
-    id<MTLBuffer> src1_buffer = [g_device newBufferWithBytes:src1 length:ncols*sizeof(__fp16) options:MTLResourceStorageModeManaged];
-    id<MTLBuffer> dst_buffer  = [g_device newBufferWithLength:nrows*sizeof(float) options:MTLResourceStorageModeManaged];
-
-    // MPSMatrixDescriptor
-    MPSMatrixDescriptor *src0_desc = [MPSMatrixDescriptor matrixDescriptorWithRows:nrows columns:ncols rowBytes:ncols*sizeof(__fp16) dataType:MPSDataTypeFloat16];
-    MPSVectorDescriptor *src1_desc = [MPSVectorDescriptor vectorDescriptorWithLength:ncols dataType:MPSDataTypeFloat16];
-    MPSVectorDescriptor *dst_desc  = [MPSVectorDescriptor vectorDescriptorWithLength:nrows dataType:MPSDataTypeFloat32];
-
-    // MPSMatrix
-    MPSMatrix *src0_mat = [[MPSMatrix alloc] initWithBuffer:g_buffers[src0.id] descriptor:src0_desc];
-    MPSVector *src1_vec = [[MPSVector alloc] initWithBuffer:src1_buffer        descriptor:src1_desc];
-    MPSVector *dst_vec  = [[MPSVector alloc] initWithBuffer:dst_buffer         descriptor:dst_desc];
-
-    // MPSMatrixVectorMultiplication
-    MPSMatrixVectorMultiplication *mul_mat_vec = [[MPSMatrixVectorMultiplication alloc] initWithDevice:g_device transpose:NO rows:nrows columns:ncols alpha:1.0 beta:0.0];
-
-    // encode
-    [mul_mat_vec encodeToCommandBuffer:commandBuffer
-                           inputMatrix:src0_mat
-                           inputVector:src1_vec
-                          resultVector:dst_vec];
-
-    [commandBuffer commit];
-    [commandBuffer waitUntilCompleted];
-
-    // copy GPU result to CPU
-    memcpy(dst, [dst_buffer contents], nrows*sizeof(float));
-}
-
-// multiply matrix with a matrix using MPSMatrixMultiplication
-void ggml_mtl_mul_mat_f16(
-        struct ggml_mtl_context * ctx,
-        struct ggml_mtl_object    src0,
-        const __fp16            * src1,
-        float                   * dst,
-        int                       nrows0,
-        int                       nrows1,
-        int                       ncols) {
-    (void) ctx; // unused
-
-    // Create a command buffer to hold commands.
-    id<MTLCommandBuffer> commandBuffer = [g_command_queue commandBuffer];
-    assert(commandBuffer != nil);
-
-    // make managed device buffer to store src1
-    id<MTLBuffer> src1_buffer = [g_device newBufferWithBytes:src1 length:ncols*nrows1*sizeof(__fp16) options:MTLResourceStorageModeManaged];
-    id<MTLBuffer> dst_buffer  = [g_device newBufferWithLength:nrows0*nrows1*sizeof(float) options:MTLResourceStorageModeManaged];
-
-    // MPSMatrixDescriptor
-    MPSMatrixDescriptor *src0_desc = [MPSMatrixDescriptor matrixDescriptorWithRows:nrows0 columns:ncols  rowBytes:ncols*sizeof(__fp16) dataType:MPSDataTypeFloat16];
-    MPSMatrixDescriptor *src1_desc = [MPSMatrixDescriptor matrixDescriptorWithRows:nrows1 columns:ncols  rowBytes:ncols*sizeof(__fp16) dataType:MPSDataTypeFloat16];
-    MPSMatrixDescriptor *dst_desc  = [MPSMatrixDescriptor matrixDescriptorWithRows:nrows1 columns:nrows0 rowBytes:nrows0*sizeof(float) dataType:MPSDataTypeFloat32];
-
-    // MPSMatrix
-    MPSMatrix *src0_mat = [[MPSMatrix alloc] initWithBuffer:g_buffers[src0.id] descriptor:src0_desc];
-    MPSMatrix *src1_mat = [[MPSMatrix alloc] initWithBuffer:src1_buffer        descriptor:src1_desc];
-    MPSMatrix *dst_mat  = [[MPSMatrix alloc] initWithBuffer:dst_buffer         descriptor:dst_desc];
-
-    //// MPSMatrixMultiplication z = x * yT
-    //MPSMatrixMultiplication *mul_mat = [[MPSMatrixMultiplication alloc] initWithDevice:g_device transposeLeft:NO transposeRight:YES resultRows:nrows resultColumns:nrows interiorColumns:ncols alpha:1.0 beta:0.0];
-
-    //// encode
-    //[mul_mat encodeToCommandBuffer:commandBuffer
-    //                   leftMatrix:src0_mat
-    //                  rightMatrix:src1_mat
-    //                 resultMatrix:dst_mat];
-
-    // MPSMatrixMultiplication zT = xT * y
-    MPSMatrixMultiplication *mul_mat = [[MPSMatrixMultiplication alloc] initWithDevice:g_device transposeLeft:NO transposeRight:YES resultRows:nrows1 resultColumns:nrows0 interiorColumns:ncols alpha:1.0 beta:0.0];
-
-    // encode
-    [mul_mat encodeToCommandBuffer:commandBuffer
-                       leftMatrix:src1_mat
-                      rightMatrix:src0_mat
-                     resultMatrix:dst_mat];
-
-    [commandBuffer commit];
-    [commandBuffer waitUntilCompleted];
-
-    // copy GPU result to CPU
-    memcpy(dst, [dst_buffer contents], nrows0*nrows1*sizeof(float));
-}

ggml.c

@@ -1,7 +1,5 @@
 #include "ggml.h"
 
-#include "ggml-mtl.h"
-
 #if defined(_MSC_VER) || defined(__MINGW32__)
 #include <malloc.h> // using malloc.h with MSC/MINGW
 #elif !defined(__FreeBSD__)
@@ -329,6 +327,45 @@ inline static void ggml_vec_dot_f32(const int n, float * restrict s, const float
     for (int i = n16; i < n; ++i) {
         sumf += x[i]*y[i];
     }
+#elif defined(__AVX512F__)
+    const int n64 = (n & ~63);
+
+    __m512 sum0 = _mm512_setzero_ps();
+    __m512 sum1 = _mm512_setzero_ps();
+    __m512 sum2 = _mm512_setzero_ps();
+    __m512 sum3 = _mm512_setzero_ps();
+
+    __m512 x0, x1, x2, x3;
+    __m512 y0, y1, y2, y3;
+
+    for (int i = 0; i < n64; i += 64) {
+        x0 = _mm512_loadu_ps(x + i + 0);
+        x1 = _mm512_loadu_ps(x + i + 16);
+        x2 = _mm512_loadu_ps(x + i + 32);
+        x3 = _mm512_loadu_ps(x + i + 48);
+
+        y0 = _mm512_loadu_ps(y + i + 0);
+        y1 = _mm512_loadu_ps(y + i + 16);
+        y2 = _mm512_loadu_ps(y + i + 32);
+        y3 = _mm512_loadu_ps(y + i + 48);
+
+        sum0 = _mm512_fmadd_ps(x0, y0, sum0);
+        sum1 = _mm512_fmadd_ps(x1, y1, sum1);
+        sum2 = _mm512_fmadd_ps(x2, y2, sum2);
+        sum3 = _mm512_fmadd_ps(x3, y3, sum3);
+    }
+
+    sum0 = _mm512_add_ps(sum0, sum1);
+    sum2 = _mm512_add_ps(sum2, sum3);
+    sum0 = _mm512_add_ps(sum0, sum2);
+
+    sumf = sum0[0] + sum0[1] + sum0[2] + sum0[3] + sum0[4] + sum0[5] + sum0[6] + sum0[7] +
+           sum0[8] + sum0[9] + sum0[10] + sum0[11] + sum0[12] + sum0[13] + sum0[14] + sum0[15];
+
+    // leftovers
+    for (int i = n64; i < n; ++i) {
+        sumf += x[i]*y[i];
+    }
 #elif defined(__AVX2__)
     // AVX 256-bit
     const int n32 = (n & ~31);
@@ -526,6 +563,47 @@ inline static void ggml_vec_dot_f16(const int n, float * restrict s, ggml_fp16_t
     for (int i = n32; i < n; ++i) {
         sumf += ggml_fp16_to_fp32(x[i])*ggml_fp16_to_fp32(y[i]);
     }
+#elif defined(__AVX512F__)
+    // AVX 512-bit
+    const int n64 = (n & ~63);
+
+    __m512 sum0 = _mm512_setzero_ps();
+    __m512 sum1 = _mm512_setzero_ps();
+    __m512 sum2 = _mm512_setzero_ps();
+    __m512 sum3 = _mm512_setzero_ps();
+
+    __m512 x0, x1, x2, x3;
+    __m512 y0, y1, y2, y3;
+
+    for (int i = 0; i < n64; i += 64) {
+        x0 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(x + i + 0 )));
+        x1 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(x + i + 16)));
+        x2 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(x + i + 32)));
+        x3 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(x + i + 48)));
+
+        y0 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(y + i + 0 )));
+        y1 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(y + i + 16)));
+        y2 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(y + i + 32)));
+        y3 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(y + i + 48)));
+
+        sum0 = _mm512_fmadd_ps(x0, y0, sum0);
+        sum1 = _mm512_fmadd_ps(x1, y1, sum1);
+        sum2 = _mm512_fmadd_ps(x2, y2, sum2);
+        sum3 = _mm512_fmadd_ps(x3, y3, sum3);
+    }
+
+    const __m512 sum01 = _mm512_add_ps(sum0, sum1);
+    const __m512 sum23 = _mm512_add_ps(sum2, sum3);
+    const __m512 sum0123 = _mm512_add_ps(sum01, sum23);
+
+    sumf = sum0123[0] + sum0123[1] + sum0123[2]  + sum0123[3]  + sum0123[4]  + sum0123[5]  + sum0123[6]  + sum0123[7] +
+           sum0123[8] + sum0123[9] + sum0123[10] + sum0123[11] + sum0123[12] + sum0123[13] + sum0123[14] + sum0123[15];
+
+    // leftovers
+    for (int i = n64; i < n; ++i) {
+        //GGML_ASSERT(false);
+        sumf += ggml_fp16_to_fp32(x[i])*ggml_fp16_to_fp32(y[i]);
+    }
 #elif defined(__AVX2__)
     // AVX 256-bit
     const int n32 = (n & ~31);
@@ -632,7 +710,7 @@ inline static void ggml_vec_mad_f32(const int n, float * restrict y, const float
     // NEON 128-bit
     const int n16 = (n & ~15);
 
-    const float32x4_t v4 = vdupq_n_f32(v);
+    const float32x4_t v0 = vdupq_n_f32(v);
 
     float32x4_t x0, x1, x2, x3;
     float32x4_t y0, y1, y2, y3;
@@ -648,14 +726,14 @@ inline static void ggml_vec_mad_f32(const int n, float * restrict y, const float
         y2 = vld1q_f32(y + i + 8);
         y3 = vld1q_f32(y + i + 12);
 
-        y0 = vfmaq_f32(y0, x0, v4);
+        y0 = vfmaq_f32(y0, x0, v0);
-        y1 = vfmaq_f32(y1, x1, v4);
+        y1 = vfmaq_f32(y1, x1, v0);
-        y2 = vfmaq_f32(y2, x2, v4);
+        y2 = vfmaq_f32(y2, x2, v0);
-        y3 = vfmaq_f32(y3, x3, v4);
+        y3 = vfmaq_f32(y3, x3, v0);
 
-        vst1q_f32(y + i + 0, y0);
+        vst1q_f32(y + i + 0,  y0);
-        vst1q_f32(y + i + 4, y1);
+        vst1q_f32(y + i + 4,  y1);
-        vst1q_f32(y + i + 8, y2);
+        vst1q_f32(y + i + 8,  y2);
         vst1q_f32(y + i + 12, y3);
     }
 
@@ -663,11 +741,46 @@ inline static void ggml_vec_mad_f32(const int n, float * restrict y, const float
     for (int i = n16; i < n; ++i) {
         y[i] += x[i]*v;
     }
+#elif defined(__AVX512F__)
+    // AVX512 512-bit
+    const int n64 = (n & ~63);
+
+    const __m512 v0 = _mm512_set1_ps(v);
+
+    __m512 x0, x1, x2, x3;
+    __m512 y0, y1, y2, y3;
+
+    for (int i = 0; i < n64; i += 64) {
+        x0 = _mm512_loadu_ps(x + i + 0);
+        x1 = _mm512_loadu_ps(x + i + 16);
+        x2 = _mm512_loadu_ps(x + i + 32);
+        x3 = _mm512_loadu_ps(x + i + 48);
+
+        y0 = _mm512_loadu_ps(y + i + 0);
+        y1 = _mm512_loadu_ps(y + i + 16);
+        y2 = _mm512_loadu_ps(y + i + 32);
+        y3 = _mm512_loadu_ps(y + i + 48);
+
+        y0 = _mm512_fmadd_ps(x0, v0, y0);
+        y1 = _mm512_fmadd_ps(x1, v0, y1);
+        y2 = _mm512_fmadd_ps(x2, v0, y2);
+        y3 = _mm512_fmadd_ps(x3, v0, y3);
+
+        _mm512_storeu_ps(y + i + 0,  y0);
+        _mm512_storeu_ps(y + i + 16, y1);
+        _mm512_storeu_ps(y + i + 32, y2);
+        _mm512_storeu_ps(y + i + 48, y3);
+    }
+
+    // leftovers
+    for (int i = n64; i < n; ++i) {
+        y[i] += x[i]*v;
+    }
 #elif defined(__AVX2__)
     // AVX 256-bit
     const int n32 = (n & ~31);
 
-    const __m256 v4 = _mm256_set1_ps(v);
+    const __m256 v0 = _mm256_set1_ps(v);
 
     __m256 x0, x1, x2, x3;
     __m256 y0, y1, y2, y3;
@@ -683,13 +796,13 @@ inline static void ggml_vec_mad_f32(const int n, float * restrict y, const float
         y2 = _mm256_loadu_ps(y + i + 16);
         y3 = _mm256_loadu_ps(y + i + 24);
 
-        y0 = _mm256_fmadd_ps(x0, v4, y0);
+        y0 = _mm256_fmadd_ps(x0, v0, y0);
-        y1 = _mm256_fmadd_ps(x1, v4, y1);
+        y1 = _mm256_fmadd_ps(x1, v0, y1);
-        y2 = _mm256_fmadd_ps(x2, v4, y2);
+        y2 = _mm256_fmadd_ps(x2, v0, y2);
-        y3 = _mm256_fmadd_ps(x3, v4, y3);
+        y3 = _mm256_fmadd_ps(x3, v0, y3);
 
-        _mm256_storeu_ps(y + i + 0, y0);
+        _mm256_storeu_ps(y + i + 0,  y0);
-        _mm256_storeu_ps(y + i + 8, y1);
+        _mm256_storeu_ps(y + i + 8,  y1);
         _mm256_storeu_ps(y + i + 16, y2);
         _mm256_storeu_ps(y + i + 24, y3);
     }
@@ -702,7 +815,7 @@ inline static void ggml_vec_mad_f32(const int n, float * restrict y, const float
     // WASM SIMD 128-bit
     const int n16 = (n & ~15);
 
-    const v128_t v4 = wasm_f32x4_splat(v);
+    const v128_t v0 = wasm_f32x4_splat(v);
 
     v128_t x0, x1, x2, x3;
     v128_t y0, y1, y2, y3;
@@ -718,10 +831,10 @@ inline static void ggml_vec_mad_f32(const int n, float * restrict y, const float
         y2 = wasm_v128_load(y + i + 8);
         y3 = wasm_v128_load(y + i + 12);
 
-        y0 = wasm_f32x4_add(y0, wasm_f32x4_mul(x0, v4));
+        y0 = wasm_f32x4_add(y0, wasm_f32x4_mul(x0, v0));
-        y1 = wasm_f32x4_add(y1, wasm_f32x4_mul(x1, v4));
+        y1 = wasm_f32x4_add(y1, wasm_f32x4_mul(x1, v0));
-        y2 = wasm_f32x4_add(y2, wasm_f32x4_mul(x2, v4));
+        y2 = wasm_f32x4_add(y2, wasm_f32x4_mul(x2, v0));
-        y3 = wasm_f32x4_add(y3, wasm_f32x4_mul(x3, v4));
+        y3 = wasm_f32x4_add(y3, wasm_f32x4_mul(x3, v0));
 
         wasm_v128_store(y + i + 0, y0);
         wasm_v128_store(y + i + 4, y1);
@@ -747,7 +860,7 @@ inline static void ggml_vec_mad_f16(const int n, ggml_fp16_t * restrict y, ggml_
     const int n32 = (n & ~31);
 
 #if defined(__ARM_FEATURE_FP16_VECTOR_ARITHMETIC)
-    const float16x8_t v8 = vdupq_n_f16(v);
+    const float16x8_t v0 = vdupq_n_f16(v);
 
     float16x8_t x0, x1, x2, x3;
     float16x8_t y0, y1, y2, y3;
@@ -763,10 +876,10 @@ inline static void ggml_vec_mad_f16(const int n, ggml_fp16_t * restrict y, ggml_
         x2 = vld1q_f16(x + i + 16);
         x3 = vld1q_f16(x + i + 24);
 
-        y0 = vfmaq_f16(y0, x0, v8);
+        y0 = vfmaq_f16(y0, x0, v0);
-        y1 = vfmaq_f16(y1, x1, v8);
+        y1 = vfmaq_f16(y1, x1, v0);
-        y2 = vfmaq_f16(y2, x2, v8);
+        y2 = vfmaq_f16(y2, x2, v0);
-        y3 = vfmaq_f16(y3, x3, v8);
+        y3 = vfmaq_f16(y3, x3, v0);
 
         vst1q_f16(y + i + 0 , y0);
         vst1q_f16(y + i + 8 , y1);
@@ -774,8 +887,7 @@ inline static void ggml_vec_mad_f16(const int n, ggml_fp16_t * restrict y, ggml_
         vst1q_f16(y + i + 24, y3);
     }
 #else
-    const float32x4_t v40 = vdupq_n_f32(v);
+    const float32x4_t v0 = vdupq_n_f32(v);
-    const float32x4_t v41 = vdupq_n_f32(v);
 
     float32x4_t x0, x1, x2, x3, x4, x5, x6, x7;
     float32x4_t y0, y1, y2, y3, y4, y5, y6, y7;
@@ -799,14 +911,14 @@ inline static void ggml_vec_mad_f16(const int n, ggml_fp16_t * restrict y, ggml_
         x6 = vcvt_f32_f16(vld1_f16(x + i + 24));
         x7 = vcvt_f32_f16(vld1_f16(x + i + 28));
 
-        y0 = vfmaq_f32(y0, x0, v40);
+        y0 = vfmaq_f32(y0, x0, v0);
-        y1 = vfmaq_f32(y1, x1, v40);
+        y1 = vfmaq_f32(y1, x1, v0);
-        y2 = vfmaq_f32(y2, x2, v40);
+        y2 = vfmaq_f32(y2, x2, v0);
-        y3 = vfmaq_f32(y3, x3, v40);
+        y3 = vfmaq_f32(y3, x3, v0);
-        y4 = vfmaq_f32(y4, x4, v41);
+        y4 = vfmaq_f32(y4, x4, v0);
-        y5 = vfmaq_f32(y5, x5, v41);
+        y5 = vfmaq_f32(y5, x5, v0);
-        y6 = vfmaq_f32(y6, x6, v41);
+        y6 = vfmaq_f32(y6, x6, v0);
-        y7 = vfmaq_f32(y7, x7, v41);
+        y7 = vfmaq_f32(y7, x7, v0);
 
         vst1_f16(y + i + 0 , vcvt_f16_f32(y0));
         vst1_f16(y + i + 4 , vcvt_f16_f32(y1));
@@ -824,11 +936,47 @@ inline static void ggml_vec_mad_f16(const int n, ggml_fp16_t * restrict y, ggml_
         GGML_ASSERT(false);
         y[i] = ggml_fp32_to_fp16(ggml_fp16_to_fp32(y[i]) + ggml_fp16_to_fp32(x[i])*v);
     }
+#elif defined(__AVX512F__)
+    // AVX 512-bit
+    const int n64 = (n & ~63);
+
+    const __m512 v0 = _mm512_set1_ps(v);
+
+    __m512 x0, x1, x2, x3;
+    __m512 y0, y1, y2, y3;
+
+    for (int i = 0; i < n64; i += 64) {
+        x0 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(x + i + 0 )));
+        x1 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(x + i + 16)));
+        x2 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(x + i + 32)));
+        x3 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(x + i + 48)));
+
+        y0 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(y + i + 0 )));
+        y1 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(y + i + 16)));
+        y2 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(y + i + 32)));
+        y3 = _mm512_cvtph_ps(_mm256_loadu_si256((__m256i*)(y + i + 48)));
+
+        y0 = _mm512_fmadd_ps(x0, v0, y0);
+        y1 = _mm512_fmadd_ps(x1, v0, y1);
+        y2 = _mm512_fmadd_ps(x2, v0, y2);
+        y3 = _mm512_fmadd_ps(x3, v0, y3);
+
+        _mm256_storeu_si256((__m256i*)(y + i + 0 ), _mm512_cvtps_ph(y0, 0));
+        _mm256_storeu_si256((__m256i*)(y + i + 16), _mm512_cvtps_ph(y1, 0));
+        _mm256_storeu_si256((__m256i*)(y + i + 32), _mm512_cvtps_ph(y2, 0));
+        _mm256_storeu_si256((__m256i*)(y + i + 48), _mm512_cvtps_ph(y3, 0));
+    }
+
+    // leftovers
+    for (int i = n64; i < n; ++i) {
+        GGML_ASSERT(false);
+        y[i] = ggml_fp32_to_fp16(ggml_fp16_to_fp32(y[i]) + ggml_fp16_to_fp32(x[i])*v);
+    }
 #elif defined(__AVX2__)
     // AVX 256-bit
     const int n32 = (n & ~31);
 
-    const __m256 v8 = _mm256_set1_ps(v);
+    const __m256 v0 = _mm256_set1_ps(v);
 
     __m256 x0, x1, x2, x3;
     __m256 y0, y1, y2, y3;
@@ -844,10 +992,10 @@ inline static void ggml_vec_mad_f16(const int n, ggml_fp16_t * restrict y, ggml_
         x2 = _mm256_cvtph_ps(_mm_loadu_si128((__m128i*)(x + i + 16)));
         x3 = _mm256_cvtph_ps(_mm_loadu_si128((__m128i*)(x + i + 24)));
 
-        y0 = _mm256_fmadd_ps(x0, v8, y0);
+        y0 = _mm256_fmadd_ps(x0, v0, y0);
-        y1 = _mm256_fmadd_ps(x1, v8, y1);
+        y1 = _mm256_fmadd_ps(x1, v0, y1);
-        y2 = _mm256_fmadd_ps(x2, v8, y2);
+        y2 = _mm256_fmadd_ps(x2, v0, y2);
-        y3 = _mm256_fmadd_ps(x3, v8, y3);
+        y3 = _mm256_fmadd_ps(x3, v0, y3);
 
         _mm_storeu_si128((__m128i*)(y + i + 0 ), _mm256_cvtps_ph(y0, 0));
         _mm_storeu_si128((__m128i*)(y + i + 8 ), _mm256_cvtps_ph(y1, 0));
@@ -864,7 +1012,7 @@ inline static void ggml_vec_mad_f16(const int n, ggml_fp16_t * restrict y, ggml_
     // WASM SIMD 128-bit
     const int n16 = (n & ~15);
 
-    const v128_t v4 = wasm_f32x4_splat(v);
+    const v128_t v0 = wasm_f32x4_splat(v);
 
     v128_t x0, x1, x2, x3;
     v128_t y0, y1, y2, y3;
@@ -888,10 +1036,10 @@ inline static void ggml_vec_mad_f16(const int n, ggml_fp16_t * restrict y, ggml_
         y2 = wasm_v128_load(ty + 8);
         y3 = wasm_v128_load(ty + 12);
 
-        y0 = wasm_f32x4_add(y0, wasm_f32x4_mul(x0, v4));
+        y0 = wasm_f32x4_add(y0, wasm_f32x4_mul(x0, v0));
-        y1 = wasm_f32x4_add(y1, wasm_f32x4_mul(x1, v4));
+        y1 = wasm_f32x4_add(y1, wasm_f32x4_mul(x1, v0));
-        y2 = wasm_f32x4_add(y2, wasm_f32x4_mul(x2, v4));
+        y2 = wasm_f32x4_add(y2, wasm_f32x4_mul(x2, v0));
-        y3 = wasm_f32x4_add(y3, wasm_f32x4_mul(x3, v4));
+        y3 = wasm_f32x4_add(y3, wasm_f32x4_mul(x3, v0));
 
         wasm_v128_store(ty + 0, y0);
         wasm_v128_store(ty + 4, y1);
@@ -1309,8 +1457,6 @@ struct ggml_context * ggml_init(struct ggml_init_params params) {
 
     static bool first_time = true;
     if (first_time) {
-        ggml_mtl_init(); // TODO: fix this
-
         for (int i = 0; i < GGML_MAX_CONTEXTS; i++) {
             g_state.contexts[i].used = false;
         }
@@ -1466,104 +1612,6 @@ struct ggml_tensor * ggml_new_tensor_impl(
         /*.perf_cycles  =*/ 0,
         /*.perf_time_us =*/ 0,
         /*.data         =*/ data == NULL ? (void *)(result + 1) : data,
-        /*.id           =*/ -1,
-        /*.pad          =*/ { 0 },
-    };
-
-    ggml_assert_aligned(result->data);
-
-    for (int i = 0; i < n_dims; i++) {
-        result->ne[i] = ne[i];
-    }
-
-    result->nb[0] = GGML_TYPE_SIZE[type];
-    for (int i = 1; i < GGML_MAX_DIMS; i++) {
-        result->nb[i] = result->nb[i - 1]*result->ne[i - 1];
-    }
-
-    ctx->n_objects++;
-
-    return result;
-}
-
-struct ggml_tensor * ggml_new_tensor_mtl_impl(
-        struct ggml_context * ctx,
-        enum   ggml_type type,
-        int    n_dims,
-        const int* ne,
-        void*  data) {
-    // always insert objects at the end of the context's memory pool
-    struct ggml_object * obj_cur = ctx->objects_end;
-
-    const size_t cur_offset = obj_cur == NULL ? 0 : obj_cur->offset;
-    const size_t cur_size   = obj_cur == NULL ? 0 : obj_cur->size;
-    const size_t cur_end    = cur_offset + cur_size;
-
-    struct ggml_mtl_object obj_mtl;
-    {
-        assert(data == NULL); // TODO: in-place metal buffer, need page aligned memory
-        size_t size_needed_mtl = 0;
-        if (data == NULL) {
-            size_needed_mtl += GGML_TYPE_SIZE[type];
-            for (int i = 0; i < n_dims; i++) {
-                size_needed_mtl *= ne[i];
-            }
-        }
-
-        obj_mtl = ggml_mtl_alloc(size_needed_mtl);
-    }
-
-    size_t size_needed = 0;
-    size_needed += sizeof(struct ggml_tensor);
-
-    if (cur_end + size_needed + GGML_OBJECT_SIZE > ctx->mem_size) {
-        GGML_PRINT("%s: not enough space in the context's memory pool\n", __func__);
-        assert(false);
-        return NULL;
-    }
-
-    char * const mem_buffer = ctx->mem_buffer;
-
-    struct ggml_object * const obj_new = (struct ggml_object *)(mem_buffer + cur_end);
-
-    *obj_new = (struct ggml_object) {
-        .offset = cur_end + GGML_OBJECT_SIZE,
-        .size   = size_needed,
-        .next   = NULL,
-    };
-
-    if (obj_cur != NULL) {
-        obj_cur->next = obj_new;
-    } else {
-        // this is the first object in this context
-        ctx->objects_begin = obj_new;
-    }
-
-    ctx->objects_end = obj_new;
-
-    //GGML_PRINT_DEBUG("%s: inserted new object at %zu\n", __func__, cur_end);
-
-    struct ggml_tensor * const result = (struct ggml_tensor *)(mem_buffer + obj_new->offset);
-
-    ggml_assert_aligned(result);
-
-    *result = (struct ggml_tensor) {
-        /*.type         =*/ type,
-        /*.n_dims       =*/ n_dims,
-        /*.ne           =*/ { 1, 1, 1, 1 },
-        /*.nb           =*/ { 0, 0, 0, 0 },
-        /*.op           =*/ GGML_OP_NONE,
-        /*.is_param     =*/ false,
-        /*.grad         =*/ NULL,
-        /*.src0         =*/ NULL,
-        /*.src1         =*/ NULL,
-        /*.opt          =*/ { NULL },
-        /*.n_tasks      =*/ 0,
-        /*.perf_runs    =*/ 0,
-        /*.perf_cycles  =*/ 0,
-        /*.perf_time_us =*/ 0,
-        /*.data         =*/ obj_mtl.data,
-        /*.id           =*/ obj_mtl.id,
         /*.pad          =*/ { 0 },
     };
 
@@ -1591,14 +1639,6 @@ struct ggml_tensor * ggml_new_tensor(
     return ggml_new_tensor_impl(ctx, type, n_dims, ne, NULL);
 }
 
-struct ggml_tensor * ggml_new_tensor_mtl(
-        struct ggml_context * ctx,
-        enum   ggml_type type,
-        int    n_dims,
-        const int* ne) {
-    return ggml_new_tensor_mtl_impl(ctx, type, n_dims, ne, NULL);
-}
-
 struct ggml_tensor * ggml_new_tensor_1d(
         struct ggml_context * ctx,
         enum   ggml_type type,
@@ -1615,15 +1655,6 @@ struct ggml_tensor * ggml_new_tensor_2d(
     return ggml_new_tensor(ctx, type, 2, ne);
 }
 
-struct ggml_tensor * ggml_new_tensor_2d_mtl(
-        struct ggml_context * ctx,
-        enum   ggml_type type,
-        int    ne0,
-        int    ne1) {
-    const int ne[2] = { ne0, ne1 };
-    return ggml_new_tensor_mtl(ctx, type, 2, ne);
-}
-
 struct ggml_tensor * ggml_new_tensor_3d(
         struct ggml_context * ctx,
         enum   ggml_type type,
@@ -4462,11 +4493,8 @@ void ggml_compute_forward_mul_mat_f16_f32(
     // nb00 <  nb01 - src0 is transposed
     //   compute by src0 columns
 
-    // are we using Metal?
-    const bool is_mtl = src0->id >= 0;
-
 #if defined(GGML_USE_ACCELERATE) || defined(GGML_USE_OPENBLAS)
-    if (ggml_compute_forward_mul_mat_use_blas(src0, src1, dst) && !is_mtl) {
+    if (ggml_compute_forward_mul_mat_use_blas(src0, src1, dst)) {
         GGML_ASSERT(nb10 == sizeof(float));
 
         if (params->ith != 0) return;
@@ -4594,20 +4622,6 @@ void ggml_compute_forward_mul_mat_f16_f32(
 
         // parallelize by src0 rows using ggml_vec_dot_f32
 
-        if (is_mtl) {
-            assert(ne02 == 1);
-            assert(ne03 == 1);
-
-            if (params->ith == 0) {
-                printf("XXXXXXXXXXX src0->ne[0] = %d, src0->ne[1] = %d\n", src0->ne[0], src0->ne[1]);
-                printf("XXXXXXXXXXX src1->ne[0] = %d, src1->ne[1] = %d\n", src1->ne[0], src1->ne[1]);
-                struct ggml_mtl_object src0_mtl = { src0->id, src0->data };
-                ggml_fp16_t * src1_fp16 = params->wdata;
-                ggml_mtl_mul_mat_f16(NULL, src0_mtl, src1_fp16, dst->data, ne01, ne11, ne00);
-            }
-            return;
-        }
-
         // total rows in src0
         const int nr = ne01*ne02*ne03;
 

ggml.h

@@ -108,8 +108,7 @@ struct ggml_tensor {
     int64_t perf_time_us;
 
     void * data;
-    int32_t id; // TODO: mtl buffer id
+    char padding[8];
-    char pad[4];
 };
 
 // computation graph
@@ -174,12 +173,6 @@ struct ggml_tensor * ggml_new_tensor_2d(
         int    ne0,
         int    ne1);
 
-struct ggml_tensor * ggml_new_tensor_2d_mtl(
-        struct ggml_context * ctx,
-        enum   ggml_type type,
-        int    ne0,
-        int    ne1);
-
 struct ggml_tensor * ggml_new_tensor_3d(
         struct ggml_context * ctx,
         enum   ggml_type type,

whisper.cpp

@@ -788,10 +788,10 @@ static bool whisper_model_load(const std::string & fname, whisper_context & wctx
                 layer.mlp_ln_w = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_audio_state);
                 layer.mlp_ln_b = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_audio_state);
 
-                layer.mlp_0_w = ggml_new_tensor_2d_mtl(ctx, wtype,       n_audio_state, 4*n_audio_state); // offload to GPU
+                layer.mlp_0_w = ggml_new_tensor_2d(ctx, wtype,           n_audio_state, 4*n_audio_state);
                 layer.mlp_0_b = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 4*n_audio_state);
 
-                layer.mlp_1_w = ggml_new_tensor_2d_mtl(ctx, wtype,     4*n_audio_state, n_audio_state);   // offload to GPU
+                layer.mlp_1_w = ggml_new_tensor_2d(ctx, wtype,         4*n_audio_state, n_audio_state);
                 layer.mlp_1_b = ggml_new_tensor_1d(ctx, GGML_TYPE_F32,   n_audio_state);
 
                 layer.attn_ln_0_w = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, n_audio_state);
@@ -1342,7 +1342,7 @@ static bool whisper_encode(
             ggml_build_forward_expand(&gf, inpO);
             ggml_graph_compute       (ctxL, &gf);
 
-            ggml_graph_print(&gf);
+            //ggml_graph_print(&gf);
         }
 
         // TODO: this is a hack to have per-layer computation graphs - need to come up with something better
@@ -2339,7 +2339,6 @@ struct whisper_full_params whisper_full_default_params(enum whisper_sampling_str
                     /*.n_threads            =*/ std::min(4, (int32_t) std::thread::hardware_concurrency()),
                     /*.n_max_text_ctx       =*/ 16384,
                     /*.offset_ms            =*/ 0,
-                    /*.duration_ms          =*/ 0,
 
                     /*.translate            =*/ false,
                     /*.no_context           =*/ false,
@@ -2377,7 +2376,6 @@ struct whisper_full_params whisper_full_default_params(enum whisper_sampling_str
                     /*.n_threads            =*/ std::min(4, (int32_t) std::thread::hardware_concurrency()),
                     /*.n_max_text_ctx       =*/ 16384,
                     /*.offset_ms            =*/ 0,
-                    /*.duration_ms          =*/ 0,
 
                     /*.translate            =*/ false,
                     /*.no_context           =*/ false,
@@ -2498,12 +2496,11 @@ int whisper_full(
     }
 
     const int seek_start = params.offset_ms/10;
-    const int seek_end = seek_start + (params.duration_ms == 0 ? whisper_n_len(ctx) : params.duration_ms/10);
 
     // if length of spectrogram is less than 1s (100 samples), then return
     // basically don't process anything that is less than 1s
     // see issue #39: https://github.com/ggerganov/whisper.cpp/issues/39
-    if (seek_end < 100 + seek_start) {
+    if (whisper_n_len(ctx) < 100 + seek_start) {
         return 0;
     }
 
@@ -2536,7 +2533,7 @@ int whisper_full(
     // main loop
     int seek = seek_start;
     while (true) {
-        const int progress_cur = (100*(seek - seek_start))/(seek_end - seek_start);
+        int progress_cur = (100*seek)/whisper_n_len(ctx);
         while (progress_cur >= progress_prev + progress_step) {
             progress_prev += progress_step;
             if (params.print_progress) {
@@ -2544,7 +2541,7 @@ int whisper_full(
             }
         }
 
-        if (seek + 100 >= seek_end) {
+        if (seek + 100 >= whisper_n_len(ctx)) {
             break;
         }
 
@@ -2625,7 +2622,7 @@ int whisper_full(
                 // end of text token
                 if (token.id == whisper_token_eot(ctx)) {
                     if (result_len == 0) {
-                        if (seek + seek_delta + 100 >= seek_end) {
+                        if (seek + seek_delta + 100 >= whisper_n_len(ctx)) {
                             result_len = i + 1;
                         } else {
                             // TODO: figure out how to resolve this

whisper.h

@@ -186,8 +186,7 @@ extern "C" {
 
         int n_threads;
         int n_max_text_ctx;
-        int offset_ms;      // start offset in ms
+        int offset_ms;
-        int duration_ms;    // audio duration to process in ms
 
         bool translate;
         bool no_context;