wehub-resource-sync
46 lines
1.8 KiB
C++
46 lines
1.8 KiB
C++
//
|
|
// MetalConvolutionDepthwise.hpp
|
|
// MNN
|
|
//
|
|
// Created by MNN on 2019/02/25.
|
|
// Copyright © 2018, Alibaba Group Holding Limited
|
|
//
|
|
|
|
#ifndef MetalConvolutionDepthwise_hpp
|
|
#define MetalConvolutionDepthwise_hpp
|
|
|
|
#import "MetalConvolutionCommon.hpp"
|
|
|
|
#if MNN_METAL_ENABLED
|
|
namespace MNN {
|
|
|
|
class MetalConvolutionDepthwise : public MetalConvolutionCommon {
|
|
public:
|
|
MetalConvolutionDepthwise(Backend *backend, const MNN::Op *op);
|
|
MetalConvolutionDepthwise(Backend *backend, const MNN::Op *op, bool dynamicWeight);
|
|
MetalConvolutionDepthwise(Backend *backend, const MNN::Op *op, std::shared_ptr<MNN::Tensor> weight,
|
|
std::shared_ptr<MNN::Tensor> bias);
|
|
virtual ~MetalConvolutionDepthwise() = 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:
|
|
bool mDynamicWeight = false;
|
|
id<MTLComputePipelineState> mWeightTransformPipeline = nil;
|
|
id<MTLComputePipelineState> mBiasTransformPipeline = nil;
|
|
id<MTLBuffer> mWeightTransformConstBuffer = nil;
|
|
id<MTLBuffer> mBiasTransformConstBuffer = nil;
|
|
std::pair<MTLSize, MTLSize> mWeightTransformThreads;
|
|
std::pair<MTLSize, MTLSize> mBiasTransformThreads;
|
|
id<MTLComputePipelineState> mPipeline;
|
|
std::pair<MTLSize, MTLSize> mThreads;
|
|
};
|
|
|
|
} // namespace MNN
|
|
#endif /* MNN_METAL_ENABLED */
|
|
|
|
#endif /* MetalConvolutionDepthwise_hpp */
|