61 lines
1.8 KiB
C++
61 lines
1.8 KiB
C++
//
|
|
// MetalConvolutionWinograd.hpp
|
|
// MNN
|
|
//
|
|
// Created by MNN on 2019/01/31.
|
|
// Copyright © 2018, Alibaba Group Holding Limited
|
|
//
|
|
|
|
#ifndef MetalConvolutionWinograd_hpp
|
|
#define MetalConvolutionWinograd_hpp
|
|
|
|
#import "MetalConvolutionCommon.hpp"
|
|
|
|
#if MNN_METAL_ENABLED
|
|
namespace MNN {
|
|
struct TransformBuffer {
|
|
int inputSize[4];
|
|
int outputSize[4];
|
|
int padX;
|
|
int padY;
|
|
int unitWidth;
|
|
int unitHeight;
|
|
int unit;
|
|
int activation;
|
|
int remain[2];
|
|
};
|
|
class MetalConvolutionWinograd : public MetalConvolutionCommon {
|
|
public:
|
|
static bool isValid(Backend *backend, const Convolution2D *conv, const Tensor *input, const Tensor* output);
|
|
MetalConvolutionWinograd(Backend *backend, const MNN::Op *op);
|
|
virtual ~MetalConvolutionWinograd() = default;
|
|
virtual ErrorCode onResize(const std::vector<Tensor *> &inputs, const std::vector<Tensor *> &outputs) override;
|
|
|
|
virtual bool onClone(Backend* bn, const Op* op, Execution** dst) override;
|
|
|
|
protected:
|
|
virtual void onEncode(const std::vector<Tensor *> &inputs, const std::vector<Tensor *> &outputs, id<MTLComputeCommandEncoder> encoder) override;
|
|
virtual std::shared_ptr<MNN::Tensor> weightTransform(int group, int oc, int ic, int kh, int kw, const float *src, bool int8Weight=false, bool int4Weight=false, id<MTLBuffer> srcGpuBuffer=nil, int subBits=0) override;
|
|
|
|
private:
|
|
MetalConvolutionWinograd(Backend *backend, const MNN::Op *op, std::shared_ptr<Tensor> weight, std::shared_ptr<Tensor> bias);
|
|
|
|
|
|
id<MTLBuffer> mShapeBuffer = nil;
|
|
|
|
int mSrcUnit;
|
|
int mDstUnit;
|
|
|
|
std::shared_ptr<Tensor> mTempSrc;
|
|
std::shared_ptr<Tensor> mTempDst;
|
|
|
|
MTLSize mInputTransformThreads;
|
|
MTLSize mMatMulThreads;
|
|
MTLSize mOutputTransformThreads;
|
|
int mSplitNum = 1;
|
|
};
|
|
|
|
} // namespace MNN
|
|
#endif /* MNN_METAL_ENABLED */
|
|
#endif /* MetalConvolutionWinograd_hpp */
|