|
|
#import "main-mtl.h" |
|
|
|
|
|
#import "ggml/ggml.h" |
|
|
|
|
|
#import <Foundation/Foundation.h> |
|
|
#import <Metal/Metal.h> |
|
|
#import <MetalPerformanceShaders/MetalPerformanceShaders.h> |
|
|
|
|
|
// TODO: couldn't get this to work |
|
|
//#define GGML_MTL_HEAP |
|
|
|
|
|
struct ggml_mtl_context { |
|
|
struct ggml_context * ctx_data; |
|
|
struct ggml_context * ctx_eval; |
|
|
struct ggml_context * ctx_work; |
|
|
|
|
|
id<MTLDevice> device; |
|
|
id<MTLCommandQueue> queue; |
|
|
id<MTLLibrary> library; |
|
|
|
|
|
#ifdef GGML_MTL_HEAP |
|
|
id<MTLHeap> heap_data; |
|
|
id<MTLHeap> heap_eval; |
|
|
#else |
|
|
id<MTLBuffer> buffer_data; |
|
|
id<MTLBuffer> buffer_eval; |
|
|
#endif |
|
|
|
|
|
id<MTLBuffer> out; |
|
|
|
|
|
// custom kernels |
|
|
id<MTLFunction> function_add; |
|
|
id<MTLComputePipelineState> pipeline_add; |
|
|
|
|
|
id<MTLFunction> function_relu; |
|
|
id<MTLComputePipelineState> pipeline_relu; |
|
|
|
|
|
id<MTLFunction> function_soft_max; |
|
|
id<MTLComputePipelineState> pipeline_soft_max; |
|
|
}; |
|
|
|
|
|
// MSL code |
|
|
NSString * const msl_library_mnist = @"\ |
|
|
#include <metal_stdlib> \n\ |
|
|
using namespace metal; \n\ |
|
|
\n\ |
|
|
#define MAX(x, y) ((x) > (y) ? (x) : (y)) \n\ |
|
|
\n\ |
|
|
constant int k_digits [[function_constant(0)]]; \n\ |
|
|
\n\ |
|
|
kernel void kernel_add( \n\ |
|
|
device const float * src0, \n\ |
|
|
device const float * src1, \n\ |
|
|
device float * dst, \n\ |
|
|
uint gid[[thread_position_in_grid]]) { \n\ |
|
|
dst[gid] = src0[gid] + src1[gid]; \n\ |
|
|
} \n\ |
|
|
\n\ |
|
|
kernel void kernel_relu( \n\ |
|
|
device const float * src, \n\ |
|
|
device float * dst, \n\ |
|
|
uint gid[[thread_position_in_grid]]) { \n\ |
|
|
dst[gid] = max(0.0f, src[gid]); \n\ |
|
|
} \n\ |
|
|
\n\ |
|
|
kernel void kernel_soft_max( \n\ |
|
|
device const float * src, \n\ |
|
|
device float * dst, \n\ |
|
|
uint gid[[thread_position_in_grid]]) { \n\ |
|
|
float max = 0.0f; \n\ |
|
|
for (int i = 0; i < k_digits; i++) { \n\ |
|
|
max = MAX(max, src[i]); \n\ |
|
|
} \n\ |
|
|
float sum = 0.0f; \n\ |
|
|
for (int i = 0; i < k_digits; i++) { \n\ |
|
|
dst[i] = exp(src[i] - max); \n\ |
|
|
sum += dst[i]; \n\ |
|
|
} \n\ |
|
|
for (int i = 0; i < k_digits; i++) { \n\ |
|
|
dst[i] /= sum; \n\ |
|
|
} \n\ |
|
|
} \n\ |
|
|
"; |
|
|
|
|
|
struct ggml_mtl_context * mnist_mtl_init( |
|
|
struct ggml_context * ctx_data, |
|
|
struct ggml_context * ctx_eval, |
|
|
struct ggml_context * ctx_work, |
|
|
struct ggml_cgraph * gf) { |
|
|
fprintf(stderr, "%s: allocating\n", __func__); |
|
|
|
|
|
struct ggml_mtl_context * ctx = malloc(sizeof(struct ggml_mtl_context)); |
|
|
|
|
|
ctx->ctx_data = ctx_data; |
|
|
ctx->ctx_eval = ctx_eval; |
|
|
ctx->ctx_work = ctx_work; |
|
|
|
|
|
ctx->device = MTLCreateSystemDefaultDevice(); |
|
|
ctx->queue = [ctx->device newCommandQueue]; |
|
|
|
|
|
// determine if we can use MPS |
|
|
if (MPSSupportsMTLDevice(ctx->device)) { |
|
|
fprintf(stderr, "%s: using MPS\n", __func__); |
|
|
} else { |
|
|
fprintf(stderr, "%s: not using MPS\n", __func__); |
|
|
GGML_ASSERT(false && "MPS not supported"); |
|
|
} |
|
|
|
|
|
// compile from source string and show compile log |
|
|
{ |
|
|
NSError * error = nil; |
|
|
ctx->library = [ctx->device newLibraryWithSource:msl_library_mnist options:nil error:&error]; |
|
|
if (error) { |
|
|
fprintf(stderr, "%s: error: %s\n", __func__, [[error description] UTF8String]); |
|
|
exit(1); |
|
|
} |
|
|
} |
|
|
|
|
|
// load kernels |
|
|
{ |
|
|
const int k_digits = ggml_graph_get_tensor(gf, "probs")->ne[0]; |
|
|
|
|
|
MTLFunctionConstantValues * constants = [MTLFunctionConstantValues new]; |
|
|
[constants setConstantValue:&k_digits type:MTLDataTypeInt withName:@"k_digits"]; |
|
|
|
|
|
ctx->function_add = [ctx->library newFunctionWithName:@"kernel_add"]; |
|
|
ctx->pipeline_add = [ctx->device newComputePipelineStateWithFunction:ctx->function_add error:nil]; |
|
|
fprintf(stderr, "%s: loaded kernel_add: %p\n", __func__, (void *) ctx->pipeline_add); |
|
|
|
|
|
ctx->function_relu = [ctx->library newFunctionWithName:@"kernel_relu"]; |
|
|
ctx->pipeline_relu = [ctx->device newComputePipelineStateWithFunction:ctx->function_relu error:nil]; |
|
|
fprintf(stderr, "%s: loaded kernel_relu: %p\n", __func__, (void *) ctx->pipeline_relu); |
|
|
|
|
|
ctx->function_soft_max = [ctx->library newFunctionWithName:@"kernel_soft_max" constantValues:constants error:nil]; |
|
|
ctx->pipeline_soft_max = [ctx->device newComputePipelineStateWithFunction:ctx->function_soft_max error:nil]; |
|
|
fprintf(stderr, "%s: loaded kernel_soft_max: %p\n", __func__, (void *) ctx->pipeline_soft_max); |
|
|
} |
|
|
|
|
|
#ifdef GGML_MTL_HEAP |
|
|
// MTLHeap approach |
|
|
|
|
|
// pin ctx_data memory to GPU |
|
|
// use MTLStorageModeShared to allow us to initialize the weights from the CPU |
|
|
// TODO: how to use MTLStorageModeManaged? |
|
|
// TODO: see if we can avoid this copy somehow |
|
|
{ |
|
|
const void * mem_buffer = ggml_get_mem_buffer(ctx_data); |
|
|
const size_t mem_size = ggml_get_mem_size(ctx_data); |
|
|
|
|
|
MTLHeapDescriptor * heap_desc = [MTLHeapDescriptor new]; |
|
|
heap_desc.storageMode = MTLStorageModeShared; |
|
|
heap_desc.size = mem_size; |
|
|
|
|
|
printf("heap_desc.size = %zu\n", mem_size); |
|
|
|
|
|
ctx->heap_data = [ctx->device newHeapWithDescriptor:heap_desc]; |
|
|
[ctx->heap_data setPurgeableState:MTLPurgeableStateNonVolatile]; // TODO: is this needed? |
|
|
ctx->heap_data.label = @"heap_data"; |
|
|
|
|
|
printf("ctx->heap_data.size = %zu\n", [ctx->heap_data size]); |
|
|
|
|
|
id<MTLBuffer> buffer = [ctx->heap_data newBufferWithLength:mem_size options:MTLResourceStorageModeShared]; |
|
|
if (!buffer) { |
|
|
fprintf(stderr, "%s: error: failed to allocate buffer\n", __func__); |
|
|
exit(1); |
|
|
} |
|
|
|
|
|
// copy data from CPU to GPU |
|
|
memcpy([buffer contents], mem_buffer, mem_size); |
|
|
|
|
|
fprintf(stderr, "%s: allocated data heap, size = %zu\n", __func__, mem_size); |
|
|
} |
|
|
|
|
|
// pin ctx_eval memory to GPU |
|
|
// this heap will be used for the intermediate results of the evaluation |
|
|
{ |
|
|
const size_t mem_size = ggml_get_mem_size(ctx_eval); |
|
|
|
|
|
MTLHeapDescriptor * heap_desc = [MTLHeapDescriptor new]; |
|
|
heap_desc.storageMode = MTLStorageModePrivate; // GPU only |
|
|
heap_desc.size = mem_size; |
|
|
|
|
|
ctx->heap_eval = [ctx->device newHeapWithDescriptor:heap_desc]; |
|
|
[ctx->heap_eval setPurgeableState:MTLPurgeableStateNonVolatile]; // TODO: is this needed? |
|
|
|
|
|
fprintf(stderr, "%s: allocated eval heap, size = %zu\n", __func__, mem_size); |
|
|
} |
|
|
#else |
|
|
// MTLBuffer approach |
|
|
|
|
|
// pin ctx_data memory to GPU |
|
|
// use MTLStorageModeShared to allow us to initialize the weights from the CPU |
|
|
// TODO: how to use MTLStorageModeManaged? |
|
|
// TODO: see if we can avoid this copy somehow |
|
|
{ |
|
|
const void * mem_buffer = ggml_get_mem_buffer(ctx_data); |
|
|
const size_t mem_size = ggml_get_mem_size(ctx_data); |
|
|
|
|
|
ctx->buffer_data = [ctx->device newBufferWithBytes:mem_buffer length:mem_size options:MTLResourceStorageModeShared]; |
|
|
|
|
|
fprintf(stderr, "%s: allocated data buffer, size = %zu\n", __func__, mem_size); |
|
|
} |
|
|
|
|
|
// pin ctx_eval memory to GPU |
|
|
// this buffer will be used for the intermediate results of the evaluation |
|
|
{ |
|
|
const size_t mem_size = ggml_get_mem_size(ctx_eval); |
|
|
|
|
|
ctx->buffer_eval = [ctx->device newBufferWithLength:mem_size options:MTLResourceStorageModePrivate]; |
|
|
|
|
|
fprintf(stderr, "%s: allocated eval buffer, size = %zu\n", __func__, mem_size); |
|
|
} |
|
|
#endif |
|
|
|
|
|
// allocate buffer for result extraction |
|
|
{ |
|
|
const size_t mem_size = ggml_nbytes(gf->nodes[gf->n_nodes - 1]); |
|
|
|
|
|
ctx->out = [ctx->device newBufferWithLength:mem_size options:MTLResourceStorageModeShared]; |
|
|
|
|
|
fprintf(stderr, "%s: allocated out buffer, size = %zu\n", __func__, mem_size); |
|
|
} |
|
|
|
|
|
return ctx; |
|
|
} |
|
|
|
|
|
void mnist_mtl_free(struct ggml_mtl_context * ctx) { |
|
|
fprintf(stderr, "%s: deallocating\n", __func__); |
|
|
|
|
|
free(ctx); |
|
|
} |
|
|
|
|
|
#ifdef GGML_MTL_HEAP |
|
|
|
|
|
// make a view of the respective MTL heap |
|
|
id<MTLBuffer> mnist_mtl_get_buffer_on_heap(struct ggml_mtl_context * ctx, struct ggml_tensor * t) { |
|
|
const int64_t offs_data = (int64_t) t->data - (int64_t) ggml_get_mem_buffer(ctx->ctx_data); |
|
|
const int64_t offs_eval = (int64_t) t->data - (int64_t) ggml_get_mem_buffer(ctx->ctx_eval); |
|
|
|
|
|
const bool is_data = (offs_eval < 0) || (offs_data >= 0 && offs_data < offs_eval); |
|
|
|
|
|
const size_t t_size = ggml_nbytes(t); |
|
|
const size_t t_offs = is_data ? offs_data : offs_eval; |
|
|
|
|
|
id<MTLBuffer> result; |
|
|
|
|
|
if (is_data) { |
|
|
fprintf(stderr, "%s: data tensor ' |
|
|
result = [ctx->heap_data newBufferWithLength:t_size options:MTLResourceStorageModeShared offset:t_offs]; |
|
|
} else { |
|
|
fprintf(stderr, "%s: eval tensor '%16s', offs = %8ld, size = %8ld\n", __func__, t->name, t_offs, t_size); |
|
|
result = [ctx->heap_eval newBufferWithLength:t_size options:MTLResourceStorageModePrivate offset:t_offs]; |
|
|
} |
|
|
|
|
|
if (result == nil) { |
|
|
fprintf(stderr, "%s: error: buffer is nil\n", __func__); |
|
|
GGML_ASSERT(false); |
|
|
} |
|
|
|
|
|
return result; |
|
|
} |
|
|
|
|
|
#else |
|
|
|
|
|
// get data / eval buffer + offset |
|
|
id<MTLBuffer> mnist_mtl_get_buffer(struct ggml_mtl_context * ctx, struct ggml_tensor * t, size_t * offs) { |
|
|
const int64_t offs_data = (int64_t) t->data - (int64_t) ggml_get_mem_buffer(ctx->ctx_data); |
|
|
const int64_t offs_eval = (int64_t) t->data - (int64_t) ggml_get_mem_buffer(ctx->ctx_eval); |
|
|
|
|
|
const bool is_data = (offs_eval < 0) || (offs_data >= 0 && offs_data < offs_eval); |
|
|
|
|
|
const size_t t_size = ggml_nbytes(t); |
|
|
const size_t t_offs = is_data ? offs_data : offs_eval; |
|
|
|
|
|
id<MTLBuffer> result; |
|
|
|
|
|
if (is_data) { |
|
|
fprintf(stderr, "%s: data tensor '%16s', offs = %8ld, size = %8ld\n", __func__, t->name, t_offs, t_size); |
|
|
result = ctx->buffer_data; |
|
|
} else { |
|
|
fprintf(stderr, "%s: eval tensor '%16s', offs = %8ld, size = %8ld\n", __func__, t->name, t_offs, t_size); |
|
|
result = ctx->buffer_eval; |
|
|
} |
|
|
|
|
|
if (result == nil) { |
|
|
fprintf(stderr, "%s: error: buffer is nil\n", __func__); |
|
|
GGML_ASSERT(false); |
|
|
} |
|
|
|
|
|
if (offs != nil) { |
|
|
*offs = t_offs; |
|
|
} |
|
|
|
|
|
return result; |
|
|
} |
|
|
|
|
|
#endif |
|
|
|
|
|
int mnist_mtl_eval( |
|
|
struct ggml_mtl_context * ctx, |
|
|
struct ggml_cgraph * gf) { |
|
|
fprintf(stderr, "%s: evaluating\n", __func__); |
|
|
|
|
|
id<MTLCommandBuffer> command_buffer = [ctx->queue commandBuffer]; |
|
|
id<MTLComputeCommandEncoder> encoder = nil; |
|
|
|
|
|
size_t offs_src0; |
|
|
size_t offs_src1; |
|
|
size_t offs_dst; |
|
|
|
|
|
// copy the input data to the GPU |
|
|
{ |
|
|
struct ggml_tensor * inp = ggml_graph_get_tensor(gf, "input"); |
|
|
|
|
|
id<MTLBuffer> id_dst = mnist_mtl_get_buffer(ctx, inp, &offs_src0); |
|
|
|
|
|
memcpy((char *) id_dst.contents + offs_src0, inp->data, ggml_nbytes(inp)); |
|
|
} |
|
|
|
|
|
for (int i = 0; i < gf->n_nodes; ++i) { |
|
|
fprintf(stderr, "%s: encoding node %3d, op = %8s\n", __func__, i, ggml_op_name(gf->nodes[i]->op)); |
|
|
|
|
|
switch (gf->nodes[i]->op) { |
|
|
case GGML_OP_ADD: |
|
|
{ |
|
|
if (encoder == nil) { |
|
|
encoder = [command_buffer computeCommandEncoder]; |
|
|
} |
|
|
|
|
|
id<MTLBuffer> id_src0 = mnist_mtl_get_buffer(ctx, gf->nodes[i]->src[0], &offs_src0); |
|
|
id<MTLBuffer> id_src1 = mnist_mtl_get_buffer(ctx, gf->nodes[i]->src[1], &offs_src1); |
|
|
id<MTLBuffer> id_dst = mnist_mtl_get_buffer(ctx, gf->nodes[i], &offs_dst); |
|
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_add]; |
|
|
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0]; |
|
|
[encoder setBuffer:id_src1 offset:offs_src1 atIndex:1]; |
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:2]; |
|
|
|
|
|
const int64_t n = ggml_nelements(gf->nodes[i]); |
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)]; |
|
|
} break; |
|
|
case GGML_OP_UNARY: |
|
|
switch (ggml_get_unary_op(gf->nodes[i])) { |
|
|
case GGML_UNARY_OP_RELU: |
|
|
{ |
|
|
if (encoder == nil) { |
|
|
encoder = [command_buffer computeCommandEncoder]; |
|
|
} |
|
|
|
|
|
id<MTLBuffer> id_src = mnist_mtl_get_buffer(ctx, gf->nodes[i]->src[0], &offs_src0); |
|
|
id<MTLBuffer> id_dst = mnist_mtl_get_buffer(ctx, gf->nodes[i], &offs_dst); |
|
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_relu]; |
|
|
[encoder setBuffer:id_src offset:offs_src0 atIndex:0]; |
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1]; |
|
|
|
|
|
const int64_t n = ggml_nelements(gf->nodes[i]); |
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)]; |
|
|
} break; |
|
|
default: |
|
|
{ |
|
|
fprintf(stderr, "%s: node %3d, op = %8s, unary op %d not implemented\n", __func__, i, ggml_op_name(gf->nodes[i]->op), (int) ggml_get_unary_op(gf->nodes[i])); |
|
|
GGML_ASSERT(false); |
|
|
return -1; |
|
|
} |
|
|
break; |
|
|
} break; |
|
|
case GGML_OP_SOFT_MAX: |
|
|
{ |
|
|
#if 0 |
|
|
// NOTE: MPSMatrixSoftMax is not working properly, probably there is a bug |
|
|
|
|
|
if (encoder != nil) { |
|
|
[encoder endEncoding]; |
|
|
encoder = nil; |
|
|
} |
|
|
|
|
|
// use MPSMatrixSoftMax |
|
|
id<MTLBuffer> id_src = mnist_mtl_get_buffer(ctx, gf->nodes[i]->src0, &offs_src0); |
|
|
id<MTLBuffer> id_dst = mnist_mtl_get_buffer(ctx, gf->nodes[i], &offs_dst); |
|
|
|
|
|
MPSMatrixDescriptor * desc = [MPSMatrixDescriptor |
|
|
matrixDescriptorWithRows:1 columns:gf->nodes[i]->ne[0] rowBytes:gf->nodes[i]->nb[1] dataType:MPSDataTypeFloat32]; |
|
|
|
|
|
MPSMatrix * mat_src = [[MPSMatrix alloc] initWithBuffer:id_src offset:offs_src0 descriptor:desc]; |
|
|
MPSMatrix * mat_dst = [[MPSMatrix alloc] initWithBuffer:id_dst offset:offs_dst descriptor:desc]; |
|
|
|
|
|
MPSMatrixSoftMax * softmax = [[MPSMatrixSoftMax alloc] initWithDevice:ctx->device]; |
|
|
|
|
|
[softmax encodeToCommandBuffer:command_buffer inputMatrix:mat_src resultMatrix:mat_dst]; |
|
|
#else |
|
|
if (encoder == nil) { |
|
|
encoder = [command_buffer computeCommandEncoder]; |
|
|
} |
|
|
|
|
|
id<MTLBuffer> id_src = mnist_mtl_get_buffer(ctx, gf->nodes[i]->src[0], &offs_src0); |
|
|
id<MTLBuffer> id_dst = mnist_mtl_get_buffer(ctx, gf->nodes[i], &offs_dst); |
|
|
|
|
|
[encoder setComputePipelineState:ctx->pipeline_soft_max]; |
|
|
[encoder setBuffer:id_src offset:offs_src0 atIndex:0]; |
|
|
[encoder setBuffer:id_dst offset:offs_dst atIndex:1]; |
|
|
|
|
|
[encoder dispatchThreadgroups:MTLSizeMake(1, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)]; |
|
|
#endif |
|
|
} break; |
|
|
case GGML_OP_MUL_MAT: |
|
|
{ |
|
|
if (encoder != nil) { |
|
|
[encoder endEncoding]; |
|
|
encoder = nil; |
|
|
} |
|
|
|
|
|
// use MPSMatrixMultiplication |
|
|
id<MTLBuffer> id_src0 = mnist_mtl_get_buffer(ctx, gf->nodes[i]->src[0], &offs_src0); |
|
|
id<MTLBuffer> id_src1 = mnist_mtl_get_buffer(ctx, gf->nodes[i]->src[1], &offs_src1); |
|
|
id<MTLBuffer> id_dst = mnist_mtl_get_buffer(ctx, gf->nodes[i], &offs_dst); |
|
|
|
|
|
const int64_t ncols0 = gf->nodes[i]->src[0]->ne[0]; |
|
|
const int64_t nrows0 = gf->nodes[i]->src[0]->ne[1]; |
|
|
|
|
|
const int64_t ncols1 = gf->nodes[i]->src[1]->ne[0]; |
|
|
const int64_t nrows1 = gf->nodes[i]->src[1]->ne[1]; |
|
|
|
|
|
const int64_t ncols2 = gf->nodes[i]->ne[0]; |
|
|
const int64_t nrows2 = gf->nodes[i]->ne[1]; |
|
|
|
|
|
GGML_ASSERT(ncols0 == ncols1); |
|
|
|
|
|
MPSMatrixDescriptor * desc0 = [MPSMatrixDescriptor |
|
|
matrixDescriptorWithRows:nrows0 columns:ncols0 rowBytes:gf->nodes[i]->src[0]->nb[1] dataType:MPSDataTypeFloat32]; |
|
|
MPSMatrixDescriptor * desc1 = [MPSMatrixDescriptor |
|
|
matrixDescriptorWithRows:nrows1 columns:ncols1 rowBytes:gf->nodes[i]->src[1]->nb[1] dataType:MPSDataTypeFloat32]; |
|
|
MPSMatrixDescriptor * desc2 = [MPSMatrixDescriptor |
|
|
matrixDescriptorWithRows:nrows2 columns:ncols2 rowBytes:gf->nodes[i]->nb[1] dataType:MPSDataTypeFloat32]; |
|
|
|
|
|
MPSMatrix * mat_src0 = [[MPSMatrix alloc] initWithBuffer:id_src0 offset:offs_src0 descriptor:desc0]; |
|
|
MPSMatrix * mat_src1 = [[MPSMatrix alloc] initWithBuffer:id_src1 offset:offs_src1 descriptor:desc1]; |
|
|
MPSMatrix * mat_dst = [[MPSMatrix alloc] initWithBuffer:id_dst offset:offs_dst descriptor:desc2]; |
|
|
|
|
|
MPSMatrixMultiplication * mul = [[MPSMatrixMultiplication alloc] initWithDevice:ctx->device |
|
|
transposeLeft:false transposeRight:true resultRows:nrows1 resultColumns:nrows0 interiorColumns:ncols0 alpha:1.0 beta:0.0]; |
|
|
|
|
|
[mul encodeToCommandBuffer:command_buffer leftMatrix:mat_src1 rightMatrix:mat_src0 resultMatrix:mat_dst]; |
|
|
} break; |
|
|
default: |
|
|
{ |
|
|
fprintf(stderr, "%s: node %3d, op = %8s not implemented\n", __func__, i, ggml_op_name(gf->nodes[i]->op)); |
|
|
GGML_ASSERT(false); |
|
|
return -1; |
|
|
} |
|
|
} |
|
|
} |
|
|
|
|
|
// extract results from the GPU |
|
|
{ |
|
|
if (encoder != nil) { |
|
|
[encoder endEncoding]; |
|
|
encoder = nil; |
|
|
} |
|
|
|
|
|
struct ggml_tensor * out = gf->nodes[gf->n_nodes - 1]; |
|
|
|
|
|
id<MTLBuffer> id_src = mnist_mtl_get_buffer(ctx, out, &offs_src0); |
|
|
id<MTLBuffer> id_dst = ctx->out; |
|
|
|
|
|
id<MTLBlitCommandEncoder> encoder_blit = [command_buffer blitCommandEncoder]; |
|
|
[encoder_blit copyFromBuffer:id_src sourceOffset:offs_src0 toBuffer:id_dst destinationOffset:0 size:ggml_nbytes(out)]; |
|
|
[encoder_blit endEncoding]; |
|
|
} |
|
|
|
|
|
[command_buffer commit]; |
|
|
[command_buffer waitUntilCompleted]; |
|
|
|
|
|
{ |
|
|
const double time_elapsed = [command_buffer GPUEndTime] - [command_buffer GPUStartTime]; |
|
|
fprintf(stderr, "%s: time elapsed = %f\n", __func__, time_elapsed); |
|
|
} |
|
|
|
|
|
// select the most probable digit |
|
|
int result = -1; |
|
|
{ |
|
|
const float * probs = ctx->out.contents; |
|
|
|
|
|
float prob = probs[0]; |
|
|
|
|
|
for (int i = 0; i < 10; ++i) { |
|
|
fprintf(stderr, "%s: probs[%2d] = %f\n", __func__, i, probs[i]); |
|
|
|
|
|
if (probs[i] > prob) { |
|
|
result = i; |
|
|
prob = probs[i]; |
|
|
} |
|
|
} |
|
|
} |
|
|
|
|
|
return result; |
|
|
} |
|
|
|
