//
//  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 */