| 52 | namespace { |
| 53 | |
| 54 | std::vector<ComputeTaskDescriptorPtr> SelectConvolution( |
| 55 | const GraphFloat32& graph, int id, ValueId input_id, ValueId output_id, |
| 56 | const Convolution2DAttributes& attr, const metal::RuntimeOptions& options) { |
| 57 | // Special precise version, in case we cover dst_shape poorly with standard |
| 58 | // work group size. |
| 59 | const auto dst_shape = graph.FindOutputs(id)[0]->tensor.shape; |
| 60 | if (GetThreadsRatioUsualToPreciseConvolution(dst_shape) >= 1.2f) { |
| 61 | // Special version for PowerVR >= IPhone6S/SE |
| 62 | // Metal has bad driver for PowerVR in IPhone6, so for Iphone6 we should use |
| 63 | // default kernel with shared memory. |
| 64 | if ((GetAppleSocVersion() == 9 || GetAppleSocVersion() == 10) && |
| 65 | CheckConvolutionPrecise1x1Support(attr)) { |
| 66 | return ConvolutionPrecise1x1PowerVR(id, input_id, output_id, attr, |
| 67 | options); |
| 68 | } |
| 69 | if (GetAppleSocVersion() >= 11 && |
| 70 | GetThreadsRatioUsualToPreciseConvolution(dst_shape) >= 1.2f) { |
| 71 | return ConvolutionPrecise(id, input_id, output_id, attr, options); |
| 72 | } |
| 73 | } |
| 74 | if (GetAppleSocVersion() >= 11) { |
| 75 | if (CheckConvolution1x1Support(attr)) { |
| 76 | return Convolution1x1(id, input_id, output_id, attr, options); |
| 77 | } else { |
| 78 | return ConvolutionGeneric(id, input_id, output_id, attr, options); |
| 79 | } |
| 80 | } else { |
| 81 | return Convolution(id, input_id, output_id, attr, options); |
| 82 | } |
| 83 | } |
| 84 | |
| 85 | std::vector<ComputeTaskDescriptorPtr> SelectDepthWiseConv( |
| 86 | int id, ValueId input_id, ValueId output_id, |
no test coverage detected