wehub-resource-sync
1438 lines
51 KiB
Plaintext
1438 lines
51 KiB
Plaintext
//
|
||
// MetalBackend.mm
|
||
// MNN
|
||
//
|
||
// Created by MNN on 2019/01/30.
|
||
// Copyright © 2018, Alibaba Group Holding Limited
|
||
//
|
||
|
||
#import "backend/metal/MetalBackend.hpp"
|
||
#define MNN_METAL
|
||
#import <MNN/MNNSharedContext.h>
|
||
#define METAL_CONST_BUFFER_LIMIT 128
|
||
#define METAL_SEPERATE_MAX_COUNT 2
|
||
// overload of MTLGPUFamilyMetal3/MTLGPUFamilyMetal4 (not available in some environments)
|
||
#define MTLGPUFamilyMetal3_MNN 5001
|
||
#define MTLGPUFamilyMetal4_MNN 5002
|
||
|
||
#define CHECK_IOS_UI_STATUS
|
||
#if MNN_METAL_ENABLED
|
||
#include <mutex>
|
||
#import "backend/metal/MNNMetalContext.h"
|
||
#import "core/Macro.h"
|
||
#import "core/TensorUtils.hpp"
|
||
#include "MetalCache_generated.h"
|
||
#include "core/MNNFileUtils.h"
|
||
#if defined(TARGET_OS_IPHONE) && TARGET_OS_IPHONE
|
||
#import <UIKit/UIKit.h>
|
||
#endif
|
||
int MNNMetalGetTensorContent(MNNMetalTensorContent* content, void* tensor) {
|
||
if (nullptr == content || nullptr == tensor) {
|
||
return 0;
|
||
}
|
||
auto t = (MNN::Tensor*)tensor;
|
||
auto des = MNN::TensorUtils::getDescribeOrigin(t);
|
||
content->buffer = ((MNN::MetalRuntimeAllocator::MetalBufferAlloc*)t->deviceId())->getBuffer();
|
||
content->texture = nil;
|
||
content->offset = des->offset;
|
||
return 0;
|
||
}
|
||
|
||
namespace MNN {
|
||
|
||
static void _MetalApplyTensor(uint8_t* host, size_t offset, Tensor* t) {
|
||
// ptr of MetalBufferAlloc
|
||
t->buffer().device = (uint64_t)host;
|
||
auto des = TensorUtils::getDescribeOrigin(t);
|
||
des->offset = offset;
|
||
}
|
||
BufferAllocator* MetalRuntime::createDynamicAllocator(int index, bool secondResize) const {
|
||
if (hint().memoryAllocatorType == Runtime::Allocator_Defer && secondResize) {
|
||
return new DeferBufferAllocator(buffer(index), 1024, _MetalApplyTensor);
|
||
}
|
||
if (mStaticAllocatorRaw.get() != nullptr) {
|
||
return new EagerBufferAllocator(BufferAllocator::Allocator::createRecurse(mStaticAllocatorRaw.get()), 1024);
|
||
}
|
||
return new EagerBufferAllocator(BufferAllocator::Allocator::createRecurse(mStaticAllocator.get()), 1024);
|
||
}
|
||
|
||
struct TunedInfo {
|
||
std::vector<std::unique_ptr<MetalCache::OpInfoT>> mInfos;
|
||
};
|
||
|
||
void registerMetalOps();
|
||
#ifdef MNN_SUPPORT_RENDER
|
||
extern void registerMetalRenderOps();
|
||
#endif
|
||
|
||
static inline std::map<OpType, MetalBackend::Creator *> *getCreatorMap() {
|
||
static std::once_flag of;
|
||
static std::map<OpType, MetalBackend::Creator *> *ret = nullptr;
|
||
std::call_once(of, [&]() { ret = new std::map<OpType, MetalBackend::Creator *>; });
|
||
return ret;
|
||
}
|
||
|
||
void MetalBackend::addCreator(OpType t, Creator *c) {
|
||
auto map = getCreatorMap();
|
||
if (map->find(t) != map->end()) {
|
||
MNN_PRINT("Error: %d type has be added\n", t);
|
||
}
|
||
map->insert(std::make_pair(t, c));
|
||
}
|
||
|
||
MetalBackend::MetalBackend(const MetalRuntime* runtime, bool usefp16AsFp32, BackendConfig::MemoryMode mode) : Backend(MNN_FORWARD_METAL),
|
||
mEmptyMem(nil)
|
||
{
|
||
mRuntime = runtime;
|
||
auto ctx = (__bridge MNNMetalContext *)runtime->context();
|
||
mBufferPool.reset(runtime->createDynamicAllocator(0, false));
|
||
mExecutionBufferPool.reset(new EagerBufferAllocator(runtime->buffer(0)->root, 1024));
|
||
mCurrentAllocator = mBufferPool.get();
|
||
mUseFloatAsFp16 = usefp16AsFp32;
|
||
mMemoryMode = mode;
|
||
mIsIphone = ctx.isIphone;
|
||
if (runtime->getCommandQueue() == nil) {
|
||
// one command queue can create only a few command buffer, so let each backend own a command queue
|
||
_commandQueue = [[ctx device] newCommandQueue];
|
||
} else {
|
||
// otherwise forbid defer encode optimize
|
||
_commandQueue = runtime->getCommandQueue();
|
||
}
|
||
if(((MetalRuntime *)mRuntime)->supportTensorOps()) {
|
||
mSupportTensorApi = true;
|
||
const char * src_tensor_f16 = "\n"
|
||
"#include <metal_stdlib> \n"
|
||
"#include <metal_tensor> \n"
|
||
"#include <MetalPerformancePrimitives/MetalPerformancePrimitives.h> \n"
|
||
" \n"
|
||
"using namespace metal; \n"
|
||
"using namespace mpp::tensor_ops; \n"
|
||
" \n"
|
||
"kernel void dummy_kernel( \n"
|
||
" tensor<device half, dextents<int32_t, 2>> A [[buffer(0)]], \n"
|
||
" tensor<device half, dextents<int32_t, 2>> B [[buffer(1)]], \n"
|
||
" device float * C [[buffer(2)]], \n"
|
||
" uint2 tgid [[threadgroup_position_in_grid]]) \n"
|
||
"{ \n"
|
||
" auto tA = A.slice(0, (int)tgid.y); \n"
|
||
" auto tB = B.slice((int)tgid.x, 0); \n"
|
||
" \n"
|
||
" matmul2d< \n"
|
||
" matmul2d_descriptor(8, 8, dynamic_extent), \n"
|
||
" execution_simdgroups<4>> mm; \n"
|
||
" \n"
|
||
" auto cT = mm.get_destination_cooperative_tensor<decltype(tA), decltype(tB), float>(); \n"
|
||
" \n"
|
||
" auto sA = tA.slice(0, 0); \n"
|
||
" auto sB = tB.slice(0, 0); \n"
|
||
" mm.run(sB, sA, cT); \n"
|
||
" \n"
|
||
" auto tC = tensor<device float, dextents<int32_t, 2>, tensor_inline>(C, dextents<int32_t, 2>(4, 4)); \n"
|
||
" \n"
|
||
" cT.store(tC); \n"
|
||
"}";
|
||
|
||
auto pipeline = makeComputePipelineWithSourceOption(src_tensor_f16, "dummy_kernel", nullptr);
|
||
if(pipeline == nullptr) {
|
||
MNN_PRINT("Metal4 Tensor api compile err, disable tensor api.\n");
|
||
mSupportTensorApi = false;
|
||
}
|
||
}
|
||
_commandBuffer = nil;
|
||
setUpGPUEnabledSwitch();
|
||
}
|
||
MetalBackend::~MetalBackend() {
|
||
flushEncoder();
|
||
removeNotificationsObservers();
|
||
}
|
||
|
||
void MetalBackend::setUpGPUEnabledSwitch() {
|
||
#if defined(TARGET_OS_IPHONE) && TARGET_OS_IPHONE
|
||
__block UIApplicationState state;
|
||
if ([NSThread isMainThread]) {
|
||
state = [UIApplication sharedApplication].applicationState;
|
||
} else {
|
||
dispatch_semaphore_t latch = dispatch_semaphore_create(0);
|
||
dispatch_async(dispatch_get_main_queue(), ^{
|
||
state = [UIApplication sharedApplication].applicationState;
|
||
dispatch_semaphore_signal(latch);
|
||
});
|
||
dispatch_semaphore_wait(latch, DISPATCH_TIME_FOREVER);
|
||
}
|
||
mGPUEnabledSwitch.store(state == UIApplicationStateActive);
|
||
mForegroundObserver = [[NSNotificationCenter defaultCenter] addObserverForName:UIApplicationWillEnterForegroundNotification object:nil queue:nil usingBlock:^(NSNotification * _Nonnull notification) {
|
||
mGPUEnabledSwitch.store(true);
|
||
}];
|
||
mBackgroundObserver = [[NSNotificationCenter defaultCenter] addObserverForName:UIApplicationDidEnterBackgroundNotification object:nil queue:nil usingBlock:^(NSNotification * _Nonnull notification) {
|
||
mGPUEnabledSwitch.store(false);
|
||
}];
|
||
#endif
|
||
}
|
||
|
||
void MetalBackend::removeNotificationsObservers() {
|
||
#if defined(TARGET_OS_IPHONE) && TARGET_OS_IPHONE
|
||
[[NSNotificationCenter defaultCenter] removeObserver:mForegroundObserver];
|
||
[[NSNotificationCenter defaultCenter] removeObserver:mBackgroundObserver];
|
||
#endif
|
||
}
|
||
|
||
id<MTLComputeCommandEncoder> MetalBackend::encoder_net() const {
|
||
id<MTLComputeCommandEncoder> result = [getCommandBufferForNet() computeCommandEncoder];
|
||
#if MNN_METAL_DEBUG || MNN_METAL_BENCHMARK
|
||
result.label = nil;
|
||
#endif
|
||
return result;
|
||
}
|
||
|
||
void *MetalBackend::context() const {
|
||
return mRuntime->context();
|
||
}
|
||
|
||
class MetalMemRelease : public Backend::MemObj {
|
||
public:
|
||
MetalMemRelease(MemChunk buffer, BufferAllocator* allocator) {
|
||
mBuffer = buffer;
|
||
mAllocator = allocator;
|
||
}
|
||
virtual ~ MetalMemRelease() {
|
||
mAllocator->free(mBuffer);
|
||
}
|
||
MemChunk chunk() override {
|
||
return mBuffer;
|
||
}
|
||
private:
|
||
MemChunk mBuffer;
|
||
BufferAllocator* mAllocator;
|
||
};
|
||
size_t MetalBackend::getTensorSizeInBytes(const Tensor* tensor) const {
|
||
auto format = TensorUtils::getDescribe(tensor)->dimensionFormat;
|
||
size_t size;
|
||
if (MNN_DATA_FORMAT_NC4HW4 == format && tensor->dimensions() >= 2) {
|
||
int width = 1;
|
||
int height = 1;
|
||
int batch = tensor->length(0);
|
||
int channel = tensor->length(1);
|
||
if (tensor->dimensions() >= 3) {
|
||
height = tensor->length(2);
|
||
}
|
||
for (int i=3; i<tensor->dimensions(); ++i) {
|
||
width *= tensor->length(i);
|
||
}
|
||
int alignC = ROUND_UP(channel, 4);
|
||
int hR = ROUND_UP(height, 4) - height;
|
||
// width parallel 4, may exceed 3 elements
|
||
int wR = ROUND_UP(width + 3, 4) - width;
|
||
int bhw = batch * width * height;
|
||
int bhwR = UP_DIV(bhw, 16) * 16 - bhw;
|
||
int extraPadding = ALIMAX(bhwR, (hR * width + wR));
|
||
size = batch * alignC * width * height;
|
||
size = size + extraPadding * 4;
|
||
} else {
|
||
size = 1;
|
||
for (int i=0; i<tensor->dimensions(); ++i) {
|
||
size *= tensor->length(i);
|
||
}
|
||
size = ROUND_UP(size, 4);
|
||
}
|
||
if (0 == size) {
|
||
return 0;
|
||
}
|
||
// use metal_float when meets float
|
||
if (halide_type_float == tensor->buffer().type.code && tensor->buffer().type.bits == 32 && mUseFloatAsFp16) {
|
||
size *= 2;
|
||
} else {
|
||
size *= tensor->getType().bytes();
|
||
}
|
||
size_t align = 4 * sizeof(int);
|
||
size = ROUND_UP(size, align);
|
||
return size;
|
||
}
|
||
|
||
Backend::MemObj* MetalBackend::onAcquire(const Tensor *_tensor, StorageType storageType) {
|
||
auto tensor = const_cast<Tensor *>(_tensor);
|
||
size_t size = getTensorSizeInBytes(_tensor);
|
||
if (0 == size) {
|
||
return nullptr;
|
||
}
|
||
// reuse if possible
|
||
MemChunk buffer;
|
||
BufferAllocator* allocator = nullptr;
|
||
switch (storageType) {
|
||
case Backend::STATIC: {
|
||
buffer = mRuntime->mStaticAllocator->alloc(size, false);
|
||
allocator = mRuntime->mStaticAllocator.get();
|
||
} break;
|
||
case Backend::DYNAMIC: {
|
||
buffer = mCurrentAllocator->alloc(size, false);
|
||
allocator = mCurrentAllocator;
|
||
} break;
|
||
case Backend::DYNAMIC_SEPERATE: {
|
||
buffer = mCurrentAllocator->alloc(size, true);
|
||
allocator = mCurrentAllocator;
|
||
} break;
|
||
case Backend::DYNAMIC_IN_EXECUTION: {
|
||
buffer = mExecutionBufferPool->alloc(size, false);
|
||
allocator = mExecutionBufferPool.get();
|
||
} break;
|
||
default:{
|
||
break;
|
||
}
|
||
}
|
||
if (storageType == Backend::STATIC) {
|
||
if(nullptr == buffer.first) {
|
||
MNN_ERROR("onAcquireBuffer error!\n");
|
||
return nullptr;
|
||
}
|
||
} else {
|
||
buffer.attach(tensor);
|
||
}
|
||
if (nullptr == buffer.first) {
|
||
_MetalApplyTensor((uint8_t*)(&mEmptyMem), 0, (Tensor*)_tensor);
|
||
} else {
|
||
_MetalApplyTensor((uint8_t*)buffer.first, buffer.second, (Tensor*)_tensor);
|
||
}
|
||
return new MetalMemRelease(buffer, allocator);
|
||
}
|
||
|
||
bool MetalBackend::onClearBuffer() {
|
||
mCurrentAllocator->release(true);
|
||
if (mExecutionBufferPool.get() != nullptr) {
|
||
mExecutionBufferPool->release(true);
|
||
}
|
||
if (nullptr != mRuntime->mStaticAllocatorRaw.get()) {
|
||
mRuntime->mStaticAllocator->sync();
|
||
mRuntime->mStaticAllocator = mRuntime->mStaticAllocatorRaw;
|
||
mRuntime->mStaticAllocatorRaw = nullptr;
|
||
}
|
||
return true;
|
||
}
|
||
|
||
Execution *MetalBackend::onCreate(const std::vector<Tensor *> &inputs, const std::vector<Tensor *> &outputs,
|
||
const Op *op) {
|
||
auto map = getCreatorMap();
|
||
|
||
auto iter = map->find(op->type());
|
||
if (iter == map->end()) {
|
||
if (nullptr != op->name()) {
|
||
MNN_PRINT("Don't support type [%s], %s\n", EnumNameOpType(op->type()), op->name()->c_str());
|
||
} else {
|
||
MNN_PRINT("Don't support type [%s]\n", EnumNameOpType(op->type()));
|
||
}
|
||
return NULL;
|
||
}
|
||
//MNN_PRINT("support type [%s]\n", EnumNameOpType(op->type()));
|
||
|
||
auto exe = iter->second->onCreate(inputs, op, this, outputs);
|
||
if (NULL == exe) {
|
||
MNN_PRINT("The Creator Don't support type [%s], %s\n", MNN::EnumNameOpType(op->type()), op->name() ? op->name()->c_str() : "");
|
||
return NULL;
|
||
}
|
||
return exe;
|
||
}
|
||
void MetalBackend::flushEncoder() const {
|
||
if (nil != mComputeEncoder) {
|
||
[mComputeEncoder endEncoding];
|
||
mComputeEncoder = nil;
|
||
}
|
||
}
|
||
void MetalBackend::_resetDynamicMemory() const {
|
||
mRuntime->pCurrentStatus = mCurrentAllocator->apply();
|
||
if (NO_ERROR != mRuntime->pCurrentStatus) {
|
||
return;
|
||
}
|
||
if (nullptr != mBufferPoolShapeImmutable.get()) {
|
||
mRuntime->pCurrentStatus = mBufferPoolShapeImmutable->apply();
|
||
}
|
||
}
|
||
|
||
void MetalBackend::onExecuteBegin() const {
|
||
_resetDynamicMemory();
|
||
mEncoderCount = 0;
|
||
}
|
||
void MetalBackend::onExecuteEnd() const {
|
||
flushEncoder();
|
||
commit_net();
|
||
}
|
||
|
||
BufferAllocator* MetalBackend::getBufferPool() const {
|
||
return mCurrentAllocator;
|
||
}
|
||
|
||
bool MetalBackend::onSelectDynamicAllocator(int index, int maxIndex) {
|
||
if (maxIndex > 2) {
|
||
return false;
|
||
}
|
||
if (maxIndex == 2 && mBufferPoolShapeImmutable.get() == nullptr) {
|
||
mBufferPoolShapeImmutable.reset(mRuntime->createDynamicAllocator(1, true));
|
||
mBufferPool.reset(mRuntime->createDynamicAllocator(0, true));
|
||
}
|
||
if (1 == index) {
|
||
mCurrentAllocator = mBufferPoolShapeImmutable.get();
|
||
} else {
|
||
mCurrentAllocator = mBufferPool.get();
|
||
}
|
||
return true;
|
||
}
|
||
|
||
bool MetalBackend::onGetTensorInfo(const Tensor* tensor, void* dstInfo) {
|
||
if (nullptr == dstInfo) {
|
||
return true;
|
||
}
|
||
auto dst = (MNNMetalTensorContent*)dstInfo;
|
||
dst->type.code = halide_type_float;
|
||
if (mUseFloatAsFp16) {
|
||
dst->type.bits = 16;
|
||
} else {
|
||
dst->type.bits = 32;
|
||
}
|
||
MNNMetalGetTensorContent(dst, (void*)tensor);
|
||
return true;
|
||
}
|
||
|
||
bool MetalBackend::isCmdBufferCommit() {
|
||
auto ctx = (__bridge MNNMetalContext *)context();
|
||
|
||
//TODO: set magic number
|
||
const int magicNum = mRuntime->hint().encorderNumForCommit;
|
||
mEncoderCount++;
|
||
if(mEncoderCount != 0 && mEncoderCount % magicNum == 0) {
|
||
return true;
|
||
}
|
||
return false;
|
||
}
|
||
|
||
id<MTLBuffer> MetalBackend::getHostBuffer(size_t size) const {
|
||
size = UP_DIV(size, METAL_CONST_BUFFER_LIMIT) * METAL_CONST_BUFFER_LIMIT;
|
||
// reuse
|
||
if (nullptr != mHostBuffer && mHostBuffer.length >= size) {
|
||
return mHostBuffer;
|
||
}
|
||
|
||
// create larger
|
||
auto context = (__bridge MNNMetalContext *)this->context();
|
||
mHostBuffer = [context newDeviceBuffer:size access:CPUReadWrite];
|
||
return mHostBuffer;
|
||
}
|
||
|
||
id<MTLBuffer> MetalBackend::getConstBuffer(size_t size) const {
|
||
if (size < METAL_CONST_BUFFER_LIMIT) {
|
||
if (!mHoldBuffers.empty()) {
|
||
auto res = mHoldBuffers.front();
|
||
mHoldBuffers.pop();
|
||
return res;
|
||
}
|
||
size = METAL_CONST_BUFFER_LIMIT;
|
||
}
|
||
auto context = (__bridge MNNMetalContext *)this->context();
|
||
auto buffer = [context newDeviceBuffer:size access:CPUReadWrite];
|
||
return buffer;
|
||
}
|
||
void MetalBackend::returnConstBuffer(id<MTLBuffer> buffer) const {
|
||
mHoldBuffers.push(buffer);
|
||
}
|
||
static inline void _getNCPlane(const Tensor* tensor, int& s, int& c, int& b) {
|
||
auto format = TensorUtils::getDescribe(tensor)->dimensionFormat;
|
||
s = 1, c = 1, b = 1;
|
||
b = tensor->length(0);
|
||
if (format == MNN_DATA_FORMAT_NHWC) {
|
||
c = tensor->length(tensor->dimensions()-1);
|
||
for (int i=1; i<tensor->dimensions()-1; ++i) {
|
||
s *= tensor->length(i);
|
||
}
|
||
} else {
|
||
c = tensor->length(1);
|
||
for (int i=2; i<tensor->dimensions(); ++i) {
|
||
s *= tensor->length(i);
|
||
}
|
||
}
|
||
}
|
||
MTLSize getTensorShape(id<MTLBuffer> shape, const Tensor *tensor) {
|
||
auto format = TensorUtils::getDescribe(tensor)->dimensionFormat;
|
||
int s, b, c;
|
||
_getNCPlane(tensor, s, c, b);
|
||
int z = UP_DIV(c, 4);
|
||
|
||
// shape
|
||
((int *)shape.contents)[0] = b;
|
||
((int *)shape.contents)[1] = c;
|
||
((int *)shape.contents)[2] = s;
|
||
((int *)shape.contents)[3] = 1;
|
||
|
||
// stride
|
||
if (format == MNN_DATA_FORMAT_NHWC) {
|
||
((int *)shape.contents)[4] = s * c;
|
||
((int *)shape.contents)[5] = 1;
|
||
((int *)shape.contents)[6] = c;
|
||
((int *)shape.contents)[7] = 1;
|
||
} else {
|
||
((int *)shape.contents)[4] = s * c;
|
||
((int *)shape.contents)[5] = s;
|
||
((int *)shape.contents)[6] = 1;
|
||
((int *)shape.contents)[7] = 1;
|
||
}
|
||
// threads
|
||
MTLSize threads = {(NSUInteger)s * b * z, 1, 1};
|
||
return threads;
|
||
}
|
||
static const char* gTranspose = R"metal(
|
||
#include <metal_stdlib>
|
||
#include <simd/simd.h>
|
||
using namespace metal;
|
||
struct tensor_shape {
|
||
uint4 size; // n, c, plane, 1
|
||
uint4 stride;
|
||
};
|
||
kernel void main0(const device IType* in [[buffer(0)]], device OType* out [[buffer(1)]], constant tensor_shape &uConstant [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
|
||
int channel = uConstant.size.y;
|
||
if (gid < channel * uConstant.size.x * uConstant.size.z) {
|
||
int tmp = gid % (channel * uConstant.size.x);
|
||
int x = gid / (channel * uConstant.size.x);
|
||
int b = tmp / channel;
|
||
int c = tmp % channel;
|
||
int outPos = b * uConstant.size.y * uConstant.size.z + c * uConstant.size.z + x;
|
||
int inPos = b * uConstant.size.y * uConstant.size.z + c + x * uConstant.size.y;
|
||
out[outPos] = (OType)(in[inPos]);
|
||
}
|
||
})metal";
|
||
|
||
static const char* gNC4HW4Convert = R"metal(
|
||
#include <metal_stdlib>
|
||
#include <simd/simd.h>
|
||
using namespace metal;
|
||
struct tensor_shape {
|
||
uint4 size; // n, c, plane, 1
|
||
uint4 stride;
|
||
};
|
||
kernel void main0(const device IType* in [[buffer(0)]], device OType* out [[buffer(1)]], constant tensor_shape &uConstant [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
|
||
int channelC4 = (uConstant.size.y + 3) / 4;
|
||
if (gid < channelC4 * uConstant.size.x * uConstant.size.z)
|
||
{
|
||
int3 pos;
|
||
pos.z = gid % (channelC4 * uConstant.size.x);
|
||
pos.y = gid / (channelC4 * uConstant.size.x);
|
||
pos.x = 0;
|
||
int batchIndex = pos.z / channelC4;
|
||
int zDiv4 = pos.z % channelC4;
|
||
|
||
int lastZ = uConstant.size.y / 4;
|
||
int cIndex = uConstant.size.y % 4;
|
||
|
||
int z = zDiv4*4;
|
||
int basicOffset = 0
|
||
+ batchIndex*uConstant.stride.x
|
||
+ z * uConstant.stride.y
|
||
+ pos.y * uConstant.stride.z
|
||
;
|
||
#ifdef MNN_OUTPUT_C4
|
||
OType color = OType(0);
|
||
if(zDiv4 == lastZ)
|
||
{
|
||
if(cIndex == 1)
|
||
{
|
||
color.r = in[basicOffset+0];
|
||
color.g = 0.0;
|
||
color.b = 0.0;
|
||
color.a = 0.0;
|
||
}
|
||
else if(cIndex == 2)
|
||
{
|
||
color.r = in[basicOffset+0];
|
||
color.g = in[basicOffset+1*uConstant.stride.y];
|
||
color.b = 0.0;
|
||
color.a = 0.0;
|
||
}
|
||
else
|
||
{
|
||
color.r = in[basicOffset+0];
|
||
color.g = in[basicOffset+1*uConstant.stride.y];
|
||
color.b = in[basicOffset+2*uConstant.stride.y];
|
||
color.a = 0.0;
|
||
}
|
||
}
|
||
else
|
||
{
|
||
color.r = in[basicOffset+0];
|
||
color.g = in[basicOffset+1*uConstant.stride.y];
|
||
color.b = in[basicOffset+2*uConstant.stride.y];
|
||
color.a = in[basicOffset+3*uConstant.stride.y];
|
||
}
|
||
|
||
out[0
|
||
+ pos.y
|
||
+ uConstant.size.x * uConstant.size.z*zDiv4
|
||
+ batchIndex*uConstant.size.z
|
||
] = color;
|
||
#else
|
||
IType color = in[0
|
||
+ pos.y
|
||
+ uConstant.size.x * uConstant.size.z*zDiv4
|
||
+ batchIndex*uConstant.size.z
|
||
];
|
||
if(zDiv4 == lastZ)
|
||
{
|
||
if(cIndex == 1)
|
||
{
|
||
out[basicOffset+0*uConstant.stride.y] = color.r;
|
||
}
|
||
else if(cIndex == 2)
|
||
{
|
||
out[basicOffset+0*uConstant.stride.y] = color.r;
|
||
out[basicOffset+1*uConstant.stride.y] = color.g;
|
||
}
|
||
else
|
||
{
|
||
out[basicOffset+0*uConstant.stride.y] = color.r;
|
||
out[basicOffset+1*uConstant.stride.y] = color.g;
|
||
out[basicOffset+2*uConstant.stride.y] = color.b;
|
||
}
|
||
}
|
||
else
|
||
{
|
||
out[basicOffset+0*uConstant.stride.y] = color.r;
|
||
out[basicOffset+1*uConstant.stride.y] = color.g;
|
||
out[basicOffset+2*uConstant.stride.y] = color.b;
|
||
out[basicOffset+3*uConstant.stride.y] = color.a;
|
||
}
|
||
#endif
|
||
}
|
||
}
|
||
)metal";
|
||
|
||
static const char* gCopy = R"metal(
|
||
#include <metal_stdlib>
|
||
#include <simd/simd.h>
|
||
using namespace metal;
|
||
kernel void main0(const device IType *in [[buffer(0)]], device OType *out [[buffer(1)]], constant uint4& limit [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
|
||
if (gid < limit.x) {
|
||
out[int(gid)] = (OType)in[int(gid)];
|
||
}
|
||
})metal";
|
||
|
||
void MetalBackend::onResizeBegin() {
|
||
// Abort last inference task if needed
|
||
flushEncoder();
|
||
_commandBuffer = nil;
|
||
wait();
|
||
mCurrentAllocator->reset();
|
||
}
|
||
|
||
ErrorCode MetalBackend::onResizeEnd() {
|
||
auto ctx = (__bridge MNNMetalContext *)context();
|
||
return mCurrentAllocator->compute();
|
||
}
|
||
|
||
static std::string _getType(const halide_type_t& type, MNN_DATA_FORMAT format, bool useFp16AsFp32) {
|
||
std::string res;
|
||
if (type.code == halide_type_float) {
|
||
if (useFp16AsFp32) {
|
||
res = "half";
|
||
} else {
|
||
res = "float";
|
||
}
|
||
} else {
|
||
switch (type.bytes()) {
|
||
case 1:
|
||
res = "char";
|
||
break;
|
||
case 2:
|
||
res = "short";
|
||
break;
|
||
case 4:
|
||
res = "int";
|
||
break;
|
||
default:
|
||
MNN_ASSERT(false);
|
||
break;
|
||
}
|
||
}
|
||
if (format == MNN_DATA_FORMAT_NC4HW4) {
|
||
return res + "4";
|
||
}
|
||
return res;
|
||
}
|
||
MetalBackend::CopyPipeline MetalBackend::_makeCopyInfo(const Tensor *src, const Tensor *dst, id<MTLBuffer> shape, int castType) const {
|
||
auto ctx = (__bridge MNNMetalContext *)context();
|
||
MetalBackend::CopyPipeline res;
|
||
auto sfmt = TensorUtils::getDescribe(src)->dimensionFormat;
|
||
auto dfmt = TensorUtils::getDescribe(dst)->dimensionFormat;
|
||
if (shape == nil) {
|
||
shape = getConstBuffer(8 * sizeof(int));
|
||
}
|
||
res.shape = shape;
|
||
if (sfmt == dfmt || src->dimensions() <= 1) {
|
||
auto srcType = _getType(src->getType(), MNN_DATA_FORMAT_NC4HW4, mUseFloatAsFp16 && castType != 1);
|
||
auto dstType = _getType(dst->getType(), MNN_DATA_FORMAT_NC4HW4, mUseFloatAsFp16 && castType != 2);
|
||
auto size = dst->elementSize();
|
||
size = UP_DIV(size, 4);
|
||
std::vector<std::string> keys = {
|
||
"copyC4",
|
||
srcType,
|
||
dstType
|
||
};
|
||
((uint32_t*)[shape contents])[0] = size;
|
||
id<MTLComputePipelineState> pipeline = mRuntime->findPipeline(keys);
|
||
if (nil == pipeline) {
|
||
MTLCompileOptions *option = [[MTLCompileOptions alloc] init];
|
||
auto dic = [NSMutableDictionary dictionaryWithCapacity:0];
|
||
[dic setValue:@(keys[1].c_str()) forKey:@"IType"];
|
||
[dic setValue:@(keys[2].c_str()) forKey:@"OType"];
|
||
option.preprocessorMacros = dic;
|
||
pipeline = makeComputePipelineWithSourceOption(gCopy, "main0", option);
|
||
mRuntime->insertPipeline(keys, pipeline);
|
||
}
|
||
res.groupSize = MTLSizeMake(UP_DIV(size, 256), 1, 1);
|
||
res.localSize = MTLSizeMake(256, 1, 1);
|
||
res.pipeline = pipeline;
|
||
return res;
|
||
}
|
||
auto srcType = _getType(src->getType(), sfmt, mUseFloatAsFp16 && castType != 1);
|
||
auto dstType = _getType(dst->getType(), dfmt, mUseFloatAsFp16 && castType != 2);
|
||
if (sfmt == MNN_DATA_FORMAT_NC4HW4 || dfmt == MNN_DATA_FORMAT_NC4HW4) {
|
||
auto normalTensor = dst;
|
||
if (dfmt == MNN_DATA_FORMAT_NC4HW4) {
|
||
normalTensor = src;
|
||
}
|
||
// convert C4 / NCHW
|
||
std::vector<std::string> keys = {
|
||
"c4convert",
|
||
srcType,
|
||
dstType
|
||
};
|
||
if (dfmt == MNN_DATA_FORMAT_NC4HW4) {
|
||
keys.emplace_back("outputc4");
|
||
}
|
||
id<MTLComputePipelineState> pipeline = mRuntime->findPipeline(keys);
|
||
if (nil == pipeline) {
|
||
MTLCompileOptions *option = [[MTLCompileOptions alloc] init];
|
||
auto dic = [NSMutableDictionary dictionaryWithCapacity:0];
|
||
[dic setValue:@(keys[1].c_str()) forKey:@"IType"];
|
||
[dic setValue:@(keys[2].c_str()) forKey:@"OType"];
|
||
if (dfmt == MNN_DATA_FORMAT_NC4HW4) {
|
||
[dic setValue:@"1" forKey:@"MNN_OUTPUT_C4"];
|
||
}
|
||
option.preprocessorMacros = dic;
|
||
pipeline = makeComputePipelineWithSourceOption(gNC4HW4Convert, "main0", option);
|
||
mRuntime->insertPipeline(keys, pipeline);
|
||
}
|
||
res.pipeline = pipeline;
|
||
auto size = getTensorShape(shape, normalTensor);
|
||
auto gl = [ctx computeBestGroupAndLocal:pipeline threads:size];
|
||
res.groupSize = gl.first;
|
||
res.localSize = gl.second;
|
||
return res;
|
||
}
|
||
// NCHW <-> NHWC
|
||
std::vector<std::string> keys = {
|
||
"transpose",
|
||
srcType,
|
||
dstType
|
||
};
|
||
id<MTLComputePipelineState> pipeline = mRuntime->findPipeline(keys);
|
||
if (nil == pipeline) {
|
||
MTLCompileOptions *option = [[MTLCompileOptions alloc] init];
|
||
auto dic = [NSMutableDictionary dictionaryWithCapacity:0];
|
||
[dic setValue:@(keys[1].c_str()) forKey:@"IType"];
|
||
[dic setValue:@(keys[2].c_str()) forKey:@"OType"];
|
||
option.preprocessorMacros = dic;
|
||
pipeline = makeComputePipelineWithSourceOption(gTranspose, "main0", option);
|
||
mRuntime->insertPipeline(keys, pipeline);
|
||
}
|
||
res.pipeline = pipeline;
|
||
int n, c, plane;
|
||
_getNCPlane(dst, plane, c, n);
|
||
auto shapePtr = (uint32_t*)shape.contents;
|
||
shapePtr[0] = n;
|
||
shapePtr[3] = 1;
|
||
if (MNN_DATA_FORMAT_NHWC == dfmt) {
|
||
shapePtr[1] = plane;
|
||
shapePtr[2] = c;
|
||
} else {
|
||
shapePtr[1] = c;
|
||
shapePtr[2] = plane;
|
||
}
|
||
auto size = plane * n * c;
|
||
res.localSize = MTLSizeMake(256, 1, 1);
|
||
res.groupSize = MTLSizeMake(UP_DIV(size, 256), 1, 1);
|
||
return res;
|
||
}
|
||
|
||
static void _execute(id<MTLComputeCommandEncoder> encoder, const MetalBackend::CopyPipeline& info, std::pair<id<MTLBuffer>, int> src, std::pair<id<MTLBuffer>, int> dst) {
|
||
[encoder setComputePipelineState:info.pipeline];
|
||
[encoder setBuffer:src.first offset:src.second atIndex:0];
|
||
[encoder setBuffer:dst.first offset:dst.second atIndex:1];
|
||
[encoder setBuffer:info.shape offset:0 atIndex:2];
|
||
[encoder dispatchThreadgroups:info.groupSize threadsPerThreadgroup:info.localSize];
|
||
}
|
||
void MetalBackend::onCopyDeviceToDevice(const Tensor *src, const Tensor *dst,
|
||
id<MTLComputeCommandEncoder> encoder, id<MTLBuffer> shape, int castType) const {
|
||
auto ctx = (__bridge MNNMetalContext *)context();
|
||
auto info = _makeCopyInfo(src, dst, shape, castType);
|
||
auto standalone = encoder == nil;
|
||
encoder = encoder ?: [getCommandBufferForBufferCopy() computeCommandEncoder];
|
||
_execute(encoder, info, MetalBackend::getBuffer(src), MetalBackend::getBuffer(dst));
|
||
if (standalone) {
|
||
[encoder endEncoding];
|
||
MNN_PRINT_ENCODER(ctx, encoder);
|
||
}
|
||
}
|
||
|
||
void MetalBackend::onCopyBuffer(const Tensor *src, const Tensor *dst) const {
|
||
flushEncoder();
|
||
auto ctx = (__bridge MNNMetalContext *)context();
|
||
commit_net();
|
||
|
||
_resetDynamicMemory();
|
||
onCopyBuffer(src, dst, nil, nil);
|
||
}
|
||
|
||
id<MTLComputeCommandEncoder> MetalBackend::encoder_for_net() const {
|
||
if (nil == mComputeEncoder) {
|
||
mComputeEncoder = encoder_net();//TO DO :: use which cmdBuffer
|
||
}
|
||
return mComputeEncoder;
|
||
}
|
||
|
||
void MetalBackend::onCopyBuffer(const Tensor *src, const Tensor *dst, id<MTLComputeCommandEncoder> encoder, id<MTLBuffer> shape) const {
|
||
MNN_ASSERT(src->buffer().dimensions == dst->buffer().dimensions);
|
||
|
||
if (!src->buffer().host && !dst->buffer().host) {
|
||
onCopyDeviceToDevice(src, dst, encoder, shape);
|
||
return;
|
||
}
|
||
auto sfmt = TensorUtils::getDescribe(src)->dimensionFormat;
|
||
auto dfmt = TensorUtils::getDescribe(dst)->dimensionFormat;
|
||
bool formatDiff = sfmt != dfmt && src->dimensions() > 1;
|
||
auto floats = src->getType().code == halide_type_float;
|
||
bool dataTypeDiff = floats && mUseFloatAsFp16;
|
||
bool needConvert = formatDiff || dataTypeDiff;
|
||
|
||
if (!src->buffer().host && dst->buffer().host) {
|
||
auto device = (id<MTLBuffer>)((MetalRuntimeAllocator::MetalBufferAlloc *)src->deviceId())->getBuffer();
|
||
auto devicePtr = (uint8_t*)device.contents + TensorUtils::getDescribeOrigin(src)->offset;
|
||
if (needConvert) {
|
||
auto tDst = const_cast<Tensor*>(dst);
|
||
auto tmpBuffer = getHostBuffer(dst->usize());
|
||
auto info = _makeCopyInfo(src, dst, shape, 2);
|
||
auto standalone = encoder == nil;
|
||
encoder = encoder ?: [getCommandBufferForBufferCopy() computeCommandEncoder];
|
||
_execute(encoder, info, MetalBackend::getBuffer(src), std::make_pair(tmpBuffer, 0));
|
||
if (standalone) {
|
||
[encoder endEncoding];
|
||
}
|
||
commit();
|
||
devicePtr = (uint8_t*)tmpBuffer.contents;
|
||
}
|
||
wait();
|
||
::memcpy(dst->host<void>(), devicePtr, dst->usize());
|
||
return;
|
||
}
|
||
if (src->buffer().host && !dst->buffer().host) {
|
||
// For command queue from user, need user to make sure last frame's gpu work is ready
|
||
bool needWait = !mRuntime->userSync();
|
||
if (needWait) {
|
||
wait();
|
||
}
|
||
auto srcSize = src->usize();
|
||
if (needConvert) {
|
||
auto tmpBuffer = getHostBuffer(srcSize);
|
||
::memcpy(tmpBuffer.contents, src->host<void>(), srcSize);
|
||
auto info = _makeCopyInfo(src, dst, shape, 1);
|
||
auto standalone = encoder == nil;
|
||
encoder = encoder ?: [getCommandBufferForBufferCopy() computeCommandEncoder];
|
||
_execute(encoder, info, std::make_pair(tmpBuffer, 0), MetalBackend::getBuffer(dst));
|
||
if (standalone) {
|
||
[encoder endEncoding];
|
||
}
|
||
commit();
|
||
} else {
|
||
auto device = (id<MTLBuffer>)((MetalRuntimeAllocator::MetalBufferAlloc *)dst->deviceId())->getBuffer();
|
||
auto devicePtr = (uint8_t*)device.contents + TensorUtils::getDescribeOrigin(dst)->offset;
|
||
::memcpy(devicePtr, src->host<void>(), srcSize);
|
||
}
|
||
return;
|
||
}
|
||
MNN_ASSERT(false); // should not be handled here
|
||
}
|
||
int MetalBackend::onSync(Tensor::MapType mtype, bool toCpu, const Tensor* dstTensor) {
|
||
if (mRuntime->pExecutionStatus == NO_EXECUTION) {
|
||
return NO_EXECUTION;
|
||
}
|
||
flushEncoder();
|
||
auto ctx = (__bridge MNNMetalContext *)context();
|
||
commit_net();
|
||
|
||
if (toCpu) {
|
||
wait();
|
||
}
|
||
return 0;
|
||
}
|
||
id<MTLCommandBuffer> MetalBackend::getCommandBufferForBufferCopy() const {
|
||
if (nil == _commandBuffer) {
|
||
_commandBuffer = [_commandQueue commandBuffer];
|
||
}
|
||
return _commandBuffer;
|
||
}
|
||
id<MTLCommandBuffer> MetalBackend::getCommandBufferForNet() const {
|
||
return getCommandBufferForBufferCopy();
|
||
}
|
||
|
||
void MetalBackend::setTensor(const MNN::Tensor* tensor, id<MTLComputeCommandEncoder> encoder, int index) {
|
||
[encoder setBuffer:((MetalRuntimeAllocator::MetalBufferAlloc *)tensor->deviceId())->getBuffer() offset:TensorUtils::getDescribeOrigin(tensor)->offset atIndex:index];
|
||
}
|
||
void MetalBackend::setMem(const MemChunk& chunk, id<MTLComputeCommandEncoder> encoder, int index) {
|
||
[encoder setBuffer:((MetalRuntimeAllocator::MetalBufferAlloc *)chunk.first)->getBuffer() offset:chunk.second atIndex:index];
|
||
}
|
||
uint8_t* MetalBackend::getMemPtr(const MemChunk& chunk) {
|
||
return (uint8_t*)((MetalRuntimeAllocator::MetalBufferAlloc *)chunk.first)->getBuffer().contents + chunk.second;
|
||
}
|
||
void MetalBackend::setBuffer(id<MTLBuffer> buffer, int offset, id<MTLComputeCommandEncoder> encoder, int index) {
|
||
[encoder setBuffer:buffer offset:offset atIndex:index];
|
||
}
|
||
std::pair<id<MTLBuffer>, int> MetalBackend::getBuffer(const MNN::Tensor* tensor) {
|
||
return std::make_pair(((MetalRuntimeAllocator::MetalBufferAlloc *)tensor->deviceId())->getBuffer(), TensorUtils::getDescribeOrigin(tensor)->offset);
|
||
}
|
||
|
||
|
||
void MetalBackend::commit() const {
|
||
#ifdef CHECK_IOS_UI_STATUS
|
||
#if defined(TARGET_OS_IPHONE) && TARGET_OS_IPHONE
|
||
if (!mGPUEnabledSwitch) {
|
||
mRuntime->pExecutionStatus = NO_EXECUTION;
|
||
_commandBuffer = nil;
|
||
return;
|
||
}
|
||
#endif
|
||
#endif
|
||
mRuntime->pExecutionStatus = NO_ERROR;
|
||
if (nil != _commandBuffer && _commandBuffer.status < MTLCommandBufferStatusCommitted) {
|
||
[_commandBuffer commit];
|
||
mRuntime->_waiting = _commandBuffer;
|
||
_commandBuffer = nil;
|
||
}
|
||
}
|
||
|
||
void MetalBackend::commit_net() const {
|
||
commit();
|
||
}
|
||
|
||
void MetalBackend::wait() const {
|
||
if (nil != mRuntime->_waiting) {
|
||
auto buffer = mRuntime->_waiting;
|
||
if (buffer.status >= MTLCommandBufferStatusCompleted) {
|
||
mRuntime->_waiting = nil;
|
||
return;
|
||
}
|
||
|
||
#if MNN_METAL_BENCHMARK
|
||
NSTimeInterval begin = [NSDate timeIntervalSinceReferenceDate];
|
||
[buffer waitUntilCompleted];
|
||
NSTimeInterval end = [NSDate timeIntervalSinceReferenceDate];
|
||
if (@available(iOS 10.3, *)) {
|
||
printf("[METAL] commit costs: %.3fms\t(kernel: %.3fms, GPU: %.3fms)\n", (end - begin) * 1000.f,
|
||
(buffer.kernelEndTime - buffer.kernelStartTime) * 1000.f,
|
||
(buffer.GPUEndTime - buffer.GPUStartTime) * 1000.f);
|
||
} else {
|
||
printf("[METAL] commit costs: %.3fms\n", (end - begin) * 1000.f);
|
||
}
|
||
#else
|
||
[buffer waitUntilCompleted];
|
||
#endif
|
||
|
||
#if MNN_METAL_DEBUG
|
||
if (buffer.error) {
|
||
printf("[METAL] %s\n", buffer.error.localizedDescription.UTF8String);
|
||
}
|
||
#endif
|
||
}
|
||
mRuntime->_waiting = nil;
|
||
}
|
||
|
||
id<MTLComputePipelineState> MetalBackend::makeComputePipelineWithSourceOption(const char* csource, const char* cname, MTLCompileOptions *options) const{
|
||
auto ctx = (__bridge MNNMetalContext *)context();
|
||
auto source = [[NSString alloc] initWithUTF8String:csource];
|
||
auto name = [[NSString alloc] initWithUTF8String:cname];
|
||
auto pipeline = [ctx pipelineWithSourceOption:source name:name options:options];
|
||
if (nil == pipeline) {
|
||
mRuntime->pCurrentStatus = NOT_SUPPORT;
|
||
MNN_ERROR("pipelineWithSourceOption error.\n");
|
||
}
|
||
return pipeline;
|
||
}
|
||
void MetalRuntime::setCommandQueue(id<MTLCommandQueue> queue, bool userSync) {
|
||
mQueue = queue;
|
||
mUserSync = userSync;
|
||
}
|
||
id<MTLComputePipelineState> MetalRuntime::findPipeline(const std::vector<std::string>& keys) const {
|
||
auto iter = mCachePipeine.find(keys);
|
||
if (iter == mCachePipeine.end()) {
|
||
return nil;
|
||
}
|
||
return iter->second;
|
||
}
|
||
void MetalRuntime::insertPipeline(const std::vector<std::string>& keys, id<MTLComputePipelineState> pipeline) const {
|
||
if (nil != pipeline) {
|
||
mCachePipeine.insert(std::make_pair(keys, pipeline));
|
||
}
|
||
}
|
||
|
||
void MetalRuntime::setGpuMode(const int mode_num) {
|
||
int totalSet = 0;
|
||
bool isSet = (mode_num & MNN_GPU_MEMORY_BUFFER);
|
||
if(isSet) {
|
||
totalSet++;
|
||
}
|
||
isSet = (mode_num & MNN_GPU_MEMORY_IMAGE);
|
||
if(isSet) {
|
||
totalSet++;
|
||
}
|
||
if(totalSet > 0) {
|
||
MNN_PRINT("warning: set BUFFER and IMAGE mode is not useful for metal, it doesn't matter, cl_mode:%x！\n", mode_num);
|
||
}
|
||
|
||
totalSet = 0;
|
||
isSet = (mode_num & MNN_GPU_TUNING_NONE);
|
||
if(isSet) {
|
||
mTuneLevel = Never;
|
||
totalSet++;
|
||
}
|
||
|
||
isSet = (mode_num & MNN_GPU_TUNING_FAST);
|
||
if(isSet) {
|
||
mTuneLevel = Fast;
|
||
totalSet++;
|
||
}
|
||
|
||
isSet = (mode_num & MNN_GPU_TUNING_NORMAL);
|
||
if(isSet) {
|
||
mTuneLevel = Normal;
|
||
totalSet++;
|
||
}
|
||
|
||
isSet = (mode_num & MNN_GPU_TUNING_HEAVY);
|
||
if(isSet) {
|
||
mTuneLevel = Heavy;
|
||
totalSet++;
|
||
}
|
||
|
||
isSet = (mode_num & MNN_GPU_TUNING_WIDE);
|
||
if(isSet) {
|
||
mTuneLevel = Wide;
|
||
totalSet++;
|
||
}
|
||
|
||
if(totalSet != 1) {
|
||
MNN_PRINT("set multi tuning mode is not permitted, please check cl_mode:%x！\n", mode_num);
|
||
}
|
||
}
|
||
|
||
struct MetalContext {
|
||
std::mutex pLock;
|
||
MNNMetalContext* pContext;
|
||
id<MTLDevice> pDevice;
|
||
};
|
||
static MetalContext* gContext = nullptr;
|
||
MetalRuntime* MetalRuntime::create(const Backend::Info& info) {
|
||
std::unique_lock<std::mutex> _l(gContext->pLock);
|
||
MNNMetalSharedContext sharedContext;
|
||
sharedContext.device = nil;
|
||
sharedContext.queue = nil;
|
||
if (info.user != nullptr) {
|
||
if (info.user->sharedContext != nullptr) {
|
||
sharedContext.device = ((MNNMetalSharedContext*)info.user->sharedContext)->device;
|
||
sharedContext.queue = ((MNNMetalSharedContext*)info.user->sharedContext)->queue;
|
||
}
|
||
}
|
||
if (nil == sharedContext.device) {
|
||
sharedContext.device = MTLCreateSystemDefaultDevice();
|
||
}
|
||
if (nil == gContext->pContext || gContext->pDevice != sharedContext.device) {
|
||
gContext->pContext = [[MNNMetalContext alloc] init];
|
||
gContext->pDevice = sharedContext.device;
|
||
BOOL res = [gContext->pContext initWithSharedContext:&sharedContext dev:sharedContext.device];
|
||
if (!res) {
|
||
gContext->pContext = nil;
|
||
return nullptr;
|
||
}
|
||
}
|
||
auto mContext = (__bridge_retained void *)(gContext->pContext);
|
||
auto rt = new MetalRuntime(mContext);
|
||
rt->setGpuMode(info.gpuMode);
|
||
if (nil != sharedContext.queue) {
|
||
rt->setCommandQueue(sharedContext.queue, true);
|
||
}
|
||
bool supportDefer = info.numThread & MNN_GPU_RECORD_BATCH;
|
||
if ((!supportDefer) && nil == sharedContext.queue) {
|
||
id<MTLCommandQueue> queue = [sharedContext.device newCommandQueue];
|
||
rt->setCommandQueue(queue, false);
|
||
}
|
||
if (nullptr != info.user) {
|
||
rt->mDefaultConfig = *info.user;
|
||
}
|
||
return rt;
|
||
}
|
||
|
||
MetalRuntime::MetalRuntime(void* context) {
|
||
mContext = context;
|
||
auto ctx = (__bridge MNNMetalContext *)mContext;
|
||
std::shared_ptr<EagerBufferAllocator::Allocator> allocator(new MetalRuntimeAllocator([ctx device]));
|
||
// supportsFamily: is available since iOS 13.0 / macOS 10.15, must check before calling
|
||
if (@available(iOS 13.0, macOS 10.15, *)) {
|
||
mSimdGroupReduce = [[ctx device] supportsFamily:MTLGPUFamilyApple7];
|
||
mSimdGroupReduce |= [[ctx device] supportsFamily:(MTLGPUFamily)MTLGPUFamilyMetal3_MNN];
|
||
mSimdGroupMatrix = [[ctx device] supportsFamily:MTLGPUFamilyApple7];
|
||
} else {
|
||
mSimdGroupReduce = false;
|
||
mSimdGroupMatrix = false;
|
||
}
|
||
mMaxThreadSize = [[ctx device] maxThreadsPerThreadgroup].width;
|
||
// Metal4 Support M1/A14 and later chips
|
||
#ifdef MNN_METAL_TENSOR
|
||
if (@available(iOS 13.0, macOS 10.15, *)) {
|
||
mTensorOps = [[ctx device] supportsFamily:(MTLGPUFamily)MTLGPUFamilyMetal4_MNN];
|
||
} else {
|
||
mTensorOps = false;
|
||
}
|
||
|
||
// AI TensorCore device support from M5/A19
|
||
bool noAICoreDevice = [[[ctx device] name] containsString:@"M1"] || \
|
||
[[[ctx device] name] containsString:@"M2"] || \
|
||
[[[ctx device] name] containsString:@"M3"] || \
|
||
[[[ctx device] name] containsString:@"M4"] || \
|
||
[[[ctx device] name] containsString:@"A14"] || \
|
||
[[[ctx device] name] containsString:@"A15"] || \
|
||
[[[ctx device] name] containsString:@"A16"] || \
|
||
[[[ctx device] name] containsString:@"A17"] || \
|
||
[[[ctx device] name] containsString:@"A18"];
|
||
mTensorOps = mTensorOps && !noAICoreDevice;
|
||
#else
|
||
mTensorOps = false;
|
||
#endif
|
||
// MNN_PRINT("Metal device name %s, open tensor: %d\n\n", [[[ctx device] name] UTF8String], mTensorOps);
|
||
mStaticAllocator.reset(new EagerBufferAllocator(allocator));
|
||
mDynamic.resize(METAL_SEPERATE_MAX_COUNT);
|
||
for (auto& buf : mDynamic) {
|
||
buf.root = allocator;
|
||
}
|
||
mTunedInfo = new TunedInfo;
|
||
}
|
||
|
||
MetalRuntime::~ MetalRuntime() {
|
||
if(mContext) {
|
||
CFRelease(mContext);
|
||
}
|
||
delete mTunedInfo;
|
||
}
|
||
|
||
bool MetalRuntime::setCache(std::pair<const void*, size_t> cache) {//Get Cache
|
||
auto buffer = cache.first;
|
||
auto size = cache.second;
|
||
if (nullptr == buffer) {
|
||
mCacheOutside = nullptr;
|
||
mCacheOutsideSize = 0;
|
||
mBuffer.clear();
|
||
return false;//actually get nothing
|
||
}
|
||
mCacheOutsideSize = size;
|
||
mCacheOutside = buffer;
|
||
auto cacheBuffer = GetCache(buffer);
|
||
flatbuffers::Verifier verify((const uint8_t*)cache.first, cache.second);
|
||
if (false == VerifyCacheBuffer(verify)) {
|
||
return false;
|
||
}
|
||
if (nullptr == cacheBuffer->tunings()) {
|
||
return false;
|
||
}
|
||
|
||
// Load Auto Tuning Info
|
||
if (nullptr != cacheBuffer->tunings()) {
|
||
auto tuningInfo = cacheBuffer->tunings();
|
||
for (int i=0; i<tuningInfo->size(); ++i) {
|
||
auto tun = tuningInfo->GetAs<Autotuning>(i);
|
||
if (nullptr == tun->threadSize() || nullptr == tun->groupSize() || nullptr == tun->key()) {
|
||
MNN_ERROR("Error tunning info\n");
|
||
continue;
|
||
}
|
||
std::vector<uint32_t> glo(tun->threadSize()->size());
|
||
for (int v=0; v<glo.size(); ++v) {
|
||
glo[v] = tun->threadSize()->data()[v];
|
||
}
|
||
std::vector<uint32_t> grop(tun->groupNum()->size());
|
||
for (int v=0; v<grop.size(); ++v) {
|
||
grop[v] = tun->groupNum()->data()[v];
|
||
}
|
||
std::vector<uint32_t> loc(tun->groupSize()->size());
|
||
for (int v=0; v<loc.size(); ++v) {
|
||
loc[v] = tun->groupSize()->data()[v];
|
||
}
|
||
uint32_t cost = tun->timeCost();
|
||
mTunedThreadGroup.insert(std::make_pair(std::make_pair(tun->key()->str(), glo), std::make_tuple(grop, loc, cost)));
|
||
mTunedThreadGroupVec[tun->key()->str()].emplace_back(std::make_pair(glo, std::make_tuple(grop, loc, cost)));
|
||
}
|
||
}
|
||
return true;
|
||
}
|
||
|
||
std::pair<const void*, size_t> MetalRuntime::makeCache(TunedInfo* info) {//make Cache
|
||
std::unique_ptr<CacheT> cache(new CacheT);
|
||
// Get All Autotuning cache
|
||
for (auto& iter : mTunedThreadGroup) {
|
||
std::unique_ptr<AutotuningT> tuning(new AutotuningT);
|
||
tuning->key = iter.first.first;
|
||
tuning->threadSize = iter.first.second;
|
||
|
||
tuning->groupNum = std::get<0>(iter.second);
|
||
tuning->groupSize = std::get<1>(iter.second);
|
||
tuning->timeCost = std::get<2>(iter.second);
|
||
|
||
cache->tunings.emplace_back(std::move(tuning));
|
||
}
|
||
cache->tuned = std::move(info->mInfos);
|
||
|
||
flatbuffers::FlatBufferBuilder builder;
|
||
auto lastOffset = Cache::Pack(builder, cache.get());
|
||
builder.Finish(lastOffset);
|
||
mBuffer.resize(builder.GetSize());
|
||
::memcpy(mBuffer.data(), builder.GetBufferPointer(), builder.GetSize());
|
||
return std::make_pair(mBuffer.data(), mBuffer.size());
|
||
}
|
||
|
||
float MetalRuntime::onGetMemoryInMB() {
|
||
auto staticMemoryInMB = mStaticAllocator->totalSize() / 1024.0f / 1024.0f;
|
||
float dynamicMemoryInMB = 0.0f;
|
||
for (auto& buf : mDynamic) {
|
||
dynamicMemoryInMB += buf.currentSize / 1024.0f / 1024.0f;
|
||
}
|
||
return staticMemoryInMB + dynamicMemoryInMB;
|
||
}
|
||
|
||
void MetalRuntime::onMaskOpReady(const std::vector<Tensor*>& inputs, const std::vector<Tensor*>& outputs,
|
||
const MNN::Op* op) {
|
||
if (nullptr != op->name()) {
|
||
auto dstInfo = mTunedInfo;
|
||
std::unique_ptr<MetalCache::OpInfoT> opInfo(new MetalCache::OpInfoT);;
|
||
opInfo->type = op->type();
|
||
opInfo->name = op->name()->str();
|
||
opInfo->inputs.resize(inputs.size());
|
||
for (int v=0; v<opInfo->inputs.size(); ++v) {
|
||
opInfo->inputs[v].reset(new MetalCache::TensorInfoT);
|
||
opInfo->inputs[v]->shape.resize(inputs[v]->dimensions());
|
||
for (int u=0; u<opInfo->inputs[v]->shape.size(); ++u) {
|
||
opInfo->inputs[v]->shape[u] = inputs[v]->length(u);
|
||
}
|
||
}
|
||
opInfo->outputs.resize(outputs.size());
|
||
for (int v=0; v<opInfo->outputs.size(); ++v) {
|
||
opInfo->outputs[v].reset(new MetalCache::TensorInfoT);
|
||
opInfo->outputs[v]->shape.resize(outputs[v]->dimensions());
|
||
for (int u=0; u<opInfo->outputs[v]->shape.size(); ++u) {
|
||
opInfo->outputs[v]->shape[u] = outputs[v]->length(u);
|
||
}
|
||
}
|
||
dstInfo->mInfos.emplace_back(std::move(opInfo));
|
||
}
|
||
}
|
||
static bool _checkTensorInfo(const MetalCache::TensorInfoT* dst, const Tensor* src) {
|
||
if (dst->shape.size() != src->dimensions()) {
|
||
return false;
|
||
}
|
||
for (int j=0; j<dst->shape.size(); ++j) {
|
||
if (dst->shape[j] != src->length(j)) {
|
||
return false;
|
||
}
|
||
}
|
||
return true;
|
||
}
|
||
bool MetalRuntime::onMeasure(const std::vector<Tensor*>& inputs, const std::vector<Tensor*>& outputs,
|
||
const MNN::Op* op, Runtime::OpInfo& dstInfo) const {
|
||
dstInfo.initCostLong = true;
|
||
if (nullptr == op->name()) {
|
||
dstInfo.initCostLong = false;
|
||
return true;
|
||
}
|
||
for(auto& info : mTunedInfo->mInfos) {
|
||
if (info->type != op->type()) {
|
||
continue;
|
||
}
|
||
if (info->name != op->name()->str()) {
|
||
continue;
|
||
}
|
||
if (info->inputs.size() != inputs.size() || info->outputs.size() != outputs.size()) {
|
||
continue;
|
||
}
|
||
bool match = true;
|
||
for (int i=0; i<inputs.size(); ++i) {
|
||
auto& dst = info->inputs[i];
|
||
auto src = inputs[i];
|
||
if (!_checkTensorInfo(dst.get(), src)) {
|
||
match = false;
|
||
break;
|
||
}
|
||
}
|
||
if (!match) {
|
||
continue;
|
||
}
|
||
for (int i=0; i<outputs.size(); ++i) {
|
||
auto& dst = info->outputs[i];
|
||
auto src = outputs[i];
|
||
if (!_checkTensorInfo(dst.get(), src)) {
|
||
match = false;
|
||
break;
|
||
}
|
||
}
|
||
if (match) {
|
||
// All Info is match
|
||
dstInfo.initCostLong = false;
|
||
break;
|
||
}
|
||
}
|
||
return true;
|
||
}
|
||
|
||
class MetalWrapAllocator : public BufferAllocator::Allocator {
|
||
private:
|
||
std::shared_ptr<BufferAllocator::Allocator> mOrigin;
|
||
id<MTLDevice> mDevice;
|
||
public:
|
||
MetalWrapAllocator(std::shared_ptr<BufferAllocator::Allocator> origin, id<MTLDevice> device) : mOrigin(origin), mDevice(device) {}
|
||
virtual ~ MetalWrapAllocator() {
|
||
// Do nothing
|
||
}
|
||
virtual void sync() override {
|
||
mOrigin->sync();
|
||
};
|
||
virtual MemChunk onAlloc(size_t size, size_t align) override {
|
||
auto mem = mOrigin->onAlloc(size, align);
|
||
MNN_ASSERT(mem.second == 0);
|
||
id<MTLBuffer> buffer = [mDevice newBufferWithBytesNoCopy:mem.first length:size options:MTLResourceStorageModeShared deallocator:nil];
|
||
|
||
auto wrap = new MetalRuntimeAllocator::MetalBufferAlloc(buffer);
|
||
return MemChunk((void *)wrap, 0);
|
||
}
|
||
virtual void onRelease(MemChunk chunk) override {
|
||
auto mem = (MetalRuntimeAllocator::MetalBufferAlloc *)chunk.first;
|
||
mOrigin->onRelease(MemChunk(mem->getBuffer().contents));
|
||
delete mem;
|
||
}
|
||
};
|
||
Backend* MetalRuntime::onCreate(const BackendConfig* config, Backend* origin) const {
|
||
if (hint().weightMemoryPath.size() > 0 && mStaticAllocatorMMap.get() == nullptr) {
|
||
// Only support set weightmap dir once
|
||
mStaticAllocatorRaw = mStaticAllocator;
|
||
// forward_type, precision_type, memory_type, power_type
|
||
std::string prefix = "1_0_0_0_";
|
||
std::string posfix = "metal.weight";
|
||
auto syncPath = prefix + "sync." + posfix;
|
||
bool autoRemove = true;
|
||
if (hint().useCachedMmap) {
|
||
autoRemove = false;
|
||
std::string fileName = MNNFilePathConcat(hint().weightMemoryPath, syncPath);
|
||
const_cast<RuntimeHint&>(hint()).useCachedMmap += MNNFileExist(fileName.c_str());
|
||
}
|
||
auto ctx = (__bridge MNNMetalContext *)mContext;
|
||
auto mmap = BufferAllocator::Allocator::createMmap(hint().weightMemoryPath.c_str(), prefix.c_str(), posfix.c_str(), autoRemove);
|
||
std::shared_ptr<BufferAllocator::Allocator> mmapMem(new MetalWrapAllocator(mmap, [ctx device]));
|
||
mStaticAllocator.reset(new EagerBufferAllocator(mmapMem, 32, 1024 * 1024 * 1024));
|
||
mStaticAllocatorMMap = mStaticAllocator;
|
||
}
|
||
BackendConfig::PrecisionMode precision = mDefaultConfig.precision;
|
||
BackendConfig::MemoryMode memory = mDefaultConfig.memory;
|
||
if (nullptr != config) {
|
||
precision = config->precision;
|
||
memory = config->memory;
|
||
}
|
||
bool useFp16AsFp32 = precision != BackendConfig::Precision_High;
|
||
auto backend = new MetalBackend(this, useFp16AsFp32, memory);
|
||
backend->setMetaPtr(pMeta);
|
||
return backend;
|
||
}
|
||
|
||
void MetalRuntime::onGabageCollect(int level) {
|
||
mStaticAllocator->release(false);
|
||
if (nullptr != mStaticAllocatorMMap) {
|
||
mStaticAllocatorMMap->release(false);
|
||
}
|
||
if (level >= 100) {
|
||
for (auto& buf : mDynamic) {
|
||
buf.release();
|
||
}
|
||
}
|
||
}
|
||
|
||
std::pair<const void*, size_t> MetalRuntime::onGetCache() {//make Cache
|
||
return makeCache(mTunedInfo);
|
||
}
|
||
|
||
bool MetalRuntime::onSetCache(const void* buffer, size_t size) {//set Cache
|
||
if (nullptr == buffer) {
|
||
return false;
|
||
}
|
||
auto cacheBuffer = MetalCache::GetCache(buffer);
|
||
flatbuffers::Verifier verify((const uint8_t*)buffer, size);
|
||
if (false == VerifyCacheBuffer(verify)) {
|
||
return false;
|
||
}
|
||
if(nullptr != cacheBuffer->tuned()) {
|
||
for (int i=0; i<cacheBuffer->tuned()->size(); ++i) {
|
||
auto srcInfo = cacheBuffer->tuned()->GetAs<MetalCache::OpInfo>(i);
|
||
std::unique_ptr<MetalCache::OpInfoT> dst(srcInfo->UnPack());
|
||
mTunedInfo->mInfos.emplace_back(std::move(dst));
|
||
}
|
||
}
|
||
return setCache(std::make_pair(buffer, size));
|
||
}
|
||
|
||
MemChunk MetalRuntimeAllocator::onAlloc(size_t size, size_t align) {
|
||
auto buffer = [mDevice newBufferWithLength:size options:MTLCPUCacheModeDefaultCache];
|
||
auto mMetalBufferAlloc = new MetalBufferAlloc(buffer);
|
||
return MemChunk((void *)mMetalBufferAlloc, 0);
|
||
}
|
||
void MetalRuntimeAllocator::onRelease(MemChunk ptr) {
|
||
delete (MetalBufferAlloc *)ptr.first;
|
||
}
|
||
|
||
class MetalRuntimeCreator : public RuntimeCreator {
|
||
public:
|
||
MetalRuntimeCreator() {
|
||
// Do nothing
|
||
}
|
||
virtual ~ MetalRuntimeCreator() {
|
||
// Do nothing
|
||
}
|
||
virtual Runtime *onCreate(const Backend::Info &info) const {
|
||
auto rt = MetalRuntime::create(info);
|
||
return rt;
|
||
}
|
||
private:
|
||
id<MTLDevice> mDevice;
|
||
};
|
||
|
||
void registerMetalRuntimeCreator() {
|
||
// according to
|
||
// https://developer.apple.com/library/archive/documentation/DeviceInformation/Reference/iOSDeviceCompatibility/HardwareGPUInformation/HardwareGPUInformation.html
|
||
// not all device with iOS 8+ supports metal.
|
||
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
|
||
if (nil != device) {
|
||
gContext = new MetalContext;
|
||
gContext->pContext = nil;
|
||
gContext->pDevice = nil;
|
||
registerMetalOps();
|
||
#ifdef MNN_SUPPORT_RENDER
|
||
registerMetalRenderOps();
|
||
#endif
|
||
MNNInsertExtraRuntimeCreator(MNN_FORWARD_METAL, new MetalRuntimeCreator, false);
|
||
} else {
|
||
MNN_ERROR("Init Metal Error\n");
|
||
}
|
||
}
|
||
} // namespace MNN
|
||
#else
|
||
namespace MNN {
|
||
void registerMetalRuntimeCreator() {
|
||
}
|
||
};
|
||
int MNNMetalGetTensorContent(MNNMetalTensorContent* content, void* tensor) {
|
||
return -1;
|
||
}
|
||
|
||
#endif
|