mirror of https://github.com/maderix/ANE.git
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>
This commit is contained in:
manni07
parent
42eae54975
commit
f78b9439dd
|
|
@ -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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
|
|
|||
Loading…
Reference in New Issue