2019-04-17 10:49:11 +08:00
|
|
|
//
|
|
|
|
// MetalBackend.hpp
|
|
|
|
// MNN
|
|
|
|
//
|
|
|
|
// Created by MNN on 2019/01/30.
|
|
|
|
// Copyright © 2018, Alibaba Group Holding Limited
|
|
|
|
//
|
|
|
|
|
|
|
|
#ifndef MetalBackend_hpp
|
|
|
|
#define MetalBackend_hpp
|
|
|
|
|
2019-12-27 22:16:57 +08:00
|
|
|
#include "core/Backend.hpp"
|
2021-11-30 10:10:53 +08:00
|
|
|
#include "core/BufferAllocator.hpp"
|
|
|
|
#include "core/TensorUtils.hpp"
|
2019-04-17 10:49:11 +08:00
|
|
|
#include "MNN_generated.h"
|
|
|
|
#include "MetalDefine.h"
|
2020-11-05 16:41:56 +08:00
|
|
|
#include <vector>
|
2021-11-30 10:10:53 +08:00
|
|
|
//#include "MNNMetalContext.h"
|
2021-09-18 15:52:30 +08:00
|
|
|
#include "MetalCache_generated.h"
|
|
|
|
using namespace MetalCache;
|
2019-04-17 10:49:11 +08:00
|
|
|
|
|
|
|
#if MNN_METAL_ENABLED
|
|
|
|
namespace MNN {
|
2021-11-30 10:10:53 +08:00
|
|
|
|
2020-11-05 16:41:56 +08:00
|
|
|
/** MetalRuntime */
|
2021-09-18 15:52:30 +08:00
|
|
|
enum MetalTuneLevel {Never = 0, Heavy = 1, Wide = 2, Normal = 3, Fast = 4};
|
|
|
|
|
2021-11-30 10:10:53 +08:00
|
|
|
class MetalRuntime {
|
2020-11-05 16:41:56 +08:00
|
|
|
public:
|
|
|
|
friend class MetalBackend;
|
2021-11-30 10:10:53 +08:00
|
|
|
MetalRuntime(const Backend::Info info);
|
2020-11-05 16:41:56 +08:00
|
|
|
virtual ~ MetalRuntime();
|
2021-11-30 10:10:53 +08:00
|
|
|
|
2020-11-05 16:41:56 +08:00
|
|
|
void *context() const {
|
|
|
|
return mContext;
|
|
|
|
}
|
2021-11-30 10:10:53 +08:00
|
|
|
|
|
|
|
bool isCreateError() const {
|
|
|
|
return mIsCreateError;
|
|
|
|
}
|
|
|
|
void setGpuMode(const int cl_mode_num);
|
2021-09-18 15:52:30 +08:00
|
|
|
|
2021-11-30 10:10:53 +08:00
|
|
|
std::pair<const void*, size_t> makeCache();
|
|
|
|
bool setCache(std::pair<const void*, size_t> cache);
|
2021-09-18 15:52:30 +08:00
|
|
|
|
|
|
|
MetalTuneLevel getTuneLevel() {
|
|
|
|
return mTuneLevel;
|
|
|
|
}
|
2021-11-30 10:10:53 +08:00
|
|
|
std::map<std::pair<std::string, std::vector<uint32_t>>, std::tuple<std::vector<uint32_t>, std::vector<uint32_t>, uint32_t>>& getTunedThreadGroup() {
|
|
|
|
return mTunedThreadGroup;
|
|
|
|
};
|
2020-11-05 16:41:56 +08:00
|
|
|
private:
|
|
|
|
void* mContext = nullptr;
|
|
|
|
std::shared_ptr<BufferAllocator> mStatic;
|
2021-11-30 10:10:53 +08:00
|
|
|
bool mIsCreateError = false;
|
|
|
|
MetalTuneLevel mTuneLevel = Wide;
|
|
|
|
std::map<std::pair<std::string, std::vector<uint32_t>>, std::tuple<std::vector<uint32_t>, std::vector<uint32_t>, uint32_t>> mTunedThreadGroup;
|
|
|
|
|
|
|
|
private:
|
2021-09-18 15:52:30 +08:00
|
|
|
std::vector<uint8_t> mBuffer;
|
|
|
|
const void* mCacheOutside = nullptr;
|
|
|
|
size_t mCacheOutsideSize = 0;
|
2021-11-30 10:10:53 +08:00
|
|
|
};
|
|
|
|
|
|
|
|
|
|
|
|
class MetalRuntimeWrapper : public Runtime {
|
|
|
|
public:
|
|
|
|
MetalRuntimeWrapper(const Backend::Info info);
|
|
|
|
virtual ~MetalRuntimeWrapper();
|
|
|
|
virtual Backend *onCreate(const BackendConfig* config) const override;
|
|
|
|
virtual void onGabageCollect(int level) override;
|
|
|
|
bool isCreateError() const {
|
|
|
|
return mIsCreateError;
|
|
|
|
}
|
|
|
|
virtual CompilerType onGetCompilerType() const override {
|
|
|
|
return Compiler_Loop;
|
|
|
|
}
|
|
|
|
virtual float onGetMemoryInMB() override;
|
|
|
|
|
|
|
|
virtual std::pair<const void*, size_t> onGetCache() override;
|
|
|
|
virtual bool onSetCache(const void* buffer, size_t size) override;
|
|
|
|
|
|
|
|
|
|
|
|
private:
|
|
|
|
std::shared_ptr<BufferAllocator> mBufferPool;
|
|
|
|
std::shared_ptr<MetalRuntime> mMetalRuntime;
|
|
|
|
bool mIsCreateError{false};
|
|
|
|
};
|
|
|
|
|
|
|
|
class MetalRuntimeAllocator : public BufferAllocator::Allocator {
|
|
|
|
public:
|
|
|
|
class MetalBufferAlloc {
|
|
|
|
public:
|
|
|
|
MetalBufferAlloc(id<MTLBuffer> buffer) {
|
|
|
|
mBuffer = buffer;
|
|
|
|
}
|
|
|
|
id<MTLBuffer> getBuffer() {
|
|
|
|
return mBuffer;
|
|
|
|
}
|
|
|
|
~MetalBufferAlloc(){};
|
|
|
|
private:
|
|
|
|
id<MTLBuffer> mBuffer = nil;
|
|
|
|
};
|
|
|
|
|
|
|
|
MetalRuntimeAllocator(MetalRuntime *rt): mMetalRuntime(rt) {
|
|
|
|
// Do nothing
|
|
|
|
}
|
|
|
|
virtual ~ MetalRuntimeAllocator() = default;
|
|
|
|
virtual std::pair<void*, int> onAlloc(int size, int align) override;
|
|
|
|
virtual void onRelease(std::pair<void*, int> ptr) override;
|
|
|
|
|
|
|
|
private:
|
|
|
|
MetalRuntime *mMetalRuntime;
|
|
|
|
id<MTLBuffer> mBuffer = nil;
|
2020-11-05 16:41:56 +08:00
|
|
|
};
|
2019-04-17 10:49:11 +08:00
|
|
|
|
|
|
|
/** Metal backend */
|
2021-11-30 10:10:53 +08:00
|
|
|
class MetalBackend : public Backend {
|
2019-04-17 10:49:11 +08:00
|
|
|
public:
|
|
|
|
/** Metal execution creator */
|
|
|
|
class Creator {
|
|
|
|
public:
|
|
|
|
/**
|
|
|
|
* @brief create execution for given input, op on metal backend.
|
|
|
|
* @param inputs given input tensors.
|
|
|
|
* @param op given op.
|
|
|
|
* @param backend metal backend.
|
|
|
|
* @return created execution if supported, NULL otherwise.
|
|
|
|
*/
|
- dynamic computation graph (beta)
- add supports (/express)
- add tests
- add benchmarks with it (/benchmark/exprModels)
- Python
- MNN engine and tools were submitted to pip
- available on Windows/macOS/Linux
- Engine/Converter
- add supports for each op benchmarking
- refactor optimizer by separating steps
- CPU
- add supports for Conv3D, Pool3D, ELU, ReverseSequence
- fix ArgMax, Permute, Scale, BinaryOp, Slice, SliceTf
- OpenCL
- add half transform in CPU
- add broadcast supports for binary
- optimize Conv2D, Reshape, Eltwise, Gemm, etc.
- OpenGL
- add sub, real div supports for binary
- add supports for unary
- optimize Conv2D, Reshape
- Vulkan
- add max supports for eltwise
- Metal
- fix metallib missing problem
- Train/Quantization
- use express to refactor training codes
2019-09-26 21:02:07 +08:00
|
|
|
virtual Execution *onCreate(const std::vector<Tensor *> &inputs, const MNN::Op *op, Backend *backend) const = 0;
|
2019-04-17 10:49:11 +08:00
|
|
|
};
|
|
|
|
/**
|
|
|
|
* @brief register creator for given op type.
|
|
|
|
* @param type given op type.
|
|
|
|
* @param creator registering creator.
|
|
|
|
*/
|
|
|
|
static void addCreator(OpType type, Creator *creator);
|
|
|
|
|
2021-11-30 10:10:53 +08:00
|
|
|
id<MTLBuffer> getHostBuffer(size_t size) const;
|
|
|
|
id<MTLBuffer> getConstBuffer(size_t size) const;
|
|
|
|
public:
|
|
|
|
MetalBackend(std::shared_ptr<BufferAllocator> staticMem, const MetalRuntime* runtime);
|
|
|
|
virtual ~MetalBackend();
|
2020-11-05 16:41:56 +08:00
|
|
|
const MetalRuntime* runtime() const {
|
|
|
|
return mRuntime;
|
|
|
|
}
|
2021-11-30 10:10:53 +08:00
|
|
|
|
|
|
|
virtual Backend::MemObj* onAcquire(const Tensor *Tensor, StorageType storageType) override;
|
2019-04-17 10:49:11 +08:00
|
|
|
virtual bool onClearBuffer() override;
|
|
|
|
virtual void onCopyBuffer(const Tensor *srcTensor, const Tensor *dstTensor) const override;
|
|
|
|
|
|
|
|
virtual Execution *onCreate(const std::vector<Tensor *> &inputs, const std::vector<Tensor *> &outputs,
|
|
|
|
const MNN::Op *op) override;
|
2021-04-28 18:02:10 +08:00
|
|
|
|
|
|
|
virtual void onResizeBegin() override;
|
|
|
|
virtual void onResizeEnd() override;
|
2019-04-17 10:49:11 +08:00
|
|
|
virtual void onExecuteBegin() const override;
|
|
|
|
virtual void onExecuteEnd() const override;
|
- dynamic computation graph (beta)
- add supports (/express)
- add tests
- add benchmarks with it (/benchmark/exprModels)
- Python
- MNN engine and tools were submitted to pip
- available on Windows/macOS/Linux
- Engine/Converter
- add supports for each op benchmarking
- refactor optimizer by separating steps
- CPU
- add supports for Conv3D, Pool3D, ELU, ReverseSequence
- fix ArgMax, Permute, Scale, BinaryOp, Slice, SliceTf
- OpenCL
- add half transform in CPU
- add broadcast supports for binary
- optimize Conv2D, Reshape, Eltwise, Gemm, etc.
- OpenGL
- add sub, real div supports for binary
- add supports for unary
- optimize Conv2D, Reshape
- Vulkan
- add max supports for eltwise
- Metal
- fix metallib missing problem
- Train/Quantization
- use express to refactor training codes
2019-09-26 21:02:07 +08:00
|
|
|
virtual std::pair<float, bool> onMeasure(const std::vector<Tensor*>& inputs, const std::vector<Tensor*>& outputs,
|
|
|
|
const MNN::Op* op) override;
|
2019-04-17 10:49:11 +08:00
|
|
|
|
|
|
|
public:
|
|
|
|
/**
|
|
|
|
* @brief get metal context object
|
|
|
|
* @return metal context object pointer
|
|
|
|
*/
|
2020-11-05 16:41:56 +08:00
|
|
|
void *context() const;
|
2019-04-17 10:49:11 +08:00
|
|
|
|
|
|
|
/**
|
|
|
|
* @brief copy buffer content to dest tensor
|
|
|
|
* @param srcTensor source tensor
|
|
|
|
* @param dstTensor destined tensor
|
|
|
|
* @param encoder command encoder
|
|
|
|
*/
|
2020-11-13 14:27:18 +08:00
|
|
|
void onCopyBuffer(const Tensor *srcTensor, const Tensor *dstTensor,
|
|
|
|
id<MTLComputeCommandEncoder> encoder, id<MTLBuffer> shape) const;
|
2019-04-17 10:49:11 +08:00
|
|
|
|
2020-11-13 14:27:18 +08:00
|
|
|
void flushEncoder() const;
|
|
|
|
id<MTLComputeCommandEncoder> encoder() const;
|
2021-04-28 18:02:10 +08:00
|
|
|
void addOpEncoder(std::function<void(void)> opEncoder);
|
|
|
|
|
2021-09-18 15:52:30 +08:00
|
|
|
bool isCommandEncoderSet();
|
2021-04-28 18:02:10 +08:00
|
|
|
void setOpEncoder() const;
|
2021-11-30 10:10:53 +08:00
|
|
|
|
|
|
|
BufferAllocator *getBufferPool() const {
|
|
|
|
return mBufferPool.get();
|
|
|
|
}
|
|
|
|
BufferAllocator *getStaticBufferPool() const {
|
|
|
|
return mStaticBufferPool.get();
|
|
|
|
}
|
|
|
|
|
2021-09-18 15:52:30 +08:00
|
|
|
bool isCmdBufferCommit();
|
|
|
|
|
2019-04-17 10:49:11 +08:00
|
|
|
private:
|
2020-11-05 16:41:56 +08:00
|
|
|
const MetalRuntime* mRuntime;
|
|
|
|
std::vector<id<MTLBuffer>> mHoldBuffers;
|
2021-11-30 10:10:53 +08:00
|
|
|
id<MTLBuffer> mShapeH2D;
|
|
|
|
id<MTLBuffer> mShapeD2H;
|
2021-09-18 15:52:30 +08:00
|
|
|
mutable NSUInteger mEncoderCount = 0;
|
|
|
|
mutable bool mOpEncoderSet = false;//whether has set encoder
|
2021-04-28 18:02:10 +08:00
|
|
|
mutable bool mOpFullSupport = true;
|
|
|
|
mutable bool mFrameEncodeCache = false;
|
|
|
|
|
|
|
|
std::vector<std::function<void(void)>> mOpEncoders;
|
2020-11-13 14:27:18 +08:00
|
|
|
mutable id<MTLComputeCommandEncoder> mComputeEncoder = nil;
|
2021-11-30 10:10:53 +08:00
|
|
|
std::shared_ptr<BufferAllocator> mBufferPool;
|
|
|
|
std::shared_ptr<BufferAllocator> mStaticBufferPool;
|
2019-04-17 10:49:11 +08:00
|
|
|
|
|
|
|
private:
|
2021-11-30 10:10:53 +08:00
|
|
|
mutable id<MTLBuffer> mHostBuffer = nullptr;
|
2019-04-17 10:49:11 +08:00
|
|
|
void onCopyHostToDevice(const Tensor *src, const Tensor *dst) const;
|
|
|
|
void onCopyDeviceToHost(const Tensor *src, const Tensor *dst) const;
|
2020-11-13 14:27:18 +08:00
|
|
|
void onCopyDeviceToDevice(const Tensor *src, const Tensor *dst, id<MTLComputeCommandEncoder> encoder, id<MTLBuffer> shape) const;
|
2019-04-17 10:49:11 +08:00
|
|
|
};
|
|
|
|
|
2020-11-05 16:41:56 +08:00
|
|
|
|
2019-04-17 10:49:11 +08:00
|
|
|
/** Metal creator register */
|
|
|
|
template <class T>
|
|
|
|
class MetalCreatorRegister {
|
|
|
|
public:
|
|
|
|
/**
|
|
|
|
* @brief initializer. register T creator for given op type.
|
|
|
|
* @param type given op type.
|
|
|
|
*/
|
|
|
|
MetalCreatorRegister(OpType type) {
|
|
|
|
T *test = new T;
|
|
|
|
MetalBackend::addCreator(type, test);
|
|
|
|
}
|
|
|
|
};
|
|
|
|
} // namespace MNN
|
|
|
|
|
2019-05-08 15:44:57 +08:00
|
|
|
#define REGISTER_METAL_OP_CREATOR(name, opType) \
|
|
|
|
void ___##name##__##opType##__() { \
|
|
|
|
MetalBackend::addCreator(opType, new name); \
|
|
|
|
}
|
|
|
|
|
2019-04-17 10:49:11 +08:00
|
|
|
#endif /* MNN_METAL_ENABLED */
|
|
|
|
#endif /* MetalBackend_hpp */
|