From f78b9439ddb655f120d21be85090dc30a4806be7 Mon Sep 17 00:00:00 2001
From: manni07 <abend_schnack@yahoo.de>
Date: Mon, 2 Mar 2026 23:35:48 +0100
Subject: [PATCH] fix(HIGH-05): ane_eval() returns bool, add step_ok tracking

- stories_io.h: ane_eval() changed from void to bool; captures BOOL
  return from evaluateWithQoS:options:request:error: with NSError log
- train_large.m: step_ok tracks all 6 ane_eval() results per step;
  gradient update skipped if any ANE eval fails in that step

ref: docs/reports/security-audit-2026-03-02.md HIGH-05

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
---
 training/stories_io.h  | 336 +++++++++++++++++++++--------------------
 training/train_large.m |  17 ++-
 2 files changed, 181 insertions(+), 172 deletions(-)

diff --git a/training/stories_io.h b/training/stories_io.h
index 391a49e..62ca282 100644
--- a/training/stories_io.h
+++ b/training/stories_io.h
@@ -1,167 +1,171 @@
-// stories_io.h — IOSurface helpers, blob builders, NEON conversion
-#pragma once
-#include "stories_config.h"
-#include <arm_neon.h>
-
-static IOSurfaceRef make_surface(size_t bytes) {
-    return IOSurfaceCreate((__bridge CFDictionaryRef)@{
-        (id)kIOSurfaceWidth:@(bytes), (id)kIOSurfaceHeight:@1,
-        (id)kIOSurfaceBytesPerElement:@1, (id)kIOSurfaceBytesPerRow:@(bytes),
-        (id)kIOSurfaceAllocSize:@(bytes), (id)kIOSurfacePixelFormat:@0});
-}
-
-static NSData *build_blob(const float *w, int rows, int cols) {
-    size_t ws=(size_t)rows*cols*2, tot=128+ws;  // size_t prevents int overflow (CRIT-04)
-    uint8_t *b=(uint8_t*)calloc(tot,1);
-    if (!b) { fprintf(stderr, "build_blob: calloc(%zu) failed\n", tot); return nil; }
-    b[0]=1;b[4]=2;b[64]=0xEF;b[65]=0xBE;b[66]=0xAD;b[67]=0xDE;b[68]=1;
-    *(uint32_t*)(b+72)=(uint32_t)ws;*(uint32_t*)(b+80)=128;
-    _Float16 *fp16=(_Float16*)(b+128);
-    for(size_t i=0;i<(size_t)rows*cols;i++) fp16[i]=(_Float16)w[i];
-    return [NSData dataWithBytesNoCopy:b length:tot freeWhenDone:YES];
-}
-static NSData *build_blob_t(const float *w, int rows, int cols) {
-    size_t ws=(size_t)cols*rows*2, tot=128+ws;  // size_t prevents int overflow (CRIT-04)
-    uint8_t *b=(uint8_t*)calloc(tot,1);
-    if (!b) { fprintf(stderr, "build_blob_t: calloc(%zu) failed\n", tot); return nil; }
-    b[0]=1;b[4]=2;b[64]=0xEF;b[65]=0xBE;b[66]=0xAD;b[67]=0xDE;b[68]=1;
-    *(uint32_t*)(b+72)=(uint32_t)ws;*(uint32_t*)(b+80)=128;
-    _Float16 *fp16=(_Float16*)(b+128);
-    for(int i=0;i<rows;i++) for(int j=0;j<cols;j++) fp16[j*rows+i]=(_Float16)w[i*cols+j];
-    return [NSData dataWithBytesNoCopy:b length:tot freeWhenDone:YES];
-}
-static NSData *build_blob_fp16(_Float16 *d, int cnt) {
-    size_t ws=(size_t)cnt*2, tot=128+ws;  // size_t prevents int overflow (CRIT-04)
-    uint8_t *b=(uint8_t*)calloc(tot,1);
-    if (!b) { fprintf(stderr, "build_blob_fp16: calloc(%zu) failed\n", tot); return nil; }
-    b[0]=1;b[4]=2;b[64]=0xEF;b[65]=0xBE;b[66]=0xAD;b[67]=0xDE;b[68]=1;
-    *(uint32_t*)(b+72)=(uint32_t)ws;*(uint32_t*)(b+80)=128;
-    memcpy(b+128,d,ws);
-    return [NSData dataWithBytesNoCopy:b length:tot freeWhenDone:YES];
-}
-
-// MED-05: NEON alignment guarantee.
-// IOSurface base address is page-aligned (≥4096 bytes). Offset = ch_off*SEQ*sizeof(_Float16).
-// With SEQ%8==0, all offsets are multiples of 16 bytes → aligned for vld1q_f16/vst1q_f32.
-// Additionally, ARM64 handles unaligned NEON loads in hardware (unlike ARM32).
-_Static_assert(SEQ % 8 == 0,
-    "SEQ must be multiple of 8 to guarantee 16-byte alignment for NEON (MED-05)");
-
-// NEON vectorized conversion
-static void cvt_f16_f32(float *dst, const _Float16 *src, int n) {
-    int i = 0;
-    for (; i+7 < n; i += 8) {
-        float16x8_t h = vld1q_f16((const __fp16*)(src+i));
-        vst1q_f32(dst+i,   vcvt_f32_f16(vget_low_f16(h)));
-        vst1q_f32(dst+i+4, vcvt_f32_f16(vget_high_f16(h)));
-    }
-    for (; i < n; i++) dst[i] = (float)src[i];
-}
-static void cvt_f32_f16(_Float16 *dst, const float *src, int n) {
-    int i = 0;
-    for (; i+7 < n; i += 8) {
-        float16x8_t h = vcombine_f16(vcvt_f16_f32(vld1q_f32(src+i)),
-                                      vcvt_f16_f32(vld1q_f32(src+i+4)));
-        vst1q_f16((__fp16*)(dst+i), h);
-    }
-    for (; i < n; i++) dst[i] = (_Float16)src[i];
-}
-
-// IOSurface I/O (channel-first [C,S] layout)
-static void io_write_fp16(IOSurfaceRef s, const float *data, int channels, int sp) {
-    if (IOSurfaceLock(s, 0, NULL) != kIOReturnSuccess) {  // MED-01
-        fprintf(stderr, "IOSurfaceLock(write) failed — surface write skipped\n");
-        return;
-    }
-    cvt_f32_f16((_Float16*)IOSurfaceGetBaseAddress(s), data, channels * sp);
-    IOSurfaceUnlock(s, 0, NULL);
-}
-static void io_read_fp16(IOSurfaceRef s, float *data, int ch_off, int channels, int sp) {
-    if (IOSurfaceLock(s, kIOSurfaceLockReadOnly, NULL) != kIOReturnSuccess) {  // MED-01
-        fprintf(stderr, "IOSurfaceLock(read) failed — output read skipped\n");
-        return;
-    }
-    cvt_f16_f32(data, (_Float16*)IOSurfaceGetBaseAddress(s) + ch_off * sp, channels * sp);
-    IOSurfaceUnlock(s, kIOSurfaceLockReadOnly, NULL);
-}
-static void io_copy(IOSurfaceRef dst, int dst_ch, IOSurfaceRef src, int src_ch, int channels, int sp) {
-    if (IOSurfaceLock(dst, 0, NULL) != kIOReturnSuccess) {  // MED-01
-        fprintf(stderr, "IOSurfaceLock(copy dst) failed — copy skipped\n");
-        return;
-    }
-    if (IOSurfaceLock(src, kIOSurfaceLockReadOnly, NULL) != kIOReturnSuccess) {  // MED-01
-        fprintf(stderr, "IOSurfaceLock(copy src) failed — copy skipped\n");
-        IOSurfaceUnlock(dst, 0, NULL);
-        return;
-    }
-    memcpy((_Float16*)IOSurfaceGetBaseAddress(dst) + dst_ch*sp,
-           (_Float16*)IOSurfaceGetBaseAddress(src) + src_ch*sp,
-           channels * sp * sizeof(_Float16));
-    IOSurfaceUnlock(src, kIOSurfaceLockReadOnly, NULL);
-    IOSurfaceUnlock(dst, 0, NULL);
-}
-static void io_write_fp16_at(IOSurfaceRef s, int ch_off, const float *data, int channels, int sp) {
-    if (IOSurfaceLock(s, 0, NULL) != kIOReturnSuccess) {  // MED-01
-        fprintf(stderr, "IOSurfaceLock(write_at) failed — surface write skipped\n");
-        return;
-    }
-    cvt_f32_f16((_Float16*)IOSurfaceGetBaseAddress(s) + ch_off * sp, data, channels * sp);
-    IOSurfaceUnlock(s, 0, NULL);
-}
-
-// Kernel compile/eval
-static Kern *compile_kern_mil_w(NSString *mil, NSDictionary *weights, int ic_bytes, int oc_bytes) {
-    @autoreleasepool {
-    if (!g_ane_ok_large) { printf("  [compile] ANE not available\n"); return NULL; }  // CRIT-01/02
-    NSData *md = [mil dataUsingEncoding:NSUTF8StringEncoding];
-    id desc = ((id(*)(Class,SEL,id,id,id))objc_msgSend)(g_D, @selector(modelWithMILText:weights:optionsPlist:), md, weights, nil);
-    if (!desc) { printf("  [compile] desc=NULL\n"); return NULL; }
-    id mdl = ((id(*)(Class,SEL,id))objc_msgSend)(g_I, @selector(inMemoryModelWithDescriptor:), desc);
-    if (!mdl) { printf("  [compile] mdl=NULL\n"); return NULL; }  // CRIT-02
-    id hx = ((id(*)(id,SEL))objc_msgSend)(mdl, @selector(hexStringIdentifier));
-    // MED-02: pid + atomic sequence counter make the directory unique per process and
-    // per call, preventing TOCTOU conflicts when two instances compile the same model.
-    int seq = __sync_fetch_and_add(&g_compile_seq, 1);
-    NSString *td = [NSTemporaryDirectory() stringByAppendingPathComponent:
-        [NSString stringWithFormat:@"ANE_%d_%d_%@", getpid(), seq, hx]];
-    [[NSFileManager defaultManager] createDirectoryAtPath:[td stringByAppendingPathComponent:@"weights"] withIntermediateDirectories:YES attributes:nil error:nil];
-    [md writeToFile:[td stringByAppendingPathComponent:@"model.mil"] atomically:YES];
-    for (NSString *path in weights) {
-        NSString *rel = [path stringByReplacingOccurrencesOfString:@"@model_path/" withString:@""];
-        [weights[path][@"data"] writeToFile:[td stringByAppendingPathComponent:rel] atomically:YES];
-    }
-    NSError *e = nil;
-    if (!((BOOL(*)(id,SEL,unsigned int,id,NSError**))objc_msgSend)(mdl, @selector(compileWithQoS:options:error:), 21, @{}, &e)) {
-        printf("  [compile] FAIL: %s\n", e ? [[e description] UTF8String] : "no error"); return NULL;
-    }
-    if (!((BOOL(*)(id,SEL,unsigned int,id,NSError**))objc_msgSend)(mdl, @selector(loadWithQoS:options:error:), 21, @{}, &e)) {
-        printf("  [compile] load FAIL\n"); return NULL;
-    }
-    __sync_fetch_and_add(&g_compile_count, 1);
-    Kern *k = (Kern*)calloc(1, sizeof(Kern));
+// stories_io.h — IOSurface helpers, blob builders, NEON conversion
+#pragma once
+#include "stories_config.h"
+#include <arm_neon.h>
+
+static IOSurfaceRef make_surface(size_t bytes) {
+    return IOSurfaceCreate((__bridge CFDictionaryRef)@{
+        (id)kIOSurfaceWidth:@(bytes), (id)kIOSurfaceHeight:@1,
+        (id)kIOSurfaceBytesPerElement:@1, (id)kIOSurfaceBytesPerRow:@(bytes),
+        (id)kIOSurfaceAllocSize:@(bytes), (id)kIOSurfacePixelFormat:@0});
+}
+
+static NSData *build_blob(const float *w, int rows, int cols) {
+    size_t ws=(size_t)rows*cols*2, tot=128+ws;  // size_t prevents int overflow (CRIT-04)
+    uint8_t *b=(uint8_t*)calloc(tot,1);
+    if (!b) { fprintf(stderr, "build_blob: calloc(%zu) failed\n", tot); return nil; }
+    b[0]=1;b[4]=2;b[64]=0xEF;b[65]=0xBE;b[66]=0xAD;b[67]=0xDE;b[68]=1;
+    *(uint32_t*)(b+72)=(uint32_t)ws;*(uint32_t*)(b+80)=128;
+    _Float16 *fp16=(_Float16*)(b+128);
+    for(size_t i=0;i<(size_t)rows*cols;i++) fp16[i]=(_Float16)w[i];
+    return [NSData dataWithBytesNoCopy:b length:tot freeWhenDone:YES];
+}
+static NSData *build_blob_t(const float *w, int rows, int cols) {
+    size_t ws=(size_t)cols*rows*2, tot=128+ws;  // size_t prevents int overflow (CRIT-04)
+    uint8_t *b=(uint8_t*)calloc(tot,1);
+    if (!b) { fprintf(stderr, "build_blob_t: calloc(%zu) failed\n", tot); return nil; }
+    b[0]=1;b[4]=2;b[64]=0xEF;b[65]=0xBE;b[66]=0xAD;b[67]=0xDE;b[68]=1;
+    *(uint32_t*)(b+72)=(uint32_t)ws;*(uint32_t*)(b+80)=128;
+    _Float16 *fp16=(_Float16*)(b+128);
+    for(int i=0;i<rows;i++) for(int j=0;j<cols;j++) fp16[j*rows+i]=(_Float16)w[i*cols+j];
+    return [NSData dataWithBytesNoCopy:b length:tot freeWhenDone:YES];
+}
+static NSData *build_blob_fp16(_Float16 *d, int cnt) {
+    size_t ws=(size_t)cnt*2, tot=128+ws;  // size_t prevents int overflow (CRIT-04)
+    uint8_t *b=(uint8_t*)calloc(tot,1);
+    if (!b) { fprintf(stderr, "build_blob_fp16: calloc(%zu) failed\n", tot); return nil; }
+    b[0]=1;b[4]=2;b[64]=0xEF;b[65]=0xBE;b[66]=0xAD;b[67]=0xDE;b[68]=1;
+    *(uint32_t*)(b+72)=(uint32_t)ws;*(uint32_t*)(b+80)=128;
+    memcpy(b+128,d,ws);
+    return [NSData dataWithBytesNoCopy:b length:tot freeWhenDone:YES];
+}
+
+// MED-05: NEON alignment guarantee.
+// IOSurface base address is page-aligned (≥4096 bytes). Offset = ch_off*SEQ*sizeof(_Float16).
+// With SEQ%8==0, all offsets are multiples of 16 bytes → aligned for vld1q_f16/vst1q_f32.
+// Additionally, ARM64 handles unaligned NEON loads in hardware (unlike ARM32).
+_Static_assert(SEQ % 8 == 0,
+    "SEQ must be multiple of 8 to guarantee 16-byte alignment for NEON (MED-05)");
+
+// NEON vectorized conversion
+static void cvt_f16_f32(float *dst, const _Float16 *src, int n) {
+    int i = 0;
+    for (; i+7 < n; i += 8) {
+        float16x8_t h = vld1q_f16((const __fp16*)(src+i));
+        vst1q_f32(dst+i,   vcvt_f32_f16(vget_low_f16(h)));
+        vst1q_f32(dst+i+4, vcvt_f32_f16(vget_high_f16(h)));
+    }
+    for (; i < n; i++) dst[i] = (float)src[i];
+}
+static void cvt_f32_f16(_Float16 *dst, const float *src, int n) {
+    int i = 0;
+    for (; i+7 < n; i += 8) {
+        float16x8_t h = vcombine_f16(vcvt_f16_f32(vld1q_f32(src+i)),
+                                      vcvt_f16_f32(vld1q_f32(src+i+4)));
+        vst1q_f16((__fp16*)(dst+i), h);
+    }
+    for (; i < n; i++) dst[i] = (_Float16)src[i];
+}
+
+// IOSurface I/O (channel-first [C,S] layout)
+static void io_write_fp16(IOSurfaceRef s, const float *data, int channels, int sp) {
+    if (IOSurfaceLock(s, 0, NULL) != kIOReturnSuccess) {  // MED-01
+        fprintf(stderr, "IOSurfaceLock(write) failed — surface write skipped\n");
+        return;
+    }
+    cvt_f32_f16((_Float16*)IOSurfaceGetBaseAddress(s), data, channels * sp);
+    IOSurfaceUnlock(s, 0, NULL);
+}
+static void io_read_fp16(IOSurfaceRef s, float *data, int ch_off, int channels, int sp) {
+    if (IOSurfaceLock(s, kIOSurfaceLockReadOnly, NULL) != kIOReturnSuccess) {  // MED-01
+        fprintf(stderr, "IOSurfaceLock(read) failed — output read skipped\n");
+        return;
+    }
+    cvt_f16_f32(data, (_Float16*)IOSurfaceGetBaseAddress(s) + ch_off * sp, channels * sp);
+    IOSurfaceUnlock(s, kIOSurfaceLockReadOnly, NULL);
+}
+static void io_copy(IOSurfaceRef dst, int dst_ch, IOSurfaceRef src, int src_ch, int channels, int sp) {
+    if (IOSurfaceLock(dst, 0, NULL) != kIOReturnSuccess) {  // MED-01
+        fprintf(stderr, "IOSurfaceLock(copy dst) failed — copy skipped\n");
+        return;
+    }
+    if (IOSurfaceLock(src, kIOSurfaceLockReadOnly, NULL) != kIOReturnSuccess) {  // MED-01
+        fprintf(stderr, "IOSurfaceLock(copy src) failed — copy skipped\n");
+        IOSurfaceUnlock(dst, 0, NULL);
+        return;
+    }
+    memcpy((_Float16*)IOSurfaceGetBaseAddress(dst) + dst_ch*sp,
+           (_Float16*)IOSurfaceGetBaseAddress(src) + src_ch*sp,
+           channels * sp * sizeof(_Float16));
+    IOSurfaceUnlock(src, kIOSurfaceLockReadOnly, NULL);
+    IOSurfaceUnlock(dst, 0, NULL);
+}
+static void io_write_fp16_at(IOSurfaceRef s, int ch_off, const float *data, int channels, int sp) {
+    if (IOSurfaceLock(s, 0, NULL) != kIOReturnSuccess) {  // MED-01
+        fprintf(stderr, "IOSurfaceLock(write_at) failed — surface write skipped\n");
+        return;
+    }
+    cvt_f32_f16((_Float16*)IOSurfaceGetBaseAddress(s) + ch_off * sp, data, channels * sp);
+    IOSurfaceUnlock(s, 0, NULL);
+}
+
+// Kernel compile/eval
+static Kern *compile_kern_mil_w(NSString *mil, NSDictionary *weights, int ic_bytes, int oc_bytes) {
+    @autoreleasepool {
+    if (!g_ane_ok_large) { printf("  [compile] ANE not available\n"); return NULL; }  // CRIT-01/02
+    NSData *md = [mil dataUsingEncoding:NSUTF8StringEncoding];
+    id desc = ((id(*)(Class,SEL,id,id,id))objc_msgSend)(g_D, @selector(modelWithMILText:weights:optionsPlist:), md, weights, nil);
+    if (!desc) { printf("  [compile] desc=NULL\n"); return NULL; }
+    id mdl = ((id(*)(Class,SEL,id))objc_msgSend)(g_I, @selector(inMemoryModelWithDescriptor:), desc);
+    if (!mdl) { printf("  [compile] mdl=NULL\n"); return NULL; }  // CRIT-02
+    id hx = ((id(*)(id,SEL))objc_msgSend)(mdl, @selector(hexStringIdentifier));
+    // MED-02: pid + atomic sequence counter make the directory unique per process and
+    // per call, preventing TOCTOU conflicts when two instances compile the same model.
+    int seq = __sync_fetch_and_add(&g_compile_seq, 1);
+    NSString *td = [NSTemporaryDirectory() stringByAppendingPathComponent:
+        [NSString stringWithFormat:@"ANE_%d_%d_%@", getpid(), seq, hx]];
+    [[NSFileManager defaultManager] createDirectoryAtPath:[td stringByAppendingPathComponent:@"weights"] withIntermediateDirectories:YES attributes:nil error:nil];
+    [md writeToFile:[td stringByAppendingPathComponent:@"model.mil"] atomically:YES];
+    for (NSString *path in weights) {
+        NSString *rel = [path stringByReplacingOccurrencesOfString:@"@model_path/" withString:@""];
+        [weights[path][@"data"] writeToFile:[td stringByAppendingPathComponent:rel] atomically:YES];
+    }
+    NSError *e = nil;
+    if (!((BOOL(*)(id,SEL,unsigned int,id,NSError**))objc_msgSend)(mdl, @selector(compileWithQoS:options:error:), 21, @{}, &e)) {
+        printf("  [compile] FAIL: %s\n", e ? [[e description] UTF8String] : "no error"); return NULL;
+    }
+    if (!((BOOL(*)(id,SEL,unsigned int,id,NSError**))objc_msgSend)(mdl, @selector(loadWithQoS:options:error:), 21, @{}, &e)) {
+        printf("  [compile] load FAIL\n"); return NULL;
+    }
+    __sync_fetch_and_add(&g_compile_count, 1);
+    Kern *k = (Kern*)calloc(1, sizeof(Kern));
     if (!k) { fprintf(stderr, "OOM: calloc(Kern)\n"); abort(); }  // HIGH-04
-    k->model = (void*)CFBridgingRetain(mdl);
-    k->ioIn = make_surface(ic_bytes);
-    k->ioOut = make_surface(oc_bytes);
-    id wI = ((id(*)(Class,SEL,IOSurfaceRef))objc_msgSend)(g_AIO, @selector(objectWithIOSurface:), k->ioIn);
-    id wO = ((id(*)(Class,SEL,IOSurfaceRef))objc_msgSend)(g_AIO, @selector(objectWithIOSurface:), k->ioOut);
-    k->request = (void*)CFBridgingRetain(((id(*)(Class,SEL,id,id,id,id,id,id,id))objc_msgSend)(g_AR,
-        @selector(requestWithInputs:inputIndices:outputs:outputIndices:weightsBuffer:perfStats:procedureIndex:),
-        @[wI], @[@0], @[wO], @[@0], nil, nil, @0));
-    k->tmpDir = (void*)CFBridgingRetain(td);
-    return k;
-    }
-}
-static void free_kern(Kern *k) {
-    if (!k) return;
-    id mdl = (__bridge id)k->model; NSError *e = nil;
-    ((BOOL(*)(id,SEL,unsigned int,NSError**))objc_msgSend)(mdl, @selector(unloadWithQoS:error:), 21, &e);
-    CFRelease(k->ioIn); CFRelease(k->ioOut);
-    [[NSFileManager defaultManager] removeItemAtPath:(__bridge id)k->tmpDir error:nil];
-    CFRelease(k->model); CFRelease(k->request); CFRelease(k->tmpDir);
-    free(k);
-}
-static void ane_eval(Kern *k) {
-    id mdl = (__bridge id)k->model; id req = (__bridge id)k->request; NSError *e = nil;
-    ((BOOL(*)(id,SEL,unsigned int,id,id,NSError**))objc_msgSend)(mdl, @selector(evaluateWithQoS:options:request:error:), 21, @{}, req, &e);
-}
+    k->model = (void*)CFBridgingRetain(mdl);
+    k->ioIn = make_surface(ic_bytes);
+    k->ioOut = make_surface(oc_bytes);
+    id wI = ((id(*)(Class,SEL,IOSurfaceRef))objc_msgSend)(g_AIO, @selector(objectWithIOSurface:), k->ioIn);
+    id wO = ((id(*)(Class,SEL,IOSurfaceRef))objc_msgSend)(g_AIO, @selector(objectWithIOSurface:), k->ioOut);
+    k->request = (void*)CFBridgingRetain(((id(*)(Class,SEL,id,id,id,id,id,id,id))objc_msgSend)(g_AR,
+        @selector(requestWithInputs:inputIndices:outputs:outputIndices:weightsBuffer:perfStats:procedureIndex:),
+        @[wI], @[@0], @[wO], @[@0], nil, nil, @0));
+    k->tmpDir = (void*)CFBridgingRetain(td);
+    return k;
+    }
+}
+static void free_kern(Kern *k) {
+    if (!k) return;
+    id mdl = (__bridge id)k->model; NSError *e = nil;
+    ((BOOL(*)(id,SEL,unsigned int,NSError**))objc_msgSend)(mdl, @selector(unloadWithQoS:error:), 21, &e);
+    CFRelease(k->ioIn); CFRelease(k->ioOut);
+    [[NSFileManager defaultManager] removeItemAtPath:(__bridge id)k->tmpDir error:nil];
+    CFRelease(k->model); CFRelease(k->request); CFRelease(k->tmpDir);
+    free(k);
+}
+static bool ane_eval(Kern *k) {  // HIGH-05: was void
+    id mdl = (__bridge id)k->model; id req = (__bridge id)k->request; NSError *e = nil;
+    BOOL ok = ((BOOL(*)(id,SEL,unsigned int,id,id,NSError**))objc_msgSend)(
+        mdl, @selector(evaluateWithQoS:options:request:error:), 21, @{}, req, &e);
+    if (!ok) fprintf(stderr, "  [ane_eval] FAILED: %s\n",
+                     e ? [[e description] UTF8String] : "unknown error");
+    return (bool)ok;
+}
diff --git a/training/train_large.m b/training/train_large.m
index cc1ac0a..c9f31fa 100644
--- a/training/train_large.m
+++ b/training/train_large.m
@@ -408,6 +408,7 @@ int main(int argc, char *argv[]) {
             uint64_t tt = mach_absolute_time();
             double t_ane=0,t_io=0,t_elem=0,t_rms=0,t_cblas_wait=0,t_cls=0;
 
+            bool step_ok = true;  // HIGH-05: track ANE eval success across all acc steps
             for (int a=0; a<ACCUM_STEPS && step<total_steps; a++, step++) {
                 uint64_t t0,t1;
                 // Sample random position in token data
@@ -433,7 +434,7 @@ int main(int argc, char *argv[]) {
                     t1=mach_absolute_time(); t_cblas_wait+=tb_ms(t1-t0); t0=t1;
                     io_write_fp16(kern[L].fwdAttn->ioIn, x_cur, DIM, SEQ);
                     t1=mach_absolute_time(); t_io+=tb_ms(t1-t0); t0=t1;
-                    ane_eval(kern[L].fwdAttn);
+                    step_ok &= ane_eval(kern[L].fwdAttn);
                     t1=mach_absolute_time(); t_ane+=tb_ms(t1-t0); t0=t1;
                     io_read_fp16(kern[L].fwdAttn->ioOut, ac->o_out,    0,     DIM, SEQ);
                     io_read_fp16(kern[L].fwdAttn->ioOut, ac->attn_out, 4*DIM, DIM, SEQ);
@@ -446,7 +447,7 @@ int main(int argc, char *argv[]) {
                     // FFN forward
                     io_write_fp16(kern[L].fwdFFN->ioIn, ac->x2, DIM, SEQ);
                     t1=mach_absolute_time(); t_io+=tb_ms(t1-t0); t0=t1;
-                    ane_eval(kern[L].fwdFFN);
+                    step_ok &= ane_eval(kern[L].fwdFFN);
                     t1=mach_absolute_time(); t_ane+=tb_ms(t1-t0); t0=t1;
                     io_read_fp16(kern[L].fwdFFN->ioOut, ac->ffn_out,  0,              DIM,    SEQ);
                     io_read_fp16(kern[L].fwdFFN->ioOut, ac->h1,       DIM,            HIDDEN, SEQ);
@@ -509,7 +510,7 @@ int main(int argc, char *argv[]) {
                     // FFN backward (ANE)
                     io_write_fp16_at(kern[L].ffnBwd->ioIn, 0, dffn, DIM, SEQ);
                     io_copy(kern[L].ffnBwd->ioIn, DIM, kern[L].fwdFFN->ioOut, DIM, 2*HIDDEN, SEQ);
-                    ane_eval(kern[L].ffnBwd);
+                    step_ok &= ane_eval(kern[L].ffnBwd);
                     io_read_fp16(kern[L].ffnBwd->ioOut, dx_ffn, 0,           DIM,    SEQ);
                     io_read_fp16(kern[L].ffnBwd->ioOut, dh1,    DIM,         HIDDEN, SEQ);
                     io_read_fp16(kern[L].ffnBwd->ioOut, dh3,    DIM+HIDDEN,  HIDDEN, SEQ);
@@ -549,10 +550,10 @@ int main(int argc, char *argv[]) {
                     // SDPA backward (ANE)
                     io_copy(kern[L].sdpaBwd1->ioIn, 0, kern[L].fwdAttn->ioOut, DIM, 3*DIM, SEQ);
                     io_write_fp16_at(kern[L].sdpaBwd1->ioIn, 3*DIM, dx2, DIM, SEQ);
-                    ane_eval(kern[L].sdpaBwd1);
+                    step_ok &= ane_eval(kern[L].sdpaBwd1);
                     io_copy(sdpaBwd2[L]->ioIn, 0, kern[L].sdpaBwd1->ioOut, DIM, 2*SCORE_CH, SEQ);
                     io_copy(sdpaBwd2[L]->ioIn, 2*SCORE_CH, kern[L].fwdAttn->ioOut, DIM, 2*DIM, SEQ);
-                    ane_eval(sdpaBwd2[L]);
+                    step_ok &= ane_eval(sdpaBwd2[L]);
 
                     io_read_fp16(sdpaBwd2[L]->ioOut, dq, 0,   DIM, SEQ);
                     io_read_fp16(sdpaBwd2[L]->ioOut, dk, DIM,  DIM, SEQ);
@@ -576,7 +577,7 @@ int main(int argc, char *argv[]) {
                     // QKV backward (ANE)
                     io_copy(kern[L].qkvBwd->ioIn, 0, sdpaBwd2[L]->ioOut, 0, 2*DIM, SEQ);
                     io_copy(kern[L].qkvBwd->ioIn, 2*DIM, kern[L].sdpaBwd1->ioOut, 0, DIM, SEQ);
-                    ane_eval(kern[L].qkvBwd);
+                    step_ok &= ane_eval(kern[L].qkvBwd);
                     io_read_fp16(kern[L].qkvBwd->ioOut, dx_attn, 0, DIM, SEQ);
 
                     // RMSNorm1 backward (using saved layer input)
@@ -632,6 +633,10 @@ int main(int argc, char *argv[]) {
                     "\"compiles\":%d}\n",
                     step, loss, step_ane, step_io, step_cls, step_elem, step_rms, step_cbw, g_compile_count);
             }
+            if (!step_ok) {  // HIGH-05: skip gradient update if any ANE eval failed
+                fprintf(stderr, "  Step %d: ANE eval error — gradient update skipped\n", step);
+                continue;  // skip to next iteration of outer while (step < total_steps)
+            }
             double tms = tb_ms(mach_absolute_time() - tt);
             total_train_ms += tms;
             total_steps_done += steps_batch;