2019-04-17 10:49:11 +08:00
|
|
|
//
|
|
|
|
|
// MetalMatMul.hpp
|
|
|
|
|
// MNN
|
|
|
|
|
//
|
|
|
|
|
// Created by MNN on 2019/01/30.
|
|
|
|
|
// Copyright © 2018, Alibaba Group Holding Limited
|
|
|
|
|
//
|
|
|
|
|
|
|
|
|
|
#ifndef MetalMatMul_hpp
|
|
|
|
|
#define MetalMatMul_hpp
|
|
|
|
|
|
2023-12-27 17:26:44 +08:00
|
|
|
#import "MetalExecution.hpp"
|
2019-04-17 10:49:11 +08:00
|
|
|
#import "MNN_generated.h"
|
2020-11-05 16:41:56 +08:00
|
|
|
#import "MetalBackend.hpp"
|
2019-04-17 10:49:11 +08:00
|
|
|
|
|
|
|
|
#if MNN_METAL_ENABLED
|
|
|
|
|
namespace MNN {
|
|
|
|
|
|
2023-12-27 17:26:44 +08:00
|
|
|
class MetalMatMul : public MetalExecution {
|
2019-04-17 10:49:11 +08:00
|
|
|
public:
|
2024-05-11 19:17:02 +08:00
|
|
|
MetalMatMul(Backend *backend, const MatMul *matmul, bool withBias);
|
2024-04-19 11:58:21 +08:00
|
|
|
virtual ~MetalMatMul();
|
2023-12-27 17:26:44 +08:00
|
|
|
virtual void onEncode(const std::vector<Tensor *> &inputs, const std::vector<Tensor *> &outputs, id<MTLComputeCommandEncoder> encoder) override;
|
2020-11-05 16:41:56 +08:00
|
|
|
virtual ErrorCode onResize(const std::vector<Tensor *> &inputs, const std::vector<Tensor *> &outputs) override;
|
2019-04-17 10:49:11 +08:00
|
|
|
|
|
|
|
|
private:
|
2021-11-30 10:10:53 +08:00
|
|
|
id<MTLBuffer> mConstBuffer = nil;
|
2020-11-05 16:41:56 +08:00
|
|
|
bool mTransposeA = false;
|
|
|
|
|
bool mTransposeB = false;
|
2024-05-11 19:17:02 +08:00
|
|
|
id<MTLComputePipelineState> mPipeline;
|
|
|
|
|
std::pair<MTLSize, MTLSize> mThreads;
|
2019-04-17 10:49:11 +08:00
|
|
|
};
|
|
|
|
|
|
|
|
|
|
} // namespace MNN
|
|
|
|
|
#endif /* MNN_METAL_ENABLED */
|
|
|
|
|
#endif /* MetalMatMul_hpp */
|