提交 1b77c29f 编写于 作者: D dolphin8

ar

...@@ -20,6 +20,8 @@ ...@@ -20,6 +20,8 @@
FC8CFEF9213551D10094D569 /* model in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEF7213551D00094D569 /* model */; }; FC8CFEF9213551D10094D569 /* model in Resources */ = {isa = PBXBuildFile; fileRef = FC8CFEF7213551D00094D569 /* model */; };
FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */ = {isa = PBXBuildFile; fileRef = FC918190211DBC3500B6F354 /* paddle-mobile.png */; }; FC918191211DBC3500B6F354 /* paddle-mobile.png in Resources */ = {isa = PBXBuildFile; fileRef = FC918190211DBC3500B6F354 /* paddle-mobile.png */; };
FC918193211DC70500B6F354 /* iphone.JPG in Resources */ = {isa = PBXBuildFile; fileRef = FC918192211DC70500B6F354 /* iphone.JPG */; }; FC918193211DC70500B6F354 /* iphone.JPG in Resources */ = {isa = PBXBuildFile; fileRef = FC918192211DC70500B6F354 /* iphone.JPG */; };
FC9A19E72148C38400CD9CBF /* ar_model in Resources */ = {isa = PBXBuildFile; fileRef = FC9A19E52148C38400CD9CBF /* ar_model */; };
FC9A19E82148C38400CD9CBF /* ar_params in Resources */ = {isa = PBXBuildFile; fileRef = FC9A19E62148C38400CD9CBF /* ar_params */; };
FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */ = {isa = PBXBuildFile; fileRef = FCA3A16021313E1F00084FE5 /* hand.jpg */; }; FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */ = {isa = PBXBuildFile; fileRef = FCA3A16021313E1F00084FE5 /* hand.jpg */; };
FCBCCC522122EEDC00D94F7E /* ssd_hand_params in Resources */ = {isa = PBXBuildFile; fileRef = FCBCCC502122EEDC00D94F7E /* ssd_hand_params */; }; FCBCCC522122EEDC00D94F7E /* ssd_hand_params in Resources */ = {isa = PBXBuildFile; fileRef = FCBCCC502122EEDC00D94F7E /* ssd_hand_params */; };
FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */ = {isa = PBXBuildFile; fileRef = FCBCCC512122EEDC00D94F7E /* ssd_hand_model */; }; FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */ = {isa = PBXBuildFile; fileRef = FCBCCC512122EEDC00D94F7E /* ssd_hand_model */; };
...@@ -64,6 +66,8 @@ ...@@ -64,6 +66,8 @@
FC8CFEF7213551D00094D569 /* model */ = {isa = PBXFileReference; lastKnownFileType = file; path = model; sourceTree = "<group>"; }; FC8CFEF7213551D00094D569 /* model */ = {isa = PBXFileReference; lastKnownFileType = file; path = model; sourceTree = "<group>"; };
FC918190211DBC3500B6F354 /* paddle-mobile.png */ = {isa = PBXFileReference; lastKnownFileType = image.png; path = "paddle-mobile.png"; 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>"; }; FC918192211DC70500B6F354 /* iphone.JPG */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = iphone.JPG; sourceTree = "<group>"; };
FC9A19E52148C38400CD9CBF /* ar_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = ar_model; sourceTree = "<group>"; };
FC9A19E62148C38400CD9CBF /* ar_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = ar_params; sourceTree = "<group>"; };
FCA3A16021313E1F00084FE5 /* hand.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = hand.jpg; sourceTree = "<group>"; }; FCA3A16021313E1F00084FE5 /* hand.jpg */ = {isa = PBXFileReference; lastKnownFileType = image.jpeg; path = hand.jpg; sourceTree = "<group>"; };
FCBCCC502122EEDC00D94F7E /* ssd_hand_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_params; sourceTree = "<group>"; }; FCBCCC502122EEDC00D94F7E /* ssd_hand_params */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_params; sourceTree = "<group>"; };
FCBCCC512122EEDC00D94F7E /* ssd_hand_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_model; sourceTree = "<group>"; }; FCBCCC512122EEDC00D94F7E /* ssd_hand_model */ = {isa = PBXFileReference; lastKnownFileType = file; path = ssd_hand_model; sourceTree = "<group>"; };
...@@ -156,6 +160,7 @@ ...@@ -156,6 +160,7 @@
FC0E2C2020EDC03B009C1FAC /* models */ = { FC0E2C2020EDC03B009C1FAC /* models */ = {
isa = PBXGroup; isa = PBXGroup;
children = ( children = (
FC9A19E42148C38400CD9CBF /* fluid_fssd_new_ar */,
FC8CFEF5213551D00094D569 /* mobilenet */, FC8CFEF5213551D00094D569 /* mobilenet */,
FC8CFEE32135452B0094D569 /* genet */, FC8CFEE32135452B0094D569 /* genet */,
FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */, FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */,
...@@ -191,6 +196,15 @@ ...@@ -191,6 +196,15 @@
path = mobilenet; path = mobilenet;
sourceTree = "<group>"; sourceTree = "<group>";
}; };
FC9A19E42148C38400CD9CBF /* fluid_fssd_new_ar */ = {
isa = PBXGroup;
children = (
FC9A19E52148C38400CD9CBF /* ar_model */,
FC9A19E62148C38400CD9CBF /* ar_params */,
);
path = fluid_fssd_new_ar;
sourceTree = "<group>";
};
FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */ = { FCBCCC4F2122EEDC00D94F7E /* mobilenet_ssd_hand */ = {
isa = PBXGroup; isa = PBXGroup;
children = ( children = (
...@@ -276,6 +290,8 @@ ...@@ -276,6 +290,8 @@
FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */, FCA3A16121313E1F00084FE5 /* hand.jpg in Resources */,
FC8CFEE62135452C0094D569 /* genet_params in Resources */, FC8CFEE62135452C0094D569 /* genet_params in Resources */,
FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */, FCBCCC532122EEDC00D94F7E /* ssd_hand_model in Resources */,
FC9A19E72148C38400CD9CBF /* ar_model in Resources */,
FC9A19E82148C38400CD9CBF /* ar_params in Resources */,
); );
runOnlyForDeploymentPostprocessing = 0; runOnlyForDeploymentPostprocessing = 0;
}; };
......
...@@ -44,7 +44,6 @@ kernel void mobilenet_preprocess_half( ...@@ -44,7 +44,6 @@ kernel void mobilenet_preprocess_half(
outTexture.write(half4(inColor.z, inColor.y, inColor.x, 0.0f), gid); outTexture.write(half4(inColor.z, inColor.y, inColor.x, 0.0f), gid);
} }
kernel void mobilenet_ssd_preprocess( kernel void mobilenet_ssd_preprocess(
texture2d<float, access::read> inTexture [[texture(0)]], texture2d<float, access::read> inTexture [[texture(0)]],
texture2d<float, access::write> outTexture [[texture(1)]], texture2d<float, access::write> outTexture [[texture(1)]],
...@@ -73,7 +72,6 @@ kernel void mobilenet_ssd_preprocess_half( ...@@ -73,7 +72,6 @@ kernel void mobilenet_ssd_preprocess_half(
outTexture.write(half4(inColor.z, inColor.y, inColor.x, 0.0f), gid); outTexture.write(half4(inColor.z, inColor.y, inColor.x, 0.0f), gid);
} }
kernel void genet_preprocess(texture2d<float, access::read> inTexture [[texture(0)]], texture2d<float, access::write> outTexture [[texture(1)]], uint2 gid [[thread_position_in_grid]]) kernel void genet_preprocess(texture2d<float, access::read> inTexture [[texture(0)]], texture2d<float, access::write> outTexture [[texture(1)]], uint2 gid [[thread_position_in_grid]])
{ {
if (gid.x >= outTexture.get_width() || if (gid.x >= outTexture.get_width() ||
...@@ -96,6 +94,28 @@ kernel void genet_preprocess_half(texture2d<half, access::read> inTexture [[text ...@@ -96,6 +94,28 @@ kernel void genet_preprocess_half(texture2d<half, access::read> inTexture [[text
outTexture.write(half4(inColor.z, inColor.y, inColor.x, 0.0f), gid); outTexture.write(half4(inColor.z, inColor.y, inColor.x, 0.0f), gid);
} }
kernel void mobilent_ar_preprocess(texture2d<float, access::read> inTexture [[texture(0)]], texture2d<float, access::write> outTexture [[texture(1)]], uint2 gid [[thread_position_in_grid]])
{
if (gid.x >= outTexture.get_width() ||
gid.y >= outTexture.get_height()) {
return;
}
const auto means = float4(128.0f, 128.0f, 128.0f, 0.0f);
const float4 inColor = (inTexture.read(gid) * 255.0 - means) * 0.017;
outTexture.write(float4(inColor.z, inColor.y, inColor.x, 0.0f), gid);
}
kernel void mobilent_ar_preprocess_half(texture2d<half, access::read> inTexture [[texture(0)]], texture2d<half, access::write> outTexture [[texture(1)]], uint2 gid [[thread_position_in_grid]])
{
if (gid.x >= outTexture.get_width() ||
gid.y >= outTexture.get_height()) {
return;
}
const auto means = half4(128.0f, 128.0f, 128.0f, 0.0f);
const half4 inColor = (inTexture.read(gid) * 255.0 - means) * 0.017;
outTexture.write(half4(inColor.z, inColor.y, inColor.x, 0.0f), gid);
}
kernel void scale(texture2d<float, access::sample> inTexture [[texture(0)]], texture2d<float, access::write> outTexture [[texture(1)]], uint2 gid [[thread_position_in_grid]]) { kernel void scale(texture2d<float, access::sample> inTexture [[texture(0)]], texture2d<float, access::write> outTexture [[texture(1)]], uint2 gid [[thread_position_in_grid]]) {
if (gid.x >= outTexture.get_width() || if (gid.x >= outTexture.get_width() ||
gid.y >= outTexture.get_height()) return; gid.y >= outTexture.get_height()) return;
...@@ -115,4 +135,3 @@ kernel void scale_half(texture2d<float, access::sample> inTexture [[texture(0)]] ...@@ -115,4 +135,3 @@ kernel void scale_half(texture2d<float, access::sample> inTexture [[texture(0)]]
float4 input = inTexture.sample(sample, float2(gid.x * w_stride, gid.y * h_stride), 0); float4 input = inTexture.sample(sample, float2(gid.x * w_stride, gid.y * h_stride), 0);
outTexture.write(half4(input), gid); outTexture.write(half4(input), gid);
} }
...@@ -21,17 +21,20 @@ let platform: Platform = .GPU ...@@ -21,17 +21,20 @@ let platform: Platform = .GPU
let threadSupport = [1] let threadSupport = [1]
let modelHelperMap: [SupportModel : Runner] = [.mobilenet_ssd : Runner.init(inNet: MobileNet_ssd_hand.init(device: MetalHelper.shared.device), commandQueue: MetalHelper.shared.queue, inPlatform: platform), let modelHelperMap: [SupportModel : Runner] = [.mobilenet_ssd : Runner.init(inNet: MobileNet_ssd_hand.init(device: MetalHelper.shared.device), commandQueue: MetalHelper.shared.queue, inPlatform: platform),
.genet : Runner.init(inNet: Genet.init(device: MetalHelper.shared.device), commandQueue: MetalHelper.shared.queue, inPlatform: platform)] .genet : Runner.init(inNet: Genet.init(device: MetalHelper.shared.device), commandQueue: MetalHelper.shared.queue, inPlatform: platform),
.mobilenet_ssd_ar : Runner.init(inNet: MobileNet_ssd_AR.init(device: MetalHelper.shared.device), commandQueue: MetalHelper.shared.queue, inPlatform: platform)]
//, .genet : Genet.init() //, .genet : Genet.init()
//let modelHelperMap: [SupportModel : Net] = [.mobilenet : MobileNet.init(), .mobilenet_ssd : MobileNet_ssd_hand.init()] //let modelHelperMap: [SupportModel : Net] = [.mobilenet : MobileNet.init(), .mobilenet_ssd : MobileNet_ssd_hand.init()]
enum SupportModel: String{ enum SupportModel: String{
// case mobilenet = "mobilenet" // case mobilenet = "mobilenet"
case mobilenet_ssd = "mobilenetssd" case mobilenet_ssd = "mobilenetssd"
case genet = "genet" case genet = "genet"
case mobilenet_ssd_ar = "mobilenetssd_ar"
static func supportedModels() -> [SupportModel] { static func supportedModels() -> [SupportModel] {
//.mobilenet, //.mobilenet,
return [.mobilenet_ssd, .genet] return [.mobilenet_ssd, .genet, .mobilenet_ssd_ar]
} }
} }
...@@ -78,7 +81,7 @@ class ViewController: UIViewController { ...@@ -78,7 +81,7 @@ class ViewController: UIViewController {
} }
@IBAction func predictAct(_ sender: Any) { @IBAction func predictAct(_ sender: Any) {
let max = 50 let max = 1
switch platform { switch platform {
case .GPU: case .GPU:
guard let inTexture = toPredictTexture else { guard let inTexture = toPredictTexture else {
...@@ -102,9 +105,10 @@ class ViewController: UIViewController { ...@@ -102,9 +105,10 @@ class ViewController: UIViewController {
} }
} }
} }
// print("sleep before ")
// usleep(33000)
// print("sleep after ")
} }
case .CPU: case .CPU:
guard let inInputPointer = inputPointer else { guard let inInputPointer = inputPointer else {
fatalError( " need input pointer " ) fatalError( " need input pointer " )
......
...@@ -15,7 +15,6 @@ ...@@ -15,7 +15,6 @@
4AA1EA90214664CD00D0F791 /* Split.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA8F214664CD00D0F791 /* Split.metal */; }; 4AA1EA90214664CD00D0F791 /* Split.metal in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA8F214664CD00D0F791 /* Split.metal */; };
4AA1EA92214665D700D0F791 /* ShapeOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA91214665D700D0F791 /* ShapeOp.swift */; }; 4AA1EA92214665D700D0F791 /* ShapeOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA91214665D700D0F791 /* ShapeOp.swift */; };
4AA1EA942146661500D0F791 /* ShapeKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA932146661500D0F791 /* ShapeKernel.swift */; }; 4AA1EA942146661500D0F791 /* ShapeKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA932146661500D0F791 /* ShapeKernel.swift */; };
4AA1EA962146665A00D0F791 /* FlattenKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA952146665A00D0F791 /* FlattenKernel.swift */; };
4AA1EA982146666500D0F791 /* FlattenOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA972146666500D0F791 /* FlattenOp.swift */; }; 4AA1EA982146666500D0F791 /* FlattenOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA972146666500D0F791 /* FlattenOp.swift */; };
4AA1EA9E2148D6F900D0F791 /* ConcatKernel.metal.inc in Headers */ = {isa = PBXBuildFile; fileRef = 4AA1EA9D2148D6F900D0F791 /* ConcatKernel.metal.inc */; }; 4AA1EA9E2148D6F900D0F791 /* ConcatKernel.metal.inc in Headers */ = {isa = PBXBuildFile; fileRef = 4AA1EA9D2148D6F900D0F791 /* ConcatKernel.metal.inc */; };
4AA1EAA02148DEEE00D0F791 /* ReshapeKernel.metal.inc in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA9F2148DEEE00D0F791 /* ReshapeKernel.metal.inc */; }; 4AA1EAA02148DEEE00D0F791 /* ReshapeKernel.metal.inc in Sources */ = {isa = PBXBuildFile; fileRef = 4AA1EA9F2148DEEE00D0F791 /* ReshapeKernel.metal.inc */; };
...@@ -59,6 +58,7 @@ ...@@ -59,6 +58,7 @@
FC292C82214255BD00CF622F /* MobileNetSSD.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC292C7E214255BC00CF622F /* MobileNetSSD.swift */; }; FC292C82214255BD00CF622F /* MobileNetSSD.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC292C7E214255BC00CF622F /* MobileNetSSD.swift */; };
FC292C85214257CB00CF622F /* CPUCompute.h in Headers */ = {isa = PBXBuildFile; fileRef = FC292C7D214255BC00CF622F /* CPUCompute.h */; settings = {ATTRIBUTES = (Public, ); }; }; FC292C85214257CB00CF622F /* CPUCompute.h in Headers */ = {isa = PBXBuildFile; fileRef = FC292C7D214255BC00CF622F /* CPUCompute.h */; settings = {ATTRIBUTES = (Public, ); }; };
FC292C872142624800CF622F /* Genet.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC292C862142624800CF622F /* Genet.swift */; }; FC292C872142624800CF622F /* Genet.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC292C862142624800CF622F /* Genet.swift */; };
FC33B0F02147659000714A93 /* MobileNet.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC33B0EF2147659000714A93 /* MobileNet.swift */; };
FC3602CC2108819F00FACB58 /* PaddleMobileUnitTest.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */; }; FC3602CC2108819F00FACB58 /* PaddleMobileUnitTest.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */; };
FC4CB74920F0B954007C0C6D /* ConvKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74820F0B954007C0C6D /* ConvKernel.metal */; }; FC4CB74920F0B954007C0C6D /* ConvKernel.metal in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74820F0B954007C0C6D /* ConvKernel.metal */; };
FC4CB74B20F12C30007C0C6D /* ProgramOptimize.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */; }; FC4CB74B20F12C30007C0C6D /* ProgramOptimize.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */; };
...@@ -69,6 +69,7 @@ ...@@ -69,6 +69,7 @@
FC5163F620EF556E00636C28 /* Texture2DTo2DArrayKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */; }; FC5163F620EF556E00636C28 /* Texture2DTo2DArrayKernel.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */; };
FC60DB8920E9AAA500FF203F /* MetalExtension.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC60DB8820E9AAA500FF203F /* MetalExtension.swift */; }; FC60DB8920E9AAA500FF203F /* MetalExtension.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC60DB8820E9AAA500FF203F /* MetalExtension.swift */; };
FC82735920E3C04200BE430A /* OpCreator.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC82735820E3C04200BE430A /* OpCreator.swift */; }; FC82735920E3C04200BE430A /* OpCreator.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC82735820E3C04200BE430A /* OpCreator.swift */; };
FC9A19E32148C31300CD9CBF /* MobilenetSSD_AR.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9A19E22148C31300CD9CBF /* MobilenetSSD_AR.swift */; };
FC9D037920E229E4000F735A /* OpParam.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D037820E229E4000F735A /* OpParam.swift */; }; FC9D037920E229E4000F735A /* OpParam.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D037820E229E4000F735A /* OpParam.swift */; };
FC9D038020E22FBB000F735A /* FeedOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D037F20E22FBB000F735A /* FeedOp.swift */; }; FC9D038020E22FBB000F735A /* FeedOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D037F20E22FBB000F735A /* FeedOp.swift */; };
FC9D038220E2312E000F735A /* FetchOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D038120E2312E000F735A /* FetchOp.swift */; }; FC9D038220E2312E000F735A /* FetchOp.swift in Sources */ = {isa = PBXBuildFile; fileRef = FC9D038120E2312E000F735A /* FetchOp.swift */; };
...@@ -124,7 +125,6 @@ ...@@ -124,7 +125,6 @@
4AA1EA8F214664CD00D0F791 /* Split.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Split.metal; sourceTree = "<group>"; }; 4AA1EA8F214664CD00D0F791 /* Split.metal */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.metal; path = Split.metal; sourceTree = "<group>"; };
4AA1EA91214665D700D0F791 /* ShapeOp.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = ShapeOp.swift; sourceTree = "<group>"; }; 4AA1EA91214665D700D0F791 /* ShapeOp.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = ShapeOp.swift; sourceTree = "<group>"; };
4AA1EA932146661500D0F791 /* ShapeKernel.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = ShapeKernel.swift; sourceTree = "<group>"; }; 4AA1EA932146661500D0F791 /* ShapeKernel.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = ShapeKernel.swift; sourceTree = "<group>"; };
4AA1EA952146665A00D0F791 /* FlattenKernel.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = FlattenKernel.swift; sourceTree = "<group>"; };
4AA1EA972146666500D0F791 /* FlattenOp.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = FlattenOp.swift; sourceTree = "<group>"; }; 4AA1EA972146666500D0F791 /* FlattenOp.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = FlattenOp.swift; sourceTree = "<group>"; };
4AA1EA9D2148D6F900D0F791 /* ConcatKernel.metal.inc */ = {isa = PBXFileReference; explicitFileType = sourcecode.metal; fileEncoding = 4; path = ConcatKernel.metal.inc; sourceTree = "<group>"; }; 4AA1EA9D2148D6F900D0F791 /* ConcatKernel.metal.inc */ = {isa = PBXFileReference; explicitFileType = sourcecode.metal; fileEncoding = 4; path = ConcatKernel.metal.inc; sourceTree = "<group>"; };
4AA1EA9F2148DEEE00D0F791 /* ReshapeKernel.metal.inc */ = {isa = PBXFileReference; explicitFileType = sourcecode.metal; fileEncoding = 4; path = ReshapeKernel.metal.inc; sourceTree = "<group>"; }; 4AA1EA9F2148DEEE00D0F791 /* ReshapeKernel.metal.inc */ = {isa = PBXFileReference; explicitFileType = sourcecode.metal; fileEncoding = 4; path = ReshapeKernel.metal.inc; sourceTree = "<group>"; };
...@@ -173,6 +173,7 @@ ...@@ -173,6 +173,7 @@
FC292C7D214255BC00CF622F /* CPUCompute.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = CPUCompute.h; sourceTree = "<group>"; }; FC292C7D214255BC00CF622F /* CPUCompute.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = CPUCompute.h; sourceTree = "<group>"; };
FC292C7E214255BC00CF622F /* MobileNetSSD.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobileNetSSD.swift; sourceTree = "<group>"; }; FC292C7E214255BC00CF622F /* MobileNetSSD.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobileNetSSD.swift; sourceTree = "<group>"; };
FC292C862142624800CF622F /* Genet.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = Genet.swift; sourceTree = "<group>"; }; FC292C862142624800CF622F /* Genet.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = Genet.swift; sourceTree = "<group>"; };
FC33B0EF2147659000714A93 /* MobileNet.swift */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.swift; path = MobileNet.swift; sourceTree = "<group>"; };
FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PaddleMobileUnitTest.swift; sourceTree = "<group>"; }; FC3602CB2108819F00FACB58 /* PaddleMobileUnitTest.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = PaddleMobileUnitTest.swift; sourceTree = "<group>"; };
FC4CB74820F0B954007C0C6D /* ConvKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvKernel.metal; sourceTree = "<group>"; }; FC4CB74820F0B954007C0C6D /* ConvKernel.metal */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.metal; path = ConvKernel.metal; sourceTree = "<group>"; };
FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ProgramOptimize.swift; sourceTree = "<group>"; }; FC4CB74A20F12C30007C0C6D /* ProgramOptimize.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = ProgramOptimize.swift; sourceTree = "<group>"; };
...@@ -183,6 +184,7 @@ ...@@ -183,6 +184,7 @@
FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Texture2DTo2DArrayKernel.swift; sourceTree = "<group>"; }; FC5163F520EF556E00636C28 /* Texture2DTo2DArrayKernel.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = Texture2DTo2DArrayKernel.swift; sourceTree = "<group>"; };
FC60DB8820E9AAA500FF203F /* MetalExtension.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MetalExtension.swift; sourceTree = "<group>"; }; FC60DB8820E9AAA500FF203F /* MetalExtension.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MetalExtension.swift; sourceTree = "<group>"; };
FC82735820E3C04200BE430A /* OpCreator.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = OpCreator.swift; sourceTree = "<group>"; }; FC82735820E3C04200BE430A /* OpCreator.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = OpCreator.swift; sourceTree = "<group>"; };
FC9A19E22148C31300CD9CBF /* MobilenetSSD_AR.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = MobilenetSSD_AR.swift; sourceTree = "<group>"; };
FC9D037820E229E4000F735A /* OpParam.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = OpParam.swift; sourceTree = "<group>"; }; FC9D037820E229E4000F735A /* OpParam.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = OpParam.swift; sourceTree = "<group>"; };
FC9D037F20E22FBB000F735A /* FeedOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = FeedOp.swift; sourceTree = "<group>"; }; FC9D037F20E22FBB000F735A /* FeedOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = FeedOp.swift; sourceTree = "<group>"; };
FC9D038120E2312E000F735A /* FetchOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = FetchOp.swift; sourceTree = "<group>"; }; FC9D038120E2312E000F735A /* FetchOp.swift */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.swift; path = FetchOp.swift; sourceTree = "<group>"; };
...@@ -282,6 +284,8 @@ ...@@ -282,6 +284,8 @@
FC039B6C20E11C3C0081E9F8 /* paddle-mobile */ = { FC039B6C20E11C3C0081E9F8 /* paddle-mobile */ = {
isa = PBXGroup; isa = PBXGroup;
children = ( children = (
FC9A19E22148C31300CD9CBF /* MobilenetSSD_AR.swift */,
FC33B0EF2147659000714A93 /* MobileNet.swift */,
FC292C862142624800CF622F /* Genet.swift */, FC292C862142624800CF622F /* Genet.swift */,
FC292C7E214255BC00CF622F /* MobileNetSSD.swift */, FC292C7E214255BC00CF622F /* MobileNetSSD.swift */,
FC292C7C214255BC00CF622F /* CPUCompute.mm */, FC292C7C214255BC00CF622F /* CPUCompute.mm */,
...@@ -399,7 +403,6 @@ ...@@ -399,7 +403,6 @@
FCBCCC6C2123073A00D94F7E /* BoxcoderKernel.swift */, FCBCCC6C2123073A00D94F7E /* BoxcoderKernel.swift */,
4AA1EA8D2146647F00D0F791 /* SplitKernel.swift */, 4AA1EA8D2146647F00D0F791 /* SplitKernel.swift */,
4AA1EA932146661500D0F791 /* ShapeKernel.swift */, 4AA1EA932146661500D0F791 /* ShapeKernel.swift */,
4AA1EA952146665A00D0F791 /* FlattenKernel.swift */,
4AA1EA87214662BD00D0F791 /* BilinearInterpKernel.swift */, 4AA1EA87214662BD00D0F791 /* BilinearInterpKernel.swift */,
FCBCCC70212309A700D94F7E /* MulticlassNMSKernel.swift */, FCBCCC70212309A700D94F7E /* MulticlassNMSKernel.swift */,
FCDDC6C5212F9FB800E5EF74 /* PreluKernel.swift */, FCDDC6C5212F9FB800E5EF74 /* PreluKernel.swift */,
...@@ -596,6 +599,7 @@ ...@@ -596,6 +599,7 @@
FCEBC0F420F1FDD90099DBAF /* ConvAddBatchNormReluOp.swift in Sources */, FCEBC0F420F1FDD90099DBAF /* ConvAddBatchNormReluOp.swift in Sources */,
FC0E2DC020EE461F009C1FAC /* ElementwiseAddKernel.swift in Sources */, FC0E2DC020EE461F009C1FAC /* ElementwiseAddKernel.swift in Sources */,
4AF928772133F1DB005B6C3A /* BoxCoder.metal in Sources */, 4AF928772133F1DB005B6C3A /* BoxCoder.metal in Sources */,
FC33B0F02147659000714A93 /* MobileNet.swift in Sources */,
FCEB684C212F093800D2448E /* PreluOp.swift in Sources */, FCEB684C212F093800D2448E /* PreluOp.swift in Sources */,
FCA67CD92138287B00BD58AA /* ConvBNReluKernel.metal in Sources */, FCA67CD92138287B00BD58AA /* ConvBNReluKernel.metal in Sources */,
FC60DB8920E9AAA500FF203F /* MetalExtension.swift in Sources */, FC60DB8920E9AAA500FF203F /* MetalExtension.swift in Sources */,
...@@ -622,7 +626,6 @@ ...@@ -622,7 +626,6 @@
FCD04E7420F3437E0007374F /* ConvAddKernel.swift in Sources */, FCD04E7420F3437E0007374F /* ConvAddKernel.swift in Sources */,
FC039BB920E11CC20081E9F8 /* Scope.swift in Sources */, FC039BB920E11CC20081E9F8 /* Scope.swift in Sources */,
FC292C5621421B4600CF622F /* PaddleMobileGPU.m in Sources */, FC292C5621421B4600CF622F /* PaddleMobileGPU.m in Sources */,
4AA1EA962146665A00D0F791 /* FlattenKernel.swift in Sources */,
FCD04E6620F314C50007374F /* PoolOp.swift in Sources */, FCD04E6620F314C50007374F /* PoolOp.swift in Sources */,
FC039BAC20E11CBC0081E9F8 /* BatchNormOp.swift in Sources */, FC039BAC20E11CBC0081E9F8 /* BatchNormOp.swift in Sources */,
FCBCCC6F2123097100D94F7E /* MulticlassNMSOp.swift in Sources */, FCBCCC6F2123097100D94F7E /* MulticlassNMSOp.swift in Sources */,
...@@ -655,6 +658,7 @@ ...@@ -655,6 +658,7 @@
FCD04E6C20F31A280007374F /* SoftmaxKernel.swift in Sources */, FCD04E6C20F31A280007374F /* SoftmaxKernel.swift in Sources */,
FCEB684A212F00DB00D2448E /* PreluKernel.metal in Sources */, FCEB684A212F00DB00D2448E /* PreluKernel.metal in Sources */,
4AA1EAA02148DEEE00D0F791 /* ReshapeKernel.metal.inc in Sources */, 4AA1EAA02148DEEE00D0F791 /* ReshapeKernel.metal.inc in Sources */,
FC9A19E32148C31300CD9CBF /* MobilenetSSD_AR.swift in Sources */,
FCDDC6CF212FE14700E5EF74 /* PriorBoxKernel.metal in Sources */, FCDDC6CF212FE14700E5EF74 /* PriorBoxKernel.metal in Sources */,
FC4CB74B20F12C30007C0C6D /* ProgramOptimize.swift in Sources */, FC4CB74B20F12C30007C0C6D /* ProgramOptimize.swift in Sources */,
FC5163F620EF556E00636C28 /* Texture2DTo2DArrayKernel.swift in Sources */, FC5163F620EF556E00636C28 /* Texture2DTo2DArrayKernel.swift in Sources */,
......
...@@ -13,16 +13,9 @@ ...@@ -13,16 +13,9 @@
limitations under the License. */ limitations under the License. */
import Foundation import Foundation
import paddle_mobile
class MobileNet: Net{ class MobileNet: Net{
var means: [Float] = [123.68, 116.78, 103.94]
var scale: Float = 0.017
let except: Int = 0
class MobilenetPreProccess: CusomKernel { class MobilenetPreProccess: CusomKernel {
init(device: MTLDevice) { init(device: MTLDevice) {
let s = CusomKernel.Shape.init(inWidth: 224, inHeight: 224, inChannel: 3) let s = CusomKernel.Shape.init(inWidth: 224, inHeight: 224, inChannel: 3)
...@@ -49,7 +42,7 @@ class MobileNet: Net{ ...@@ -49,7 +42,7 @@ class MobileNet: Net{
let labels = PreWords.init(fileName: "synset") let labels = PreWords.init(fileName: "synset")
func resultStr(res: [Float]) -> String { override public func resultStr(res: [Float]) -> String {
var s: [String] = [] var s: [String] = []
res.top(r: 5).enumerated().forEach{ res.top(r: 5).enumerated().forEach{
s.append(String(format: "%d: %@ (%3.2f%%)", $0 + 1, labels[$1.0], $1.1 * 100)) s.append(String(format: "%d: %@ (%3.2f%%)", $0 + 1, labels[$1.0], $1.1 * 100))
...@@ -57,17 +50,18 @@ class MobileNet: Net{ ...@@ -57,17 +50,18 @@ class MobileNet: Net{
return s.joined(separator: "\n") return s.joined(separator: "\n")
} }
var preprocessKernel: CusomKernel
let dim = (n: 1, h: 224, w: 224, c: 3)
let modelPath: String
let paramPath: String
let modelDir: String
init(device: MTLDevice) { override init(device: MTLDevice) {
super.init(device: device)
means = [123.68, 116.78, 103.94]
scale = 0.017
except = 0
modelPath = Bundle.main.path(forResource: "model", ofType: nil) ?! "model null" modelPath = Bundle.main.path(forResource: "model", ofType: nil) ?! "model null"
paramPath = Bundle.main.path(forResource: "params", ofType: nil) ?! "para null" paramPath = Bundle.main.path(forResource: "params", ofType: nil) ?! "para null"
modelDir = "" modelDir = ""
preprocessKernel = MobilenetPreProccess.init(device: device) preprocessKernel = MobilenetPreProccess.init(device: device)
dim = (n: 1, h: 224, w: 224, c: 3)
} }
} }
/* 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
public class MobileNet_ssd_AR: Net{
@objc public override init(device: MTLDevice) {
super.init(device: device)
means = [103.94, 116.78, 123.68]
scale = 1
except = 2
modelPath = Bundle.main.path(forResource: "ar_model", ofType: nil) ?! "model null"
paramPath = Bundle.main.path(forResource: "ar_params", ofType: nil) ?! "para null"
modelDir = ""
preprocessKernel = MobilenetssdPreProccess.init(device: device)
dim = (n: 1, h: 160, w: 160, c: 3)
}
class MobilenetssdPreProccess: CusomKernel {
init(device: MTLDevice) {
let s = CusomKernel.Shape.init(inWidth: 160, inHeight: 160, inChannel: 3)
super.init(device: device, inFunctionName: "mobilent_ar_preprocess_half", outputDim: s, usePaddleMobileLib: false)
}
}
override public func resultStr(res: [Float]) -> String {
return " \(res)"
}
override func fetchResult(paddleMobileRes: ResultHolder) -> [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()
}
var scoreFormatArr: [Float32] = score.metalTexture.realNHWC(dim: (n: score.padToFourDim[0], h: score.padToFourDim[1], w: score.padToFourDim[2], c: score.padToFourDim[3]))
// print("score: ")
// print(scoreFormatArr.strideArray())
//
var bboxArr = bbox.metalTexture.float32Array()
// print("bbox: ")
// print(bboxArr.strideArray())
let nmsCompute = NMSCompute.init()
nmsCompute.scoreThredshold = 0.01
nmsCompute.nmsTopK = 400
nmsCompute.keepTopK = 200
nmsCompute.nmsEta = 1.0
nmsCompute.nmsThreshold = 0.45
nmsCompute.background_label = 0;
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 }
return output
}
}
...@@ -38,7 +38,8 @@ extension Runable where Self: OperatorProtocol{ ...@@ -38,7 +38,8 @@ extension Runable where Self: OperatorProtocol{
} }
func inputVariant() -> [String : [Variant]] { func inputVariant() -> [String : [Variant]] {
fatalError(" op \(type) need implement inputVariant") return [:]
// fatalError(" op \(type) need implement inputVariant")
} }
func delogOutput() { func delogOutput() {
...@@ -166,7 +167,7 @@ let opInfos = [gConvType : (inputs: ["Input"], outputs: ["Out ...@@ -166,7 +167,7 @@ let opInfos = [gConvType : (inputs: ["Input"], outputs: ["Out
gPreluType : (inputs: ["X"], outputs: ["Out"]), gPreluType : (inputs: ["X"], outputs: ["Out"]),
gConv2dTransposeType : (inputs: ["Input"], outputs: ["Output"]), gConv2dTransposeType : (inputs: ["Input"], outputs: ["Output"]),
gBilinearInterpType : (inputs: ["X"], outputs: ["Out"]), gBilinearInterpType : (inputs: ["X"], outputs: ["Out"]),
gSplit : (inputs: ["Input"], outputs: ["Out"]), gSplit : (inputs: ["X"], outputs: ["Out"]),
gShape : (inputs: ["Input"], outputs: ["Out"]), gShape : (inputs: ["Input"], outputs: ["Out"]),
gFlatten : (inputs: ["Input"], outputs: ["Out"]) gFlatten : (inputs: ["X"], outputs: ["Out"])
] ]
...@@ -14,21 +14,9 @@ ...@@ -14,21 +14,9 @@
import Foundation import Foundation
class FlattenParam<P: PrecisionType>: OpParam { class FlattenOp<P: PrecisionType>: Operator<ReshapeKernel<P>, ReshapeParam<P>>, Runable, Creator, InferShaperable{
typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws {
do {
output = try FlattenParam.output(outputs: opDesc.outputs, from: inScope)
} catch let error {
throw error
}
}
var output: Texture<P>
}
class FlattenOp<P: PrecisionType>: Operator<FlattenKernel<P>, FlattenParam<P>>, Runable, Creator, InferShaperable{
typealias OpType = SplitOp<P> typealias OpType = FlattenOp<P>
func inferShape() { func inferShape() {
// para.output.dim = para.input.dim // para.output.dim = para.input.dim
......
/* 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
struct FlattenMetalParam {
}
class FlattenKernel<P: PrecisionType>: Kernel, Computable{
func compute(commandBuffer: MTLCommandBuffer, param: FlattenParam<P>) throws {
guard let encoder = commandBuffer.makeComputeCommandEncoder() else {
throw PaddleMobileError.predictError(message: " encode is nil")
}
encoder.setTexture(param.output.metalTexture, index: 0)
encoder.endEncoding()
}
required init(device: MTLDevice, param: FlattenParam<P>) {
param.output.initTexture(device: device, computePrecision: computePrecision)
if computePrecision == .Float32 {
super.init(device: device, inFunctionName: "split")
} else if computePrecision == .Float16 {
super.init(device: device, inFunctionName: "split_half")
} else {
fatalError()
}
}
}
...@@ -18,7 +18,8 @@ class SplitParam<P: PrecisionType>: OpParam { ...@@ -18,7 +18,8 @@ class SplitParam<P: PrecisionType>: OpParam {
typealias ParamPrecisionType = P typealias ParamPrecisionType = P
required init(opDesc: OpDesc, inScope: Scope) throws { required init(opDesc: OpDesc, inScope: Scope) throws {
do { do {
output = try SplitParam.output(outputs: opDesc.outputs, from: inScope) // output = try SplitParam.output(outputs: opDesc.outputs, from: inScope)
output = try SplitParam.outputOut(outputs: opDesc.outputs, from: inScope)
} catch let error { } catch let error {
throw error throw error
} }
......
...@@ -16,7 +16,7 @@ import Foundation ...@@ -16,7 +16,7 @@ import Foundation
class ScaleKernel: CusomKernel { class ScaleKernel: CusomKernel {
init(device: MTLDevice, shape: Shape) { init(device: MTLDevice, shape: Shape) {
super.init(device: device, inFunctionName: "scale", outputDim: shape, usePaddleMobileLib: false) super.init(device: device, inFunctionName: "scale_half", outputDim: shape, usePaddleMobileLib: false)
} }
} }
...@@ -37,7 +37,7 @@ public class Net: NSObject { ...@@ -37,7 +37,7 @@ public class Net: NSObject {
fatalError() fatalError()
} }
func fetchResult(paddleMobileRes: ResultHolder) -> [Float32] { func fetchResult(paddleMobileRes: ResultHolder) -> [Float32] {
fatalError() return paddleMobileRes.resultArr
} }
@objc public init(device: MTLDevice) { @objc public init(device: MTLDevice) {
super.init() super.init()
......
...@@ -14,10 +14,9 @@ ...@@ -14,10 +14,9 @@
import Foundation import Foundation
let testTo = 161 let testTo = 3
var isTest = false var isTest = false
let computePrecision: ComputePrecision = .Float16 let computePrecision: ComputePrecision = .Float16
public class ResultHolder { public class ResultHolder {
...@@ -59,8 +58,10 @@ public class Executor<P: PrecisionType> { ...@@ -59,8 +58,10 @@ public class Executor<P: PrecisionType> {
var ops: [Runable & InferShaperable] = [] var ops: [Runable & InferShaperable] = []
let program: Program let program: Program
let device: MTLDevice let device: MTLDevice
let inflightSemaphore: DispatchSemaphore
let queue: MTLCommandQueue let queue: MTLCommandQueue
public init(inDevice:MTLDevice, inQueue: MTLCommandQueue, inProgram: Program) throws { public init(inDevice:MTLDevice, inQueue: MTLCommandQueue, inProgram: Program) throws {
self.inflightSemaphore = DispatchSemaphore(value: 3)
program = inProgram program = inProgram
device = inDevice device = inDevice
queue = inQueue queue = inQueue
...@@ -82,6 +83,8 @@ public class Executor<P: PrecisionType> { ...@@ -82,6 +83,8 @@ public class Executor<P: PrecisionType> {
guard let buffer = queue.makeCommandBuffer() else { guard let buffer = queue.makeCommandBuffer() else {
throw PaddleMobileError.predictError(message: "CommandBuffer is nil") throw PaddleMobileError.predictError(message: "CommandBuffer is nil")
} }
inflightSemaphore.wait()
let resInput: MTLTexture let resInput: MTLTexture
if let inPre = preProcessKernle { if let inPre = preProcessKernle {
do { do {
...@@ -112,12 +115,14 @@ public class Executor<P: PrecisionType> { ...@@ -112,12 +115,14 @@ public class Executor<P: PrecisionType> {
outputTextures = ops[ops.count - except].inputVariant() outputTextures = ops[ops.count - except].inputVariant()
} }
buffer.addCompletedHandler { (commandbuffer) in buffer.addCompletedHandler { [weak self] (commandbuffer) in
// let inputArr = resInput.toTensor(dim: (n: dim[0], c: dim[3], h: dim[1], w: dim[2])) // let inputArr = resInput.toTensor(dim: (n: dim[0], c: dim[3], h: dim[1], w: dim[2]))
//// print(inputArr.strideArray()) //// print(inputArr.strideArray())
//
// writeToLibrary(fileName: "test_image_ssd", array: inputArr) // print(dim)
// writeToLibrary(fileName: "test_image_ssd_ar", array: inputArr)
// print("write to library done") // print("write to library done")
// return // return
// print(inputArr) // print(inputArr)
...@@ -139,18 +144,23 @@ public class Executor<P: PrecisionType> { ...@@ -139,18 +144,23 @@ public class Executor<P: PrecisionType> {
// return // return
guard let SSelf = self else {
fatalError()
}
let afterDate = Date.init() let afterDate = Date.init()
var resultHolder: ResultHolder var resultHolder: ResultHolder
if except > 0 { if except > 0 {
resultHolder = ResultHolder.init(inDim: [], inResult: [], inElapsedTime: afterDate.timeIntervalSince(beforeDate), inIntermediateResults: outputTextures) resultHolder = ResultHolder.init(inDim: [], inResult: [], inElapsedTime: afterDate.timeIntervalSince(beforeDate), inIntermediateResults: outputTextures)
} else { } else {
let outputVar: Variant = self.program.scope.output()! let outputVar: Variant = SSelf.program.scope.output()!
let output: Texture<P> = outputVar as! Texture<P> let output: Texture<P> = outputVar as! Texture<P>
resultHolder = ResultHolder.init(inDim: output.dim.dims, inResult: output.toTensor(), inElapsedTime: afterDate.timeIntervalSince(beforeDate)) resultHolder = ResultHolder.init(inDim: output.dim.dims, inResult: output.toTensor(), inElapsedTime: afterDate.timeIntervalSince(beforeDate))
} }
completionHandle(resultHolder) completionHandle(resultHolder)
SSelf.inflightSemaphore.signal()
} }
buffer.commit() buffer.commit()
} }
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册