Merge pull request #850 from codeWorm2015/metal

 cpu nms

metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/project.pbxproj

@@ -15,8 +15,11 @@
 		FC039B8920E11C560081E9F8 /* Assets.xcassets in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8820E11C560081E9F8 /* Assets.xcassets */; };
 		FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */ = {isa = PBXBuildFile; fileRef = FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */; };
 		FC27991021341CE5000B6BAD /* Net.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC27990F21341CE5000B6BAD /* Net.swift */; };
+		FC27991321343A3A000B6BAD /* CPUCompute.mm in Sources */ = {isa = PBXBuildFile; fileRef = FC27991221343A3A000B6BAD /* CPUCompute.mm */; };
 		FC3C800F2133F46600D1295E /* MobileNetSSD.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3C800E2133F46600D1295E /* MobileNetSSD.swift */; };
 		FC3C80112133F4AB00D1295E /* MobileNet.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3C80102133F4AB00D1295E /* MobileNet.swift */; };
+		FC8CFEDB21351F5D0094D569 /* params in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFED921351F5D0094D569 /* params */; };
+		FC8CFEDC21351F5D0094D569 /* model in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEDA21351F5D0094D569 /* model */; };
 		FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */ = {isa = PBXBuildFile; fileRef = FC918190211DBC3500B6F354 /* paddle-mobile.png */; };
 		FC918193211DC70500B6F354 /* iphone.JPG in Resources */ = {isa = PBXBuildFile; fileRef = FC918192211DC70500B6F354 /* iphone.JPG */; };
 		FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */ = {isa = PBXBuildFile; fileRef = FCA3A16021313E1F00084FE5 /* hand.jpg */; };
@@ -58,8 +61,13 @@
 		FC039B8B20E11C560081E9F8 /* Base */ = {isa = PBXFileReference; lastKnownFileType = file.storyboard; name = Base; path = Base.lproj/LaunchScreen.storyboard; sourceTree = "<group>"; };
 		FC039B8D20E11C560081E9F8 /* Info.plist */ = {isa = PBXFileReference; lastKnownFileType = text.plist.xml; path = Info.plist; sourceTree = "<group>"; };
 		FC27990F21341CE5000B6BAD /* Net.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Net.swift; sourceTree = "<group>"; };
+		FC27991121343A39000B6BAD /* paddle-mobile-demo-Bridging-Header.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = "paddle-mobile-demo-Bridging-Header.h"; sourceTree = "<group>"; };
+		FC27991221343A3A000B6BAD /* CPUCompute.mm */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.objcpp; path = CPUCompute.mm; sourceTree = "<group>"; };
+		FC27991421343A46000B6BAD /* CPUCompute.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = CPUCompute.h; sourceTree = "<group>"; };
 		FC3C800E2133F46600D1295E /* MobileNetSSD.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobileNetSSD.swift; sourceTree = "<group>"; };
 		FC3C80102133F4AB00D1295E /* MobileNet.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobileNet.swift; sourceTree = "<group>"; };
+		FC8CFED921351F5D0094D569 /* params */ = {isa = PBXFileReference; lastKnownFileType = file; path = params; sourceTree = "<group>"; };
+		FC8CFEDA21351F5D0094D569 /* model */ = {isa = PBXFileReference; lastKnownFileType = file; path = model; sourceTree = "<group>"; };
 		FC918190211DBC3500B6F354 /* paddle-mobile.png */ = {isa = PBXFileReference; lastKnownFileType = image.png; path = "paddle-mobile.png"; sourceTree = "<group>"; };
 		FC918192211DC70500B6F354 /* iphone.JPG */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = iphone.JPG; sourceTree = "<group>"; };
 		FCA3A16021313E1F00084FE5 /* hand.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = hand.jpg; sourceTree = "<group>"; };
@@ -125,6 +133,7 @@
 		FC039B8020E11C550081E9F8 /* paddle-mobile-demo */ = {
 			isa = PBXGroup;
 			children = (
+				FC8CFED2213519540094D569 /* Net */,
 				FC0E2C2020EDC03B009C1FAC /* models */,
 				FC0E2C1D20EDC030009C1FAC /* images */,
 				FC039B8120E11C550081E9F8 /* AppDelegate.swift */,
@@ -134,10 +143,7 @@
 				FC039B8820E11C560081E9F8 /* Assets.xcassets */,
 				FC039B8A20E11C560081E9F8 /* LaunchScreen.storyboard */,
 				FC039B8D20E11C560081E9F8 /* Info.plist */,
-				FCBCCC542122EF5400D94F7E /* MetalHelper.swift */,
-				FC3C800E2133F46600D1295E /* MobileNetSSD.swift */,
-				FC3C80102133F4AB00D1295E /* MobileNet.swift */,
-				FC27990F21341CE5000B6BAD /* Net.swift */,
+				FC27991121343A39000B6BAD /* paddle-mobile-demo-Bridging-Header.h */,
 			);
 			path = "paddle-mobile-demo";
 			sourceTree = "<group>";
@@ -158,6 +164,7 @@
 		FC0E2C2020EDC03B009C1FAC /* models */ = {
 			isa = PBXGroup;
 			children = (
+				FC8CFED821351F5D0094D569 /* enet */,
 				FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */,
 				FCD04E6020F3146A0007374F /* mobilenet */,
 			);
@@ -165,6 +172,28 @@
 			path = ../../models;
 			sourceTree = "<group>";
 		};
+		FC8CFED2213519540094D569 /* Net */ = {
+			isa = PBXGroup;
+			children = (
+				FCBCCC542122EF5400D94F7E /* MetalHelper.swift */,
+				FC3C800E2133F46600D1295E /* MobileNetSSD.swift */,
+				FC3C80102133F4AB00D1295E /* MobileNet.swift */,
+				FC27990F21341CE5000B6BAD /* Net.swift */,
+				FC27991221343A3A000B6BAD /* CPUCompute.mm */,
+				FC27991421343A46000B6BAD /* CPUCompute.h */,
+			);
+			path = Net;
+			sourceTree = "<group>";
+		};
+		FC8CFED821351F5D0094D569 /* enet */ = {
+			isa = PBXGroup;
+			children = (
+				FC8CFED921351F5D0094D569 /* params */,
+				FC8CFEDA21351F5D0094D569 /* model */,
+			);
+			path = enet;
+			sourceTree = "<group>";
+		};
 		FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */ = {
 			isa = PBXGroup;
 			children = (
@@ -218,6 +247,7 @@
 				TargetAttributes = {
 					FC039B7D20E11C550081E9F8 = {
 						CreatedOnToolsVersion = 9.3.1;
+						LastSwiftMigration = 0940;
 					};
 				};
 			};
@@ -244,7 +274,9 @@
 			isa = PBXResourcesBuildPhase;
 			buildActionMask = 2147483647;
 			files = (
+				FC8CFEDB21351F5D0094D569 /* params in Resources */,
 				FCD04E6320F3146B0007374F /* params in Resources */,
+				FC8CFEDC21351F5D0094D569 /* model in Resources */,
 				FC039B8C20E11C560081E9F8 /* LaunchScreen.storyboard in Resources */,
 				FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */,
 				FC039B8920E11C560081E9F8 /* Assets.xcassets in Resources */,
@@ -309,6 +341,7 @@
 				FC013928210204A3008100E3 /* PreProcessKernel.metal in Sources */,
 				FC27991021341CE5000B6BAD /* Net.swift in Sources */,
 				FCBCCC552122EF5500D94F7E /* MetalHelper.swift in Sources */,
+				FC27991321343A3A000B6BAD /* CPUCompute.mm in Sources */,
 				FC3C80112133F4AB00D1295E /* MobileNet.swift in Sources */,
 				FC3C800F2133F46600D1295E /* MobileNetSSD.swift in Sources */,
 				FC039B8220E11C550081E9F8 /* AppDelegate.swift in Sources */,
@@ -456,19 +489,22 @@
 			baseConfigurationReference = 878829884E1A14D7044721D5 /* Pods-paddle-mobile-demo.debug.xcconfig */;
 			buildSettings = {
 				ASSETCATALOG_COMPILER_APPICON_NAME = AppIcon;
-				CODE_SIGN_IDENTITY = "iPhone Developer";
-				CODE_SIGN_STYLE = Automatic;
-				DEVELOPMENT_TEAM = A798K58VVL;
+				CLANG_ENABLE_MODULES = YES;
+				CODE_SIGN_IDENTITY = "iPhone Distribution";
+				CODE_SIGN_STYLE = Manual;
+				DEVELOPMENT_TEAM = 6T9LLJKSM4;
 				INFOPLIST_FILE = "paddle-mobile-demo/Info.plist";
 				IPHONEOS_DEPLOYMENT_TARGET = 9.0;
 				LD_RUNPATH_SEARCH_PATHS = (
 					"$(inherited)",
 					"@executable_path/Frameworks",
 				);
-				PRODUCT_BUNDLE_IDENTIFIER = com.paddlemobile.metal;
+				PRODUCT_BUNDLE_IDENTIFIER = com.baidu.mms.qa;
 				PRODUCT_NAME = "$(TARGET_NAME)";
-				PROVISIONING_PROFILE = "";
-				PROVISIONING_PROFILE_SPECIFIER = "";
+				PROVISIONING_PROFILE = "ba9c4b24-7bd0-49c5-93cd-e3051e775d6c";
+				PROVISIONING_PROFILE_SPECIFIER = Distribution_MMS;
+				SWIFT_OBJC_BRIDGING_HEADER = "paddle-mobile-demo/paddle-mobile-demo-Bridging-Header.h";
+				SWIFT_OPTIMIZATION_LEVEL = "-Onone";
 				SWIFT_VERSION = 4.0;
 				TARGETED_DEVICE_FAMILY = "1,2";
 			};
@@ -479,19 +515,21 @@
 			baseConfigurationReference = 081C9CF10DB06C58B8B6B039 /* Pods-paddle-mobile-demo.release.xcconfig */;
 			buildSettings = {
 				ASSETCATALOG_COMPILER_APPICON_NAME = AppIcon;
-				CODE_SIGN_IDENTITY = "iPhone Developer";
-				CODE_SIGN_STYLE = Automatic;
-				DEVELOPMENT_TEAM = A798K58VVL;
+				CLANG_ENABLE_MODULES = YES;
+				CODE_SIGN_IDENTITY = "iPhone Distribution";
+				CODE_SIGN_STYLE = Manual;
+				DEVELOPMENT_TEAM = 6T9LLJKSM4;
 				INFOPLIST_FILE = "paddle-mobile-demo/Info.plist";
 				IPHONEOS_DEPLOYMENT_TARGET = 9.0;
 				LD_RUNPATH_SEARCH_PATHS = (
 					"$(inherited)",
 					"@executable_path/Frameworks",
 				);
-				PRODUCT_BUNDLE_IDENTIFIER = com.paddlemobile.metal;
+				PRODUCT_BUNDLE_IDENTIFIER = com.baidu.mms.qa;
 				PRODUCT_NAME = "$(TARGET_NAME)";
-				PROVISIONING_PROFILE = "";
-				PROVISIONING_PROFILE_SPECIFIER = "";
+				PROVISIONING_PROFILE = "ba9c4b24-7bd0-49c5-93cd-e3051e775d6c";
+				PROVISIONING_PROFILE_SPECIFIER = Distribution_MMS;
+				SWIFT_OBJC_BRIDGING_HEADER = "paddle-mobile-demo/paddle-mobile-demo-Bridging-Header.h";
 				SWIFT_VERSION = 4.0;
 				TARGETED_DEVICE_FAMILY = "1,2";
 			};

metal/paddle-mobile-demo/paddle-mobile-demo.xcodeproj/xcuserdata/liuruilong.xcuserdatad/xcschemes/paddle-mobile-demo.xcscheme

@@ -42,7 +42,7 @@
       </AdditionalOptions>
    </TestAction>
    <LaunchAction
-      buildConfiguration = "Debug"
+      buildConfiguration = "Release"
       selectedDebuggerIdentifier = "Xcode.DebuggerFoundation.Debugger.LLDB"
       selectedLauncherIdentifier = "Xcode.DebuggerFoundation.Launcher.LLDB"
       launchStyle = "0"

metal/paddle-mobile-demo/paddle-mobile-demo/Net/CPUCompute.h

+/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
+ 
+ Licensed under the Apache License, Version 2.0 (the "License");
+ you may not use this file except in compliance with the License.
+ You may obtain a copy of the License at
+ 
+ http://www.apache.org/licenses/LICENSE-2.0
+ 
+ Unless required by applicable law or agreed to in writing, software
+ distributed under the License is distributed on an "AS IS" BASIS,
+ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ See the License for the specific language governing permissions and
+ limitations under the License. */
+
+#pragma once
+
+#import <Foundation/Foundation.h>
+
+@interface NMSCompute: NSObject
+
+@property (assign, nonatomic) float scoreThredshold;
+
+@property (assign, nonatomic) int nmsTopK;
+
+@property (assign, nonatomic) int keepTopK;
+
+@property (assign, nonatomic) float nmsEta;
+
+@property (assign, nonatomic) float nmsThreshold;
+
+@property (assign, nonatomic) int background_label;
+
+@property (strong, nonatomic) NSArray<NSNumber *> *scoreDim;
+
+@property (strong, nonatomic) NSArray<NSNumber *> *bboxDim;
+
+-(NSArray<NSNumber *> *)computeWithScore:(float *)score andBBoxs:(float *)bbox;
+
+@end

metal/paddle-mobile-demo/paddle-mobile-demo/Net/CPUCompute.mm

+/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
+ 
+ Licensed under the Apache License, Version 2.0 (the "License");
+ you may not use this file except in compliance with the License.
+ You may obtain a copy of the License at
+ 
+ http://www.apache.org/licenses/LICENSE-2.0
+ 
+ Unless required by applicable law or agreed to in writing, software
+ distributed under the License is distributed on an "AS IS" BASIS,
+ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ See the License for the specific language governing permissions and
+ limitations under the License. */
+
+
+#import "CPUCompute.h"
+
+#import <map>
+#import <vector>
+#import <utility>
+#import <algorithm>
+
+
+struct NMSParam {
+  
+  float *score_data;
+  
+  float *box_data;
+  
+  float *output;
+  
+  int output_size;
+  
+  std::vector<int> score_dim;
+  
+  std::vector<int> box_dim;
+  
+  float scoreThredshold;
+  
+  int nmsTopK;
+  
+  int keepTopK;
+  
+  float nmsEta;
+  
+  float nmsThreshold;
+  
+  int background_label;
+};
+
+
+constexpr int kOutputDim = 6;
+constexpr int kBBoxSize = 4;
+
+template <class T>
+bool SortScorePairDescend(const std::pair<float, T>& pair1,
+                          const std::pair<float, T>& pair2) {
+  return pair1.first > pair2.first;
+}
+
+template <class T>
+static inline void GetMaxScoreIndex(
+                                    const std::vector<T>& scores, const T threshold, int top_k,
+                                    std::vector<std::pair<T, int>>* sorted_indices) {
+  for (size_t i = 0; i < scores.size(); ++i) {
+    if (scores[i] > threshold) {
+      sorted_indices->push_back(std::make_pair(scores[i], i));
+    }
+  }
+  // Sort the score pair according to the scores in descending order
+  std::stable_sort(sorted_indices->begin(), sorted_indices->end(),
+                   SortScorePairDescend<int>);
+  // Keep top_k scores if needed.
+  if (top_k > -1 && top_k < static_cast<int>(sorted_indices->size())) {
+    sorted_indices->resize(top_k);
+  }
+}
+
+template <class T>
+static inline T BBoxArea(const T* box, const bool normalized) {
+  if (box[2] < box[0] || box[3] < box[1]) {
+    // If coordinate values are is invalid
+    // (e.g. xmax < xmin or ymax < ymin), return 0.
+    return static_cast<T>(0.);
+  } else {
+    const T w = box[2] - box[0];
+    const T h = box[3] - box[1];
+    if (normalized) {
+      return w * h;
+    } else {
+      // If coordinate values are not within range [0, 1].
+      return (w + 1) * (h + 1);
+    }
+  }
+}
+
+template <class T>
+static inline T JaccardOverlap(const T* box1, const T* box2,
+                               const bool normalized) {
+  if (box2[0] > box1[2] || box2[2] < box1[0] || box2[1] > box1[3] ||
+      box2[3] < box1[1]) {
+    return static_cast<T>(0.);
+  } else {
+    const T inter_xmin = std::max(box1[0], box2[0]);
+    const T inter_ymin = std::max(box1[1], box2[1]);
+    const T inter_xmax = std::min(box1[2], box2[2]);
+    const T inter_ymax = std::min(box1[3], box2[3]);
+    const T inter_w = inter_xmax - inter_xmin;
+    const T inter_h = inter_ymax - inter_ymin;
+    const T inter_area = inter_w * inter_h;
+    const T bbox1_area = BBoxArea<T>(box1, normalized);
+    const T bbox2_area = BBoxArea<T>(box2, normalized);
+    return inter_area / (bbox1_area + bbox2_area - inter_area);
+  }
+}
+
+template <typename T>
+static inline void NMSFast(
+                           const T *bbox_data,
+                           std::vector<int> bbox_dim,
+                           const T *score_data,
+                           const T score_threshold, const T nms_threshold,
+                           const T eta, const int top_k,
+                           std::vector<int>* selected_indices) {
+  // The total boxes for each instance.
+  int num_boxes = bbox_dim[0];
+  // 4: [xmin ymin xmax ymax]
+  int box_size = bbox_dim[1];
+
+  std::vector<T> scores_data(num_boxes);
+  std::copy_n(score_data, num_boxes, scores_data.begin());
+  std::vector<std::pair<T, int>> sorted_indices;
+  GetMaxScoreIndex(scores_data, score_threshold, top_k, &sorted_indices);
+
+  selected_indices->clear();
+  T adaptive_threshold = nms_threshold;
+
+  while (sorted_indices.size() != 0) {
+    const int idx = sorted_indices.front().second;
+    bool keep = true;
+    for (size_t k = 0; k < selected_indices->size(); ++k) {
+      if (keep) {
+        const int kept_idx = (*selected_indices)[k];
+        T overlap = JaccardOverlap<T>(bbox_data + idx * box_size,
+                                      bbox_data + kept_idx * box_size, true);
+        keep = overlap <= adaptive_threshold;
+      } else {
+        break;
+      }
+    }
+    if (keep) {
+      selected_indices->push_back(idx);
+    }
+    sorted_indices.erase(sorted_indices.begin());
+    if (keep && eta < 1 && adaptive_threshold > 0.5) {
+      adaptive_threshold *= eta;
+    }
+  }
+}
+
+template <typename T>
+void MultiClassNMS(const T *boxes_data,
+                   const std::vector<int> &box_dim,
+                   const T *scores_data,
+                   const std::vector<int> &score_dim,
+                   std::map<int, std::vector<int>>* indices, int* num_nmsed_out,
+                   const int& background_label, const int& nms_top_k,
+                   const int& keep_top_k, const T& nms_threshold,
+                   const T& nms_eta, const T& score_threshold) {
+  
+  int64_t class_num = score_dim[0];
+  int64_t predict_dim = score_dim[1];
+  int num_det = 0;
+  for (int c = 0; c < class_num; ++c) {
+    if (c == background_label) continue;
+    const T *score_data = scores_data + c * predict_dim;
+    
+    /// [c] is key
+    NMSFast<T>(boxes_data, box_dim, score_data, score_threshold, nms_threshold, nms_eta,
+                   nms_top_k, &((*indices)[c]));
+    num_det += (*indices)[c].size();
+  }
+
+  *num_nmsed_out = num_det;
+  if (keep_top_k > -1 && num_det > keep_top_k) {
+    std::vector<std::pair<T, std::pair<int, int>>> score_index_pairs;
+    for (const auto& it : *indices) {
+      int label = it.first;
+      const T* sdata = scores_data + label * predict_dim;
+      const std::vector<int>& label_indices = it.second;
+      for (size_t j = 0; j < label_indices.size(); ++j) {
+        int idx = label_indices[j];
+        // PADDLE_ENFORCE_LT(idx, predict_dim);
+        score_index_pairs.push_back(std::make_pair(sdata[idx], std::make_pair(label, idx)));
+      }
+    }
+    // Keep top k results per image.
+    std::stable_sort(score_index_pairs.begin(), score_index_pairs.end(),
+                     SortScorePairDescend<std::pair<int, int>>);
+    score_index_pairs.resize(keep_top_k);
+
+    // Store the new indices.
+    std::map<int, std::vector<int>> new_indices;
+    for (size_t j = 0; j < score_index_pairs.size(); ++j) {
+      int label = score_index_pairs[j].second.first;
+      int idx = score_index_pairs[j].second.second;
+      new_indices[label].push_back(idx);
+    }
+    new_indices.swap(*indices);
+    *num_nmsed_out = keep_top_k;
+  }
+}
+
+template <typename T>
+void MultiClassOutput(const T *scores_data,
+                      const std::vector<int> &score_dim,
+                      const T *bboxes_data,
+                      T *outputs_data,
+                      const std::map<int, std::vector<int>>& selected_indices) {
+  int predict_dim = score_dim[1];
+  int count = 0;
+  for (const auto& it : selected_indices) {
+    /// one batch
+    int label = it.first;
+    const T* sdata = scores_data + label * predict_dim;
+    const std::vector<int>& indices = it.second;
+    for (size_t j = 0; j < indices.size(); ++j) {
+      int idx = indices[j];
+      const T* bdata = bboxes_data + idx * kBBoxSize;
+      outputs_data[count * kOutputDim] = label;           // label
+      outputs_data[count * kOutputDim + 1] = sdata[idx];  // score
+      // xmin, ymin, xmax, ymax
+      std::memcpy(outputs_data + count * kOutputDim + 2, bdata, 4 * sizeof(T));
+      count++;
+    }
+  }
+}
+
+void MultiClassNMSCompute(NMSParam *param) {
+  assert(param->score_dim[0] == 1);
+  assert(param->box_dim[0] == 1);
+  assert (param->score_dim.size() == 3);
+  assert(param->box_dim.size() == 3);
+
+  float* outputs;
+  auto background_label = param->background_label;
+  auto nms_top_k = param->nmsTopK;
+  auto keep_top_k = param->keepTopK;
+  auto nms_threshold = param->nmsThreshold;
+  auto nms_eta = param->nmsEta;
+  auto score_threshold = param->scoreThredshold;
+
+  std::vector<int> score_dim_one_batch = {param->score_dim[1], param->score_dim[2]};
+  std::vector<int> box_dim_one_batch = {param->box_dim[1], param->box_dim[2]};
+  
+  std::vector<int> batch_starts = {0};
+  
+  std::map<int, std::vector<int>> indices;
+  int num_nmsed_out = 0;
+  
+  MultiClassNMS<float>(param->box_data, box_dim_one_batch, param->score_data, score_dim_one_batch, &indices, &num_nmsed_out,
+                       background_label, nms_top_k, keep_top_k, nms_threshold,
+                       nms_eta, score_threshold);
+  batch_starts.push_back(batch_starts.back() + num_nmsed_out);
+
+  int output_size = 0;
+  int num_kept = batch_starts.back();
+  if (num_kept == 0) {
+    outputs = new float[1];
+    outputs[0] = -1;
+    output_size = 1;
+  } else {
+    outputs = new float[num_kept * kOutputDim];
+    int64_t s = batch_starts[0];
+    int64_t e = batch_starts[1];
+    if (e > s) {
+      MultiClassOutput<float>(param->score_data, score_dim_one_batch, param->box_data, outputs, indices);
+    }
+    output_size = num_kept * kOutputDim;
+  }
+  param->output = outputs;
+  param->output_size = output_size;
+}
+
+@implementation NMSCompute
+
+-(NSArray<NSNumber *> *)computeWithScore:(float *)score andBBoxs:(float *)bbox {
+  NMSParam param;
+  param.box_data = bbox;
+  param.score_data = score;
+  param.background_label = self.background_label;
+  param.scoreThredshold = self.scoreThredshold;
+  param.nmsTopK = self.nmsTopK;
+  param.keepTopK = self.keepTopK;
+  param.nmsEta = self.nmsEta;
+  param.nmsThreshold = self.nmsThreshold;
+  std::vector<int> score_dim;
+  for (int i = 0; i < self.scoreDim.count; ++i) {
+    score_dim.push_back(self.scoreDim[i].intValue);
+  }
+  param.score_dim = score_dim;
+  
+  std::vector<int> box_dim;
+  for (int i = 0; i < self.bboxDim.count; ++i) {
+    box_dim.push_back(self.bboxDim[i].intValue);
+  }
+  param.box_dim = box_dim;
+  MultiClassNMSCompute(&param);
+  NSMutableArray<NSNumber *> *output = [NSMutableArray arrayWithCapacity:param.output_size];
+  for (int i = 0; i < param.output_size; ++i) {
+    [output addObject:[NSNumber numberWithFloat:param.output[i]]];
+  }
+  return output;
+}
+
+@end
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+

metal/paddle-mobile-demo/paddle-mobile-demo/MobileNetSSD.swift → metal/paddle-mobile-demo/paddle-mobile-demo/Net/MobileNetSSD.swift

@@ -14,11 +14,11 @@
 
 import Foundation
 import paddle_mobile
+//import 
 
 class MobileNet_ssd_hand: Net{
   
   var program: Program?
-  
   var executor: Executor<Float32>?
   
   let except: Int = 2
@@ -30,156 +30,47 @@ class MobileNet_ssd_hand: Net{
   }
   
   func resultStr(res: [Float]) -> String {
-    return "哈哈哈, 还没好"
-  }
-  
-  func bboxArea(box: [Float32], normalized: Bool) -> Float32 {
-    if box[2] < box[0] || box[3] < box[1] {
-      return 0.0
-    } else {
-      let w = box[2] - box[0]
-      let h = box[3] - box[1]
-      if normalized {
-        return w * h
-      } else {
-        return (w + 1) * (h + 1)
-      }
-    }
-  }
-  
-  func jaccardOverLap(box1: [Float32], box2: [Float32], normalized: Bool) -> Float32 {
-    if box2[0] > box1[2] || box2[2] < box1[0] || box2[1] > box1[3] ||
-      box2[3] < box1[1] {
-      return 0.0
-    } else {
-      let interXmin = max(box1[0], box2[0])
-      let interYmin = max(box1[1], box2[1])
-      let interXmax = min(box1[2], box2[2])
-      let interYmax = min(box1[3], box2[3])
-      let interW = interXmax - interXmin
-      let interH = interYmax - interYmin
-      let interArea = interW * interH
-      let bbox1Area = bboxArea(box: box1, normalized: normalized)
-      let bbox2Area = bboxArea(box: box2, normalized: normalized)
-      return interArea / (bbox1Area + bbox2Area - interArea)
-    }
+    return " \(res)"
   }
-  
+
   func fetchResult(paddleMobileRes: ResultHolder<Float32>) -> [Float32]{
+
     guard let interRes = paddleMobileRes.intermediateResults else {
       fatalError(" need have inter result ")
     }
-    
+
     guard let scores = interRes["Scores"], scores.count > 0, let score = scores[0] as?  Texture<Float32> else {
       fatalError(" need score ")
     }
-    
+
     guard let bboxs = interRes["BBoxes"], bboxs.count > 0, let bbox = bboxs[0] as? Texture<Float32> else {
       fatalError()
     }
-    
-    let score_thredshold: Float32 = 0.01
-    let nms_top_k = 400
-    let keep_top_k = 200
-    let nms_eta: Float32 = 1.0
-    var nms_threshold: Float32 = 0.45
-    
-    let bboxArr = bbox.metalTexture.floatArray { (f) -> Float32 in
+
+    var scoreFormatArr: [Float32] = score.metalTexture.realNHWC(dim: (n: score.originDim[0], h: score.originDim[1], w: score.originDim[2], c: score.originDim[3]))
+    var bboxArr = bbox.metalTexture.floatArray { (f) -> Float32 in
       return f
     }
+
+    let nmsCompute = NMSCompute.init()
+    nmsCompute.scoreThredshold = 0.01
+    nmsCompute.nmsTopK = 200
+    nmsCompute.keepTopK = 200
+    nmsCompute.nmsEta = 1.0
+    nmsCompute.nmsThreshold = 0.45
+    nmsCompute.background_label = 0;
     
-    let scoreFormatArr: [Float32] = score.metalTexture.realNHWC(dim: (n: score.originDim[0], h: score.originDim[1], w: score.originDim[2], c: score.originDim[3]))
-    var outputArr: [Float32] = []
-    let cNumOfOneClass = score.tensorDim[2]        // 1917
-    let boxSize = bbox.tensorDim[2]                 // 4
-    let classNum = score.tensorDim[1]              // 7
-    
-    var selectedIndexs: [Int : [(Int, Float32)]] = [:]
-    var numDet: Int = 0
-    for i in 0..<classNum {
-      var sliceScore = Array<Float32>(scoreFormatArr[(i * cNumOfOneClass)..<((i + 1) * cNumOfOneClass)])
-      
-      var scoreThresholdArr: [(Float32, Int)] = []
-      
-      for j in 0..<cNumOfOneClass {
-        if sliceScore[j] > score_thredshold {
-          scoreThresholdArr.append((sliceScore[j], j))
-        }
-      }
-      
-      scoreThresholdArr.sort { $0 > $1 }
-      
-      if scoreThresholdArr.count > nms_top_k {
-        scoreThresholdArr.removeLast(scoreThresholdArr.count - nms_top_k)
-      }
-      
-      var selectedIndex: [(Int, Float32)] = []
-      
-      while scoreThresholdArr.count > 0 {
-        let idx = scoreThresholdArr[0].1
-        let score = scoreThresholdArr[0].0
-        var keep = true
-        for j in 0..<selectedIndex.count {
-          if keep {
-            let keptIdx = selectedIndex[j].0
-            let box1 = Array<Float32>(bboxArr[(idx * boxSize)..<(idx * boxSize + 4)])
-            let box2 = Array<Float32>(bboxArr[(keptIdx * boxSize)..<(keptIdx * boxSize + 4)])
-            
-            let overlap = jaccardOverLap(box1: box1, box2: box2, normalized: true)
-            keep = (overlap <= nms_threshold)
-          } else {
-            break
-          }
-        }
-        
-        if keep {
-          selectedIndex.append((idx, score))
-        }
-        
-        scoreThresholdArr.removeFirst()
-        if keep && nms_eta < 1.0 && nms_threshold > 0.5 {
-          nms_threshold *= nms_eta
-        }
-      }
-      selectedIndexs[i] = selectedIndex
-      numDet += selectedIndex.count
-    }
-    
-    var scoreIndexPairs: [(Float32, (Int, Int))] = []
-    for selected in selectedIndexs {
-      for scoreIndex in selected.value {
-        scoreIndexPairs.append((scoreIndex.1, (selected.key, scoreIndex.0)))
-      }
-    }
-    
-    scoreIndexPairs.sort { $0.0 > $1.0 }
-    
-    if scoreIndexPairs.count > keep_top_k {
-      scoreIndexPairs.removeLast(scoreIndexPairs.count - keep_top_k)
+    nmsCompute.scoreDim = [NSNumber.init(value: score.tensorDim[0]), NSNumber.init(value: score.tensorDim[1]), NSNumber.init(value: score.tensorDim[2])]
+
+    nmsCompute.bboxDim = [NSNumber.init(value: bbox.tensorDim[0]), NSNumber.init(value: bbox.tensorDim[1]), NSNumber.init(value: bbox.tensorDim[2])]
+    guard let result = nmsCompute.compute(withScore: &scoreFormatArr, andBBoxs: &bboxArr) else {
+      fatalError( " result error " )
     }
+
+    let output: [Float32] = result.map { $0.floatValue }
     
-    var newIndices: [Int : [(Int, Float32)]] = [:]
-    for scoreIndexPair in scoreIndexPairs {
-      // label: scoreIndexPair.1.0
-      let label = scoreIndexPair.1.0
-      if newIndices[label] != nil {
-        newIndices[label]?.append((scoreIndexPair.1.1, scoreIndexPair.0))
-      } else {
-        newIndices[label] = [(scoreIndexPair.1.1, scoreIndexPair.0)]
-      }
-    }
     
-    for indice in newIndices {
-      let selectedIndexAndScore = indice.value
-      for indexAndScore in selectedIndexAndScore {
-        outputArr.append(Float32(indice.key))   // label
-        outputArr.append(indexAndScore.1)   // score
-        let subBox = bboxArr[(indexAndScore.0 * boxSize)..<(indexAndScore.0 * boxSize + 4)]
-        outputArr.append(contentsOf: subBox)
-      }
-    }
-    print(" fuck success !")
-    return outputArr
+    return output
   }
   
   var preprocessKernel: CusomKernel

metal/paddle-mobile-demo/paddle-mobile-demo/Net.swift → metal/paddle-mobile-demo/paddle-mobile-demo/Net/Net.swift

-//
-//  Net.swift
-//  paddle-mobile-demo
-//
-//  Created by liuRuiLong on 2018/8/27.
-//  Copyright © 2018年 orange. All rights reserved.
-//
+/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
+ 
+ Licensed under the Apache License, Version 2.0 (the "License");
+ you may not use this file except in compliance with the License.
+ You may obtain a copy of the License at
+ 
+ http://www.apache.org/licenses/LICENSE-2.0
+ 
+ Unless required by applicable law or agreed to in writing, software
+ distributed under the License is distributed on an "AS IS" BASIS,
+ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ See the License for the specific language governing permissions and
+ limitations under the License. */
 
 import Foundation
 
@@ -14,13 +20,16 @@ import Foundation
 import paddle_mobile
 import MetalPerformanceShaders
 
-let modelHelperMap: [SupportModel : Net] = [.mobilenet : MobileNet.init(), .mobilenet_ssd : MobileNet_ssd_hand.init()]
+
+let modelHelperMap: [SupportModel : Net] = [.mobilenet_ssd : MobileNet_ssd_hand.init()]
+//let modelHelperMap: [SupportModel : Net] = [.mobilenet : MobileNet.init(), .mobilenet_ssd : MobileNet_ssd_hand.init()]
 
 enum SupportModel: String{
-  case mobilenet = "mobilenet"
+//  case mobilenet = "mobilenet"
   case mobilenet_ssd = "mobilenetssd"
   static func supportedModels() -> [SupportModel] {
-    return [.mobilenet, .mobilenet_ssd]
+    //.mobilenet,
+    return [.mobilenet_ssd]
   }
 }
 
@@ -37,7 +46,8 @@ protocol Net {
   func resultStr(res: [Float]) -> String
   func fetchResult(paddleMobileRes: ResultHolder<Float32>) -> [Float32]
   mutating func load() throws
-  func predict(inTexture: MTLTexture, completion: @escaping ([Float32]) -> Void) throws
+  
+  func predict(inTexture: MTLTexture, completion: @escaping ((time:TimeInterval, resultArray: [Float32])) -> Void) throws
   mutating func clear()
 }
 
@@ -54,13 +64,16 @@ extension Net {
     }
   }
   
-  func predict(inTexture: MTLTexture, completion: @escaping ([Float32]) -> Void) throws {
+  func predict(inTexture: MTLTexture, completion: @escaping ((time:TimeInterval, resultArray: [Float32])) -> Void) throws {
     guard let inExecutor = executor else {
       fatalError(" 请先 load ")
     }
     try inExecutor.predict(input: inTexture, dim: dim, completionHandle: { (result) in
-      let resultArr = self.fetchResult(paddleMobileRes: result)
-      completion(resultArr)
+      
+      var resultArr:[Float32] = []
+      resultArr = self.fetchResult(paddleMobileRes: result)
+      completion((time: TimeInterval(result.elapsedTime), resultArray: resultArr))
+
     }, preProcessKernle: preprocessKernel, except: except)
   }
   
@@ -81,6 +94,4 @@ extension Net {
     return paddleMobileRes.resultArr
   }
   
-  //  func predict()
-  
 }

metal/paddle-mobile-demo/paddle-mobile-demo/ViewController.swift

@@ -28,6 +28,7 @@ class ViewController: UIViewController {
   var selectImage: UIImage?
   var modelType: SupportModel = SupportModel.supportedModels()[0]
   var toPredictTexture: MTLTexture?
+  
   var net: Net {
     get {
       return modelHelperMap[modelType] ?! " has no this type "
@@ -35,6 +36,7 @@ class ViewController: UIViewController {
     set {
     }
   }
+  
   var threadNum = 1
   
   @IBAction func loadAct(_ sender: Any) {
@@ -63,53 +65,26 @@ class ViewController: UIViewController {
       return
     }
     do {
-      try  net.predict(inTexture: inTexture) { [weak self] (result) in
-        guard let sSelf = self else {
-          fatalError()
-        }
-        print(result)
-        let resultStr = sSelf.net.resultStr(res: result)
-        DispatchQueue.main.async {
-            sSelf.resultTextView.text = resultStr
+      let max = 1
+      let startDate = Date.init()
+      for i in 0..<max {
+        try net.predict(inTexture: inTexture) { [weak self] (result) in
+          guard let sSelf = self else {
+            fatalError()
+          }
+          
+          if i == max - 1 {
+            let time = Date.init().timeIntervalSince(startDate)
+            DispatchQueue.main.async {
+              sSelf.resultTextView.text = sSelf.net.resultStr(res: result.resultArray)
+              sSelf.elapsedTimeLabel.text = "平均耗时: \(time/Double(max) * 1000.0) ms"
+            }
+          }
         }
       }
     } catch let error {
       print(error)
     }
-   
-    
-//    guard let inExecutor = executor else {
-//      resultTextView.text = "请先 load ! "
-//      return
-//    }
-//
-//    do {
-//      let max = 1
-//      var startDate = Date.init()
-//      for i in 0..<max {
-//        try inExecutor.predict(input: inTexture, dim: modelHelper.dim, completionHandle: { [weak self] (result) in
-//          guard let sSelf = self else {
-//            fatalError()
-//          }
-//
-//          if i == (max / 2 - 1) {
-//            startDate = Date.init()
-//          }
-//
-//          let resultArr = sSelf.modelHelper.fetchResult(paddleMobileRes: result)
-//
-//          if i == max - 1 {
-//            let time = Date.init().timeIntervalSince(startDate)
-//            DispatchQueue.main.async {
-//              sSelf.resultTextView.text = sSelf.modelHelper.resultStr(res: resultArr)
-//              sSelf.elapsedTimeLabel.text = "平均耗时: \(time/Double(max/2) * 1000.0) ms"
-//            }
-//          }
-//          }, preProcessKernle: self.modelHelper.preprocessKernel, except: modelHelper.except)
-//      }
-//    } catch let error {
-//      print(error)
-//    }
   }
   
   override func viewDidLoad() {

metal/paddle-mobile-demo/paddle-mobile-demo/paddle-mobile-demo-Bridging-Header.h

+//
+//  Use this file to import your target's public headers that you would like to expose to Swift.
+//
+
+
+
+#import "CPUCompute.h"

metal/paddle-mobile/paddle-mobile/Common/Extensions.swift

@@ -98,7 +98,7 @@ extension Array where Element: Comparable{
 }
 
 extension Array {
-  func strideArray(inCount: Int = 20) -> [(Int, Element)] {
+  public func strideArray(inCount: Int = 20) -> [(Int, Element)] {
     if count < inCount {
       return (0..<count).map{ ($0, self[$0]) }
     } else {

metal/paddle-mobile/paddle-mobile/Common/MetalExtension.swift

@@ -217,16 +217,16 @@ extension MTLComputeCommandEncoder {
     let height = computePipline.maxTotalThreadsPerThreadgroup/width
     let threadsPerGroup = MTLSize.init(width: width, height: height, depth: 1)
     
-    print(" thread: threads per group: \(threadsPerGroup) ")
-    print(" thread: out texture width: \(outTexture.width) , out texture height: \(outTexture.height)")
+//    print(" thread: threads per group: \(threadsPerGroup) ")
+//    print(" thread: out texture width: \(outTexture.width) , out texture height: \(outTexture.height)")
     
     let groupWidth = (outTexture.width + width - 1)/width
     let groupHeight = (outTexture.height + height - 1)/height
     let groupDepth = slices
     let groups = MTLSize.init(width: groupWidth, height: groupHeight, depth: groupDepth)
     
-    print("groups: \(groups) ")
-    print("threads per group: \(threadsPerGroup)")
+//    print("groups: \(groups) ")
+//    print("threads per group: \(threadsPerGroup)")
     
     setComputePipelineState(computePipline)
     
@@ -367,9 +367,9 @@ public extension MTLTexture {
   }
   
   func realNHWC(dim: (n: Int, h: Int, w: Int, c: Int)) -> [Float32] {
-    print("origin dim: \(dim)")
-    print("texture: ")
-    print(self)
+//    print("origin dim: \(dim)")
+//    print("texture: ")
+//    print(self)
     
     let textureArray = floatArray { (i : Float32) -> Float32 in
       return i
@@ -394,7 +394,7 @@ public extension MTLTexture {
         }
       }
     }
-    print(" tensor count -- \(output.count)")
+//    print(" tensor count -- \(output.count)")
     return output
   }
   

metal/paddle-mobile/paddle-mobile/Executor.swift

@@ -125,9 +125,9 @@ public class Executor<P: PrecisionType> {
       //            print(stridableInput)
       
       //            let _: Flo? = input.logDesc(header: "input: ", stridable: true)
-      for op in self.ops {
+//      for op in self.ops {
 //          op.delogOutput()
-      }
+//      }
 //                  return
       
 //      self.ops[testTo].delogOutput()

metal/paddle-mobile/paddle-mobile/Operators/Kernels/PriorBoxKernel.swift

@@ -98,16 +98,6 @@ class PriorBoxKernel<P: PrecisionType>: Kernel, Computable{
     guard let encoder = commandBuffer.makeComputeCommandEncoder() else {
       throw PaddleMobileError.predictError(message: " encode is nil")
     }
-    print("metalParam: \(metalParam)")
-    
-    print(" newAspectRatios ")
-    print(param.newAspectRatios!)
-    
-    print(" clip: \(metalParam.clip)")
-    
-    print(" metalParam.numPriors: \(metalParam.numPriors)")
-    
-    print(" aspecRatiosSize: \(metalParam.aspecRatiosSize)")
     
     encoder.setTexture(param.input.metalTexture, index: 0)
     encoder.setTexture(param.output.metalTexture, index: 1)

metal/paddle-mobile/paddle-mobile/framework/Tensor.swift

@@ -174,7 +174,7 @@ class Tensor<P: PrecisionType>: Tensorial {
             fatalError(" not support !")
         }
         //TODO: release
-//        data.release()
+        data.release()
     }
     
     var width: Int {