From d588dcd98729e8daac963ceb7fee0100c98a333e Mon Sep 17 00:00:00 2001
From: Piotr Sowa <piotrsowa@Piotr-MacBook-Pro.local>
Date: Sun, 27 Jan 2019 21:42:03 +0100
Subject: [PATCH] DarkNet OpenCL Port

---
 .idea/codeStyles/Project.xml |  29 --
 .idea/workspace.xml          | 222 ++++++----
 CMakeLists.txt               |   2 +-
 Makefile                     |   4 +-
 README.md                    |   6 +
 cfg/openimages.data          |   8 +
 cfg/resnext101-32x4d.cfg     |  19 +-
 cfg/resnext152-32x4d.cfg     |  18 +-
 cfg/vgg-16.cfg               |  10 +-
 cfg/yolov3-openimages.cfg    | 789 +++++++++++++++++++++++++++++++++++
 data/openimages.names        | 601 ++++++++++++++++++++++++++
 examples/darknet.c           |   2 -
 examples/detector.c          |   7 +-
 include/darknet.h            |   2 +-
 src/activation_kernels.cl    |  10 +-
 src/activations.c            |   7 +
 src/activations.h            |   2 +
 src/blas_kernels.cl          | 121 ++++--
 src/connected_layer.c        |   3 +-
 src/convolutional_kernels.c  |   2 +-
 src/convolutional_layer.c    |  76 +++-
 src/cost_layer.c             |  10 +-
 src/crop_layer.c             |   7 +-
 src/data.c                   |   8 +-
 src/deconvolutional_layer.c  |  14 +-
 src/detection_layer.c        |   6 +-
 src/dropout_layer.c          |   6 +-
 src/image_opencv.cpp         | 135 ++++++
 src/iseg_layer.c             |  22 +-
 src/layer.c                  | 107 ++---
 src/local_layer.c            |   5 +-
 src/lstm_layer.c             |  36 +-
 src/maxpool_layer.c          |  12 +-
 src/network.c                |  51 +--
 src/normalization_layer.c    |  12 +-
 src/opencl.c                 | 155 +++++--
 src/opencl.h                 |   1 +
 src/parser.c                 |  33 +-
 src/region_layer.c           |  84 ++--
 src/reorg_layer.c            |  10 +-
 src/rnn_layer.c              |   4 +-
 src/route_layer.c            |   4 +-
 src/shortcut_layer.c         |  19 +-
 src/softmax_layer.c          |  16 +-
 src/upsample_layer.c         |  15 +-
 src/utils.c                  |   8 +-
 src/yolo_layer.c             |  39 +-
 47 files changed, 2271 insertions(+), 488 deletions(-)
 delete mode 100644 .idea/codeStyles/Project.xml
 create mode 100644 cfg/openimages.data
 create mode 100644 cfg/yolov3-openimages.cfg
 create mode 100644 data/openimages.names
 create mode 100644 src/image_opencv.cpp

diff --git a/.idea/codeStyles/Project.xml b/.idea/codeStyles/Project.xml
deleted file mode 100644
index 30aa626c..00000000
--- a/.idea/codeStyles/Project.xml
+++ /dev/null
@@ -1,29 +0,0 @@
-<component name="ProjectCodeStyleConfiguration">
-  <code_scheme name="Project" version="173">
-    <Objective-C-extensions>
-      <file>
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="Import" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="Macro" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="Typedef" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="Enum" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="Constant" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="Global" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="Struct" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="FunctionPredecl" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="Function" />
-      </file>
-      <class>
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="Property" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="Synthesize" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="InitMethod" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="StaticMethod" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="InstanceMethod" />
-        <option name="com.jetbrains.cidr.lang.util.OCDeclarationKind" value="DeallocMethod" />
-      </class>
-      <extensions>
-        <pair source="cpp" header="h" fileNamingConvention="NONE" />
-        <pair source="c" header="h" fileNamingConvention="NONE" />
-      </extensions>
-    </Objective-C-extensions>
-  </code_scheme>
-</component>
\ No newline at end of file
diff --git a/.idea/workspace.xml b/.idea/workspace.xml
index 73615c91..6242cc33 100644
--- a/.idea/workspace.xml
+++ b/.idea/workspace.xml
@@ -1,10 +1,10 @@
 <?xml version="1.0" encoding="UTF-8"?>
 <project version="4">
-  <component name="CMakeRunConfigurationManager" shouldGenerate="true" shouldDeleteObsolete="true" buildAllGenerated="true">
+  <component name="CMakeRunConfigurationManager" shouldGenerate="true" shouldDeleteObsolete="true">
     <generated>
-      <config projectName="darknet" targetName="bindarknet" />
       <config projectName="darknet" targetName="libdarknet_s" />
       <config projectName="darknet" targetName="libdarknet" />
+      <config projectName="darknet" targetName="bindarknet" />
     </generated>
   </component>
   <component name="CMakeSettings">
@@ -13,7 +13,9 @@
     </configurations>
   </component>
   <component name="ChangeListManager">
-    <list default="true" id="5c1f2602-dc6e-4bd5-b03e-d07ee3b8a66f" name="Default Changelist" comment="" />
+    <list default="true" id="5c1f2602-dc6e-4bd5-b03e-d07ee3b8a66f" name="Default Changelist" comment="">
+      <change beforePath="$PROJECT_DIR$/src/local_layer.c" beforeDir="false" afterPath="$PROJECT_DIR$/src/local_layer.c" afterDir="false" />
+    </list>
     <ignored path="$PROJECT_DIR$/cmake-build-debug/" />
     <option name="EXCLUDED_CONVERTED_TO_IGNORED" value="true" />
     <option name="SHOW_DIALOG" value="false" />
@@ -22,49 +24,16 @@
     <option name="LAST_RESOLUTION" value="IGNORE" />
   </component>
   <component name="ExecutionTargetManager" SELECTED_TARGET="CMakeBuildProfile:Debug" />
-  <component name="FUSProjectUsageTrigger">
-    <session id="-1374244216">
-      <usages-collector id="statistics.lifecycle.project">
-        <counts>
-          <entry key="project.closed" value="4" />
-          <entry key="project.open.time.5" value="1" />
-          <entry key="project.open.time.7" value="3" />
-          <entry key="project.opened" value="4" />
-        </counts>
-      </usages-collector>
-      <usages-collector id="statistics.file.extensions.open">
-        <counts>
-          <entry key="Makefile" value="1" />
-          <entry key="c" value="5" />
-          <entry key="txt" value="1" />
-        </counts>
-      </usages-collector>
-      <usages-collector id="statistics.file.types.open">
-        <counts>
-          <entry key="CMakeLists.txt" value="1" />
-          <entry key="ObjectiveC" value="5" />
-          <entry key="PLAIN_TEXT" value="1" />
-        </counts>
-      </usages-collector>
-      <usages-collector id="statistics.file.extensions.edit">
-        <counts>
-          <entry key="c" value="16" />
-        </counts>
-      </usages-collector>
-      <usages-collector id="statistics.file.types.edit">
-        <counts>
-          <entry key="ObjectiveC" value="16" />
-        </counts>
-      </usages-collector>
-    </session>
-  </component>
   <component name="FileEditorManager">
     <leaf SIDE_TABS_SIZE_LIMIT_KEY="300">
       <file pinned="false" current-in-tab="true">
         <entry file="file://$PROJECT_DIR$/src/network.c">
           <provider selected="true" editor-type-id="text-editor">
-            <state relative-caret-position="15">
-              <caret line="772" selection-start-line="772" selection-end-line="772" />
+            <state relative-caret-position="97">
+              <caret line="784" selection-start-line="784" selection-end-line="784" />
+              <folding>
+                <element signature="e#11230#11299#0" expanded="true" />
+              </folding>
             </state>
           </provider>
         </entry>
@@ -72,23 +41,24 @@
       <file pinned="false" current-in-tab="false">
         <entry file="file://$PROJECT_DIR$/src/opencl.c">
           <provider selected="true" editor-type-id="text-editor">
-            <state relative-caret-position="15">
-              <caret line="1" selection-start-line="1" selection-end-line="1" />
-            </state>
-          </provider>
-        </entry>
-      </file>
-      <file pinned="false" current-in-tab="false">
-        <entry file="file://$PROJECT_DIR$/src/opencl.h">
-          <provider selected="true" editor-type-id="text-editor">
-            <state relative-caret-position="30">
-              <caret line="2" selection-start-line="2" selection-end-line="2" />
+            <state relative-caret-position="47">
+              <caret line="922" selection-start-line="922" selection-end-line="922" />
             </state>
           </provider>
         </entry>
       </file>
     </leaf>
   </component>
+  <component name="FindInProjectRecents">
+    <findStrings>
+      <find>free</find>
+      <find>//opencl_</find>
+      <find>size_t get_workspace_size</find>
+      <find>free(</find>
+      <find>//opencl_pull</find>
+      <find>train_net</find>
+    </findStrings>
+  </component>
   <component name="Git.Settings">
     <option name="RECENT_GIT_ROOT_PATH" value="$PROJECT_DIR$" />
   </component>
@@ -98,28 +68,26 @@
         <option value="$PROJECT_DIR$/src/parser.c" />
         <option value="$PROJECT_DIR$/CMakeLists.txt" />
         <option value="$PROJECT_DIR$/Makefile" />
+        <option value="$PROJECT_DIR$/src/data.c" />
+        <option value="$PROJECT_DIR$/src/convolutional_kernels.c" />
         <option value="$PROJECT_DIR$/src/network.c" />
+        <option value="$PROJECT_DIR$/src/detection_layer.c" />
+        <option value="$PROJECT_DIR$/src/opencl.c" />
+        <option value="$PROJECT_DIR$/src/local_layer.c" />
       </list>
     </option>
   </component>
-  <component name="JsBuildToolGruntFileManager" detection-done="true" sorting="DEFINITION_ORDER" />
-  <component name="JsBuildToolPackageJson" detection-done="true" sorting="DEFINITION_ORDER" />
-  <component name="JsGulpfileManager">
-    <detection-done>true</detection-done>
-    <sorting>DEFINITION_ORDER</sorting>
-  </component>
   <component name="OCFindUsagesOptions" text="false" ivars="false" properties="true" derivedClasses="false" />
-  <component name="ProjectFrameBounds">
+  <component name="ProjectFrameBounds" extendedState="6">
     <option name="y" value="23" />
     <option name="width" value="1440" />
-    <option name="height" value="836" />
+    <option name="height" value="839" />
   </component>
   <component name="ProjectView">
     <navigator proportions="" version="1">
       <foldersAlwaysOnTop value="true" />
     </navigator>
     <panes>
-      <pane id="Scope" />
       <pane id="ProjectPane">
         <subPane>
           <expand>
@@ -127,10 +95,16 @@
               <item name="darknet" type="b2602c69:ProjectViewProjectNode" />
               <item name="darknet" type="462c0819:PsiDirectoryNode" />
             </path>
+            <path>
+              <item name="darknet" type="b2602c69:ProjectViewProjectNode" />
+              <item name="darknet" type="462c0819:PsiDirectoryNode" />
+              <item name="src" type="462c0819:PsiDirectoryNode" />
+            </path>
           </expand>
           <select />
         </subPane>
       </pane>
+      <pane id="Scope" />
     </panes>
   </component>
   <component name="PropertiesComponent">
@@ -151,7 +125,7 @@
       </list>
     </option>
   </component>
-  <component name="RunManager" selected="Application.bindarknet (cifar)">
+  <component name="RunManager" selected="Application.bindarknet (Yv1)">
     <configuration name="Build All" type="CMakeRunConfiguration" factoryName="Application" PASS_PARENT_ENVS_2="true" CONFIG_NAME="Debug" RUN_TARGET_PROJECT_NAME="darknet" RUN_TARGET_NAME="bindarknet" EXPLICIT_BUILD_TARGET_NAME="all">
       <method v="2">
         <option name="com.jetbrains.cidr.execution.CidrBuildBeforeRunTaskProvider$BuildBeforeRunTask" enabled="true" />
@@ -207,7 +181,7 @@
         <option name="com.jetbrains.cidr.execution.CidrBuildBeforeRunTaskProvider$BuildBeforeRunTask" enabled="true" />
       </method>
     </configuration>
-    <configuration name="bindarknet (tYv2)" type="CMakeRunConfiguration" factoryName="Application" singleton="true" PROGRAM_PARAMS="detector train voc.data yolo-voc.2.0.cfg darknet19_448.conv.23 -i 1" WORKING_DIR="file://$PROJECT_DIR$/../train" PASS_PARENT_ENVS_2="true" PROJECT_NAME="darknet" TARGET_NAME="bindarknet" CONFIG_NAME="Debug" RUN_TARGET_PROJECT_NAME="darknet" RUN_TARGET_NAME="bindarknet">
+    <configuration name="bindarknet (tYv2)" type="CMakeRunConfiguration" factoryName="Application" singleton="true" PROGRAM_PARAMS="detector train voc.data yolov2-voc.cfg darknet19_448.conv.23 -i 1" WORKING_DIR="file://$PROJECT_DIR$/../train" PASS_PARENT_ENVS_2="true" PROJECT_NAME="darknet" TARGET_NAME="bindarknet" CONFIG_NAME="Debug" RUN_TARGET_PROJECT_NAME="darknet" RUN_TARGET_NAME="bindarknet">
       <method v="2">
         <option name="com.jetbrains.cidr.execution.CidrBuildBeforeRunTaskProvider$BuildBeforeRunTask" enabled="true" />
       </method>
@@ -296,23 +270,27 @@
       <workItem from="1533072395045" duration="2200000" />
       <workItem from="1533076392813" duration="90000" />
       <workItem from="1533465199569" duration="295000" />
+      <workItem from="1544973155534" duration="4931000" />
+      <workItem from="1545004229098" duration="268000" />
+      <workItem from="1545084711013" duration="8346000" />
+      <workItem from="1545094152381" duration="1702000" />
     </task>
     <servers />
   </component>
   <component name="TimeTrackingManager">
-    <option name="totallyTimeSpent" value="578816000" />
+    <option name="totallyTimeSpent" value="594063000" />
   </component>
   <component name="ToolWindowManager">
-    <frame x="0" y="23" width="1440" height="836" extended-state="0" />
+    <frame x="0" y="23" width="1440" height="839" extended-state="6" />
     <editor active="true" />
     <layout>
-      <window_info content_ui="combo" id="Project" order="0" visible="true" weight="0.2639485" />
+      <window_info content_ui="combo" id="Project" order="0" visible="true" weight="0.27825466" />
       <window_info id="Structure" order="1" side_tool="true" weight="0.25" />
       <window_info id="Favorites" order="2" side_tool="true" />
       <window_info anchor="bottom" id="Message" order="0" />
       <window_info anchor="bottom" id="Find" order="1" />
-      <window_info anchor="bottom" id="Run" order="2" weight="0.3293011" />
-      <window_info anchor="bottom" id="Debug" order="3" weight="0.39919356" />
+      <window_info anchor="bottom" id="Run" order="2" weight="0.32797858" />
+      <window_info anchor="bottom" id="Debug" order="3" weight="0.39892906" />
       <window_info anchor="bottom" id="Cvs" order="4" weight="0.25" />
       <window_info anchor="bottom" id="Inspection" order="5" weight="0.4" />
       <window_info anchor="bottom" id="TODO" order="6" />
@@ -321,7 +299,7 @@
       <window_info anchor="bottom" id="Version Control" order="9" />
       <window_info anchor="bottom" id="CMake" order="10" />
       <window_info anchor="bottom" id="Database Changes" order="11" show_stripe_button="false" />
-      <window_info anchor="bottom" id="Messages" order="12" weight="0.3293011" />
+      <window_info anchor="bottom" id="Messages" order="12" weight="0.32797858" />
       <window_info anchor="right" id="Commander" order="0" weight="0.4" />
       <window_info anchor="right" id="Ant Build" order="1" weight="0.25" />
       <window_info anchor="right" content_ui="combo" id="Hierarchy" order="2" weight="0.25" />
@@ -334,17 +312,7 @@
   <component name="UnknownFeatures">
     <option featureType="com.intellij.fileTypeFactory" implementationName="*.cfg" />
   </component>
-  <component name="VcsContentAnnotationSettings">
-    <option name="myLimit" value="2678400000" />
-  </component>
   <component name="editorHistoryManager">
-    <entry file="file://$PROJECT_DIR$/examples/yolo.c">
-      <provider selected="true" editor-type-id="text-editor">
-        <state relative-caret-position="262">
-          <caret line="329" column="43" lean-forward="true" selection-start-line="329" selection-start-column="43" selection-end-line="329" selection-end-column="43" />
-        </state>
-      </provider>
-    </entry>
     <entry file="file://$PROJECT_DIR$/src/parser.c">
       <provider selected="true" editor-type-id="text-editor">
         <state relative-caret-position="94">
@@ -387,26 +355,104 @@
         </state>
       </provider>
     </entry>
-    <entry file="file://$PROJECT_DIR$/src/network.c">
+    <entry file="file://$PROJECT_DIR$/src/data.c">
+      <provider selected="true" editor-type-id="text-editor">
+        <state relative-caret-position="118">
+          <caret line="704" column="37" lean-forward="true" selection-start-line="704" selection-start-column="37" selection-end-line="704" selection-end-column="37" />
+        </state>
+      </provider>
+    </entry>
+    <entry file="file://$PROJECT_DIR$/src/opencl.h">
       <provider selected="true" editor-type-id="text-editor">
-        <state relative-caret-position="15">
-          <caret line="772" selection-start-line="772" selection-end-line="772" />
+        <state relative-caret-position="30">
+          <caret line="2" selection-start-line="2" selection-end-line="2" />
+        </state>
+      </provider>
+    </entry>
+    <entry file="file://$PROJECT_DIR$/src/connected_layer.c">
+      <provider selected="true" editor-type-id="text-editor">
+        <state relative-caret-position="128">
+          <caret line="307" selection-start-line="307" selection-end-line="307" />
+        </state>
+      </provider>
+    </entry>
+    <entry file="file://$PROJECT_DIR$/src/dropout_layer_kernels.c">
+      <provider selected="true" editor-type-id="text-editor">
+        <state relative-caret-position="159">
+          <caret line="52" selection-start-line="52" selection-end-line="52" />
+        </state>
+      </provider>
+    </entry>
+    <entry file="file://$PROJECT_DIR$/src/maxpool_layer_kernels.c">
+      <provider selected="true" editor-type-id="text-editor">
+        <state relative-caret-position="227">
+          <caret line="53" selection-start-line="53" selection-end-line="53" />
+        </state>
+      </provider>
+    </entry>
+    <entry file="file://$PROJECT_DIR$/src/region_layer.c">
+      <provider selected="true" editor-type-id="text-editor">
+        <state relative-caret-position="93">
+          <caret line="479" selection-start-line="479" selection-end-line="481" selection-end-column="65" />
+          <folding>
+            <element signature="e#0#25#0" expanded="true" />
+          </folding>
+        </state>
+      </provider>
+    </entry>
+    <entry file="file://$PROJECT_DIR$/src/detection_layer.c">
+      <provider selected="true" editor-type-id="text-editor">
+        <state relative-caret-position="3072">
+          <caret line="268" column="36" selection-start-line="268" selection-start-column="36" selection-end-line="268" selection-end-column="36" />
+        </state>
+      </provider>
+    </entry>
+    <entry file="file://$PROJECT_DIR$/examples/yolo.c">
+      <provider selected="true" editor-type-id="text-editor">
+        <state relative-caret-position="251">
+          <caret line="59" selection-start-line="59" selection-end-line="59" />
+        </state>
+      </provider>
+    </entry>
+    <entry file="file://$PROJECT_DIR$/src/upsample_layer.c">
+      <provider selected="true" editor-type-id="text-editor">
+        <state>
+          <folding>
+            <element signature="e#0#27#0" expanded="true" />
+          </folding>
+        </state>
+      </provider>
+    </entry>
+    <entry file="file://$PROJECT_DIR$/src/convolutional_kernels.c">
+      <provider selected="true" editor-type-id="text-editor">
+        <state relative-caret-position="1115">
+          <caret line="258" selection-start-line="258" selection-end-line="258" />
+        </state>
+      </provider>
+    </entry>
+    <entry file="file://$PROJECT_DIR$/src/local_layer.c">
+      <provider selected="true" editor-type-id="text-editor">
+        <state relative-caret-position="216">
+          <caret line="242" column="55" selection-start-line="242" selection-start-column="55" selection-end-line="242" selection-end-column="55" />
         </state>
       </provider>
     </entry>
     <entry file="file://$PROJECT_DIR$/src/opencl.c">
       <provider selected="true" editor-type-id="text-editor">
-        <state relative-caret-position="15">
-          <caret line="1" selection-start-line="1" selection-end-line="1" />
+        <state relative-caret-position="47">
+          <caret line="922" selection-start-line="922" selection-end-line="922" />
         </state>
       </provider>
     </entry>
-    <entry file="file://$PROJECT_DIR$/src/opencl.h">
+    <entry file="file://$PROJECT_DIR$/src/network.c">
       <provider selected="true" editor-type-id="text-editor">
-        <state relative-caret-position="30">
-          <caret line="2" selection-start-line="2" selection-end-line="2" />
+        <state relative-caret-position="97">
+          <caret line="784" selection-start-line="784" selection-end-line="784" />
+          <folding>
+            <element signature="e#11230#11299#0" expanded="true" />
+          </folding>
         </state>
       </provider>
     </entry>
   </component>
-</project>
+</project>
\ No newline at end of file
diff --git a/CMakeLists.txt b/CMakeLists.txt
index b91f454a..e31ff938 100755
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -36,7 +36,7 @@ else()
 	endif()
 	list (APPEND DARKNET_CFLAGS "-Wfatal-errors")
 	list (APPEND DARKNET_CFLAGS "-Wno-unused-variable")
-	list (APPEND DARKNET_CFLAGS "-O0")
+	list (APPEND DARKNET_CFLAGS "-O2")
 endif()
 
 # Find OpenCV
diff --git a/Makefile b/Makefile
index 31b2f01f..085eb287 100644
--- a/Makefile
+++ b/Makefile
@@ -54,9 +54,9 @@ ifeq ($(DEBUG), 1)
 OPTS=-O0 -g
 else
 ifeq ($(RPI), 1)
-OPTS=-O0
+OPTS=-O2
 else
-OPTS=-O0
+OPTS=-O2
 endif
 endif
 
diff --git a/README.md b/README.md
index 2045df5f..85e22fe3 100644
--- a/README.md
+++ b/README.md
@@ -1 +1,7 @@
 [![DarkNet Training](https://img.youtube.com/vi/Mxw7XkFBFPc/0.jpg)](https://www.youtube.com/watch?v=Mxw7XkFBFPc)
+
+https://codingbytodesign.net/2018/05/26/darknet-training/ 
+
+https://codingbytodesign.net/2018/08/01/darknet-in-opencl/
+
+Thanks!
diff --git a/cfg/openimages.data b/cfg/openimages.data
new file mode 100644
index 00000000..fa80e5ab
--- /dev/null
+++ b/cfg/openimages.data
@@ -0,0 +1,8 @@
+classes= 601
+train  = /home/pjreddie/data/openimsv4/openimages.train.list
+#valid  = coco_testdev
+valid = data/coco_val_5k.list
+names = data/openimages.names
+backup = /home/pjreddie/backup/
+eval=coco
+
diff --git a/cfg/resnext101-32x4d.cfg b/cfg/resnext101-32x4d.cfg
index 382836d1..8538ccc3 100644
--- a/cfg/resnext101-32x4d.cfg
+++ b/cfg/resnext101-32x4d.cfg
@@ -1,21 +1,25 @@
 [net]
-batch=128
-subdivisions=8
- batch=1
- subdivisions=1
+# Training
+# batch=128
+# subdivisions=8
+
+# Testing
+batch=1
+subdivisions=1
+
 height=256
 width=256
+channels=3
 min_crop=128
 max_crop=448
-channels=3
-momentum=0.9
-decay=0.0005
 
 burn_in=1000
 learning_rate=0.1
 policy=poly
 power=4
 max_batches=800000
+momentum=0.9
+decay=0.0005
 
 angle=7
 hue=.1
@@ -24,6 +28,7 @@ exposure=.75
 aspect=.75
 
 
+
 [convolutional]
 batch_normalize=1
 filters=64
diff --git a/cfg/resnext152-32x4d.cfg b/cfg/resnext152-32x4d.cfg
index 1296cf9c..48279fd2 100644
--- a/cfg/resnext152-32x4d.cfg
+++ b/cfg/resnext152-32x4d.cfg
@@ -1,21 +1,25 @@
 [net]
-batch=128
-subdivisions=16
-# batch=1
-# subdivisions=1
+# Training
+# batch=128
+# subdivisions=16
+
+# Testing
+batch=1
+subdivisions=1
+
 height=256
 width=256
+channels=3
 min_crop=128
 max_crop=448
-channels=3
-momentum=0.9
-decay=0.0005
 
 burn_in=1000
 learning_rate=0.1
 policy=poly
 power=4
 max_batches=800000
+momentum=0.9
+decay=0.0005
 
 angle=7
 hue=.1
diff --git a/cfg/vgg-16.cfg b/cfg/vgg-16.cfg
index 79694b87..c73b17b6 100644
--- a/cfg/vgg-16.cfg
+++ b/cfg/vgg-16.cfg
@@ -1,6 +1,12 @@
 [net]
-batch=128
-subdivisions=4
+# Training
+# batch=128
+# subdivisions=4
+
+# Testing
+batch=1
+subdivisions=1
+
 height=256
 width=256
 channels=3
diff --git a/cfg/yolov3-openimages.cfg b/cfg/yolov3-openimages.cfg
new file mode 100644
index 00000000..65d241a7
--- /dev/null
+++ b/cfg/yolov3-openimages.cfg
@@ -0,0 +1,789 @@
+[net]
+# Testing
+ batch=1
+ subdivisions=1
+# Training
+batch=64
+subdivisions=16
+width=608
+height=608
+channels=3
+momentum=0.9
+decay=0.0005
+angle=0
+saturation = 1.5
+exposure = 1.5
+hue=.1
+
+learning_rate=0.001
+burn_in=5000
+max_batches = 500200
+policy=steps
+steps=400000,450000
+scales=.1,.1
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=leaky
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+######################
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=1818
+activation=linear
+
+
+[yolo]
+mask = 6,7,8
+anchors = 10,13,  16,30,  33,23,  30,61,  62,45,  59,119,  116,90,  156,198,  373,326
+classes=601
+num=9
+jitter=.3
+ignore_thresh = .7
+truth_thresh = 1
+random=1
+
+
+[route]
+layers = -4
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[upsample]
+stride=2
+
+[route]
+layers = -1, 61
+
+
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=1818
+activation=linear
+
+
+[yolo]
+mask = 3,4,5
+anchors = 10,13,  16,30,  33,23,  30,61,  62,45,  59,119,  116,90,  156,198,  373,326
+classes=601
+num=9
+jitter=.3
+ignore_thresh = .7
+truth_thresh = 1
+random=1
+
+
+
+[route]
+layers = -4
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[upsample]
+stride=2
+
+[route]
+layers = -1, 36
+
+
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=1818
+activation=linear
+
+
+[yolo]
+mask = 0,1,2
+anchors = 10,13,  16,30,  33,23,  30,61,  62,45,  59,119,  116,90,  156,198,  373,326
+classes=601
+num=9
+jitter=.3
+ignore_thresh = .7
+truth_thresh = 1
+random=1
+
diff --git a/data/openimages.names b/data/openimages.names
new file mode 100644
index 00000000..ddfd8f22
--- /dev/null
+++ b/data/openimages.names
@@ -0,0 +1,601 @@
+Tortoise
+Container
+Magpie
+Sea turtle
+Football
+Ambulance
+Ladder
+Toothbrush
+Syringe
+Sink
+Toy
+Organ
+Cassette deck
+Apple
+Human eye
+Cosmetics
+Paddle
+Snowman
+Beer
+Chopsticks
+Human beard
+Bird
+Parking meter
+Traffic light
+Croissant
+Cucumber
+Radish
+Towel
+Doll
+Skull
+Washing machine
+Glove
+Tick
+Belt
+Sunglasses
+Banjo
+Cart
+Ball
+Backpack
+Bicycle
+Home appliance
+Centipede
+Boat
+Surfboard
+Boot
+Headphones
+Hot dog
+Shorts
+Fast food
+Bus
+Boy
+Screwdriver
+Bicycle wheel
+Barge
+Laptop
+Miniskirt
+Drill
+Dress
+Bear
+Waffle
+Pancake
+Brown bear
+Woodpecker
+Blue jay
+Pretzel
+Bagel
+Tower
+Teapot
+Person
+Bow and arrow
+Swimwear
+Beehive
+Brassiere
+Bee
+Bat
+Starfish
+Popcorn
+Burrito
+Chainsaw
+Balloon
+Wrench
+Tent
+Vehicle registration plate
+Lantern
+Toaster
+Flashlight
+Billboard
+Tiara
+Limousine
+Necklace
+Carnivore
+Scissors
+Stairs
+Computer keyboard
+Printer
+Traffic sign
+Chair
+Shirt
+Poster
+Cheese
+Sock
+Fire hydrant
+Land vehicle
+Earrings
+Tie
+Watercraft
+Cabinetry
+Suitcase
+Muffin
+Bidet
+Snack
+Snowmobile
+Clock
+Medical equipment
+Cattle
+Cello
+Jet ski
+Camel
+Coat
+Suit
+Desk
+Cat
+Bronze sculpture
+Juice
+Gondola
+Beetle
+Cannon
+Computer mouse
+Cookie
+Office building
+Fountain
+Coin
+Calculator
+Cocktail
+Computer monitor
+Box
+Stapler
+Christmas tree
+Cowboy hat
+Hiking equipment
+Studio couch
+Drum
+Dessert
+Wine rack
+Drink
+Zucchini
+Ladle
+Human mouth
+Dairy
+Dice
+Oven
+Dinosaur
+Ratchet
+Couch
+Cricket ball
+Winter melon
+Spatula
+Whiteboard
+Pencil sharpener
+Door
+Hat
+Shower
+Eraser
+Fedora
+Guacamole
+Dagger
+Scarf
+Dolphin
+Sombrero
+Tin can
+Mug
+Tap
+Harbor seal
+Stretcher
+Can opener
+Goggles
+Human body
+Roller skates
+Coffee cup
+Cutting board
+Blender
+Plumbing fixture
+Stop sign
+Office supplies
+Volleyball
+Vase
+Slow cooker
+Wardrobe
+Coffee
+Whisk
+Paper towel
+Personal care
+Food
+Sun hat
+Tree house
+Flying disc
+Skirt
+Gas stove
+Salt and pepper shakers
+Mechanical fan
+Face powder
+Fax
+Fruit
+French fries
+Nightstand
+Barrel
+Kite
+Tart
+Treadmill
+Fox
+Flag
+Horn
+Window blind
+Human foot
+Golf cart
+Jacket
+Egg
+Street light
+Guitar
+Pillow
+Human leg
+Isopod
+Grape
+Human ear
+Power plugs and sockets
+Panda
+Giraffe
+Woman
+Door handle
+Rhinoceros
+Bathtub
+Goldfish
+Houseplant
+Goat
+Baseball bat
+Baseball glove
+Mixing bowl
+Marine invertebrates
+Kitchen utensil
+Light switch
+House
+Horse
+Stationary bicycle
+Hammer
+Ceiling fan
+Sofa bed
+Adhesive tape
+Harp
+Sandal
+Bicycle helmet
+Saucer
+Harpsichord
+Human hair
+Heater
+Harmonica
+Hamster
+Curtain
+Bed
+Kettle
+Fireplace
+Scale
+Drinking straw
+Insect
+Hair dryer
+Kitchenware
+Indoor rower
+Invertebrate
+Food processor
+Bookcase
+Refrigerator
+Wood-burning stove
+Punching bag
+Common fig
+Cocktail shaker
+Jaguar
+Golf ball
+Fashion accessory
+Alarm clock
+Filing cabinet
+Artichoke
+Table
+Tableware
+Kangaroo
+Koala
+Knife
+Bottle
+Bottle opener
+Lynx
+Lavender
+Lighthouse
+Dumbbell
+Human head
+Bowl
+Humidifier
+Porch
+Lizard
+Billiard table
+Mammal
+Mouse
+Motorcycle
+Musical instrument
+Swim cap
+Frying pan
+Snowplow
+Bathroom cabinet
+Missile
+Bust
+Man
+Waffle iron
+Milk
+Ring binder
+Plate
+Mobile phone
+Baked goods
+Mushroom
+Crutch
+Pitcher
+Mirror
+Lifejacket
+Table tennis racket
+Pencil case
+Musical keyboard
+Scoreboard
+Briefcase
+Kitchen knife
+Nail
+Tennis ball
+Plastic bag
+Oboe
+Chest of drawers
+Ostrich
+Piano
+Girl
+Plant
+Potato
+Hair spray
+Sports equipment
+Pasta
+Penguin
+Pumpkin
+Pear
+Infant bed
+Polar bear
+Mixer
+Cupboard
+Jacuzzi
+Pizza
+Digital clock
+Pig
+Reptile
+Rifle
+Lipstick
+Skateboard
+Raven
+High heels
+Red panda
+Rose
+Rabbit
+Sculpture
+Saxophone
+Shotgun
+Seafood
+Submarine sandwich
+Snowboard
+Sword
+Picture frame
+Sushi
+Loveseat
+Ski
+Squirrel
+Tripod
+Stethoscope
+Submarine
+Scorpion
+Segway
+Training bench
+Snake
+Coffee table
+Skyscraper
+Sheep
+Television
+Trombone
+Tea
+Tank
+Taco
+Telephone
+Torch
+Tiger
+Strawberry
+Trumpet
+Tree
+Tomato
+Train
+Tool
+Picnic basket
+Cooking spray
+Trousers
+Bowling equipment
+Football helmet
+Truck
+Measuring cup
+Coffeemaker
+Violin
+Vehicle
+Handbag
+Paper cutter
+Wine
+Weapon
+Wheel
+Worm
+Wok
+Whale
+Zebra
+Auto part
+Jug
+Pizza cutter
+Cream
+Monkey
+Lion
+Bread
+Platter
+Chicken
+Eagle
+Helicopter
+Owl
+Duck
+Turtle
+Hippopotamus
+Crocodile
+Toilet
+Toilet paper
+Squid
+Clothing
+Footwear
+Lemon
+Spider
+Deer
+Frog
+Banana
+Rocket
+Wine glass
+Countertop
+Tablet computer
+Waste container
+Swimming pool
+Dog
+Book
+Elephant
+Shark
+Candle
+Leopard
+Axe
+Hand dryer
+Soap dispenser
+Porcupine
+Flower
+Canary
+Cheetah
+Palm tree
+Hamburger
+Maple
+Building
+Fish
+Lobster
+Asparagus
+Furniture
+Hedgehog
+Airplane
+Spoon
+Otter
+Bull
+Oyster
+Horizontal bar
+Convenience store
+Bomb
+Bench
+Ice cream
+Caterpillar
+Butterfly
+Parachute
+Orange
+Antelope
+Beaker
+Moths and butterflies
+Window
+Closet
+Castle
+Jellyfish
+Goose
+Mule
+Swan
+Peach
+Coconut
+Seat belt
+Raccoon
+Chisel
+Fork
+Lamp
+Camera
+Squash
+Racket
+Human face
+Human arm
+Vegetable
+Diaper
+Unicycle
+Falcon
+Chime
+Snail
+Shellfish
+Cabbage
+Carrot
+Mango
+Jeans
+Flowerpot
+Pineapple
+Drawer
+Stool
+Envelope
+Cake
+Dragonfly
+Sunflower
+Microwave oven
+Honeycomb
+Marine mammal
+Sea lion
+Ladybug
+Shelf
+Watch
+Candy
+Salad
+Parrot
+Handgun
+Sparrow
+Van
+Grinder
+Spice rack
+Light bulb
+Corded phone
+Sports uniform
+Tennis racket
+Wall clock
+Serving tray
+Kitchen & dining room table
+Dog bed
+Cake stand
+Cat furniture
+Bathroom accessory
+Facial tissue holder
+Pressure cooker
+Kitchen appliance
+Tire
+Ruler
+Luggage and bags
+Microphone
+Broccoli
+Umbrella
+Pastry
+Grapefruit
+Band-aid
+Animal
+Bell pepper
+Turkey
+Lily
+Pomegranate
+Doughnut
+Glasses
+Human nose
+Pen
+Ant
+Car
+Aircraft
+Human hand
+Skunk
+Teddy bear
+Watermelon
+Cantaloupe
+Dishwasher
+Flute
+Balance beam
+Sandwich
+Shrimp
+Sewing machine
+Binoculars
+Rays and skates
+Ipod
+Accordion
+Willow
+Crab
+Crown
+Seahorse
+Perfume
+Alpaca
+Taxi
+Canoe
+Remote control
+Wheelchair
+Rugby ball
+Armadillo
+Maracas
+Helmet
diff --git a/examples/darknet.c b/examples/darknet.c
index 4501b07d..b2fb79d2 100644
--- a/examples/darknet.c
+++ b/examples/darknet.c
@@ -475,8 +475,6 @@ int main(int argc, char **argv)
         composite_3d(argv[2], argv[3], argv[4], (argc > 5) ? atof(argv[5]) : 0);
     } else if (0 == strcmp(argv[1], "test")){
         test_resize(argv[2]);
-    } else if (0 == strcmp(argv[1], "captcha")){
-        run_captcha(argc, argv);
     } else if (0 == strcmp(argv[1], "nightmare")){
         run_nightmare(argc, argv);
     } else if (0 == strcmp(argv[1], "rgbgr")){
diff --git a/examples/detector.c b/examples/detector.c
index 502fe353..f78d67cc 100644
--- a/examples/detector.c
+++ b/examples/detector.c
@@ -168,6 +168,11 @@ void train_detector(char *datacfg, char *cfgfile, char *weightfile, int *gpus, i
     char buff[256];
     sprintf(buff, "%s/%s_final.weights", backup_directory, base);
     save_weights(net, buff);
+    free(paths);
+    free(plist);
+    free(base);
+    free(nets);
+    free(options);
 }
 
 
@@ -556,7 +561,7 @@ void validate_detector_recall(char *cfgfile, char *weightfile)
             ++total;
             box t = {truth[j].x, truth[j].y, truth[j].w, truth[j].h};
             float best_iou = 0;
-            for(k = 0; k < nboxes; ++k){
+            for(k = 0; k < l.w*l.h*l.n; ++k){
                 float iou = box_iou(dets[k].bbox, t);
                 if(dets[k].objectness > thresh && iou > best_iou){
                     best_iou = iou;
diff --git a/include/darknet.h b/include/darknet.h
index d01accbb..9091b533 100644
--- a/include/darknet.h
+++ b/include/darknet.h
@@ -46,7 +46,7 @@ typedef struct{
 tree *read_tree(char *filename);
 
 typedef enum{
-    LOGISTIC, RELU, RELIE, LINEAR, RAMP, TANH, PLSE, LEAKY, ELU, LOGGY, STAIR, HARDTAN, LHTAN
+    LOGISTIC, RELU, RELIE, LINEAR, RAMP, TANH, PLSE, LEAKY, ELU, LOGGY, STAIR, HARDTAN, LHTAN, SELU
 } ACTIVATION;
 
 typedef enum{
diff --git a/src/activation_kernels.cl b/src/activation_kernels.cl
index fee25bf0..1ddb7053 100644
--- a/src/activation_kernels.cl
+++ b/src/activation_kernels.cl
@@ -11,6 +11,7 @@ float logistic_activate_kernel(float x);
 float loggy_activate_kernel(float x);
 float relu_activate_kernel(float x);
 float elu_activate_kernel(float x);
+float selu_activate_kernel(float x);
 float relie_activate_kernel(float x);
 float ramp_activate_kernel(float x);
 float leaky_activate_kernel(float x);
@@ -23,6 +24,7 @@ float logistic_gradient_kernel(float x);
 float loggy_gradient_kernel(float x);
 float relu_gradient_kernel(float x);
 float elu_gradient_kernel(float x);
+float selu_gradient_kernel(float x);
 float relie_gradient_kernel(float x);
 float ramp_gradient_kernel(float x);
 float leaky_gradient_kernel(float x);
@@ -31,7 +33,7 @@ float plse_gradient_kernel(float x);
 float stair_gradient_kernel(float x);
 
 typedef enum{
-    LOGISTIC, RELU, RELIE, LINEAR, RAMP, TANH, PLSE, LEAKY, ELU, LOGGY, STAIR, HARDTAN, LHTAN
+    LOGISTIC, RELU, RELIE, LINEAR, RAMP, TANH, PLSE, LEAKY, ELU, LOGGY, STAIR, HARDTAN, LHTAN, SELU
 } ACTIVATION;
 
 float activate_kernel(float x, ACTIVATION a);
@@ -60,6 +62,7 @@ float logistic_activate_kernel(float x){return 1.f/(1.f + exp(-x));}
 float loggy_activate_kernel(float x){return 2.f/(1.f + exp(-x)) - 1;}
 float relu_activate_kernel(float x){return x*(x>0);}
 float elu_activate_kernel(float x){return (x >= 0)*x + (x < 0)*(exp(x)-1);}
+float selu_activate_kernel(float x){return (x >= 0)*1.0507f*x + (x < 0)*1.0507f*1.6732f*(exp(x)-1);}
 float relie_activate_kernel(float x){return (x>0) ? x : .01f*x;}
 float ramp_activate_kernel(float x){return x*(x>0)+.1f*x;}
 float leaky_activate_kernel(float x){return (x>0) ? x : .1f*x;}
@@ -92,6 +95,7 @@ float loggy_gradient_kernel(float x)
 }
 float relu_gradient_kernel(float x){return (x>0);}
 float elu_gradient_kernel(float x){return (x >= 0) + (x < 0)*(x + 1);}
+float selu_gradient_kernel(float x){return (x >= 0)*1.0507 + (x < 0)*(x + 1.0507*1.6732);}
 float relie_gradient_kernel(float x){return (x>0) ? 1 : .01f;}
 float ramp_gradient_kernel(float x){return (x>0)+.1f;}
 float leaky_gradient_kernel(float x){return (x>0) ? 1 : .1f;}
@@ -116,6 +120,8 @@ float activate_kernel(float x, ACTIVATION a)
             return relu_activate_kernel(x);
         case ELU:
             return elu_activate_kernel(x);
+        case SELU:
+            return selu_activate_kernel(x);
         case RELIE:
             return relie_activate_kernel(x);
         case RAMP:
@@ -149,6 +155,8 @@ float gradient_kernel(float x, ACTIVATION a)
             return relu_gradient_kernel(x);
         case ELU:
             return elu_gradient_kernel(x);
+        case SELU:
+            return selu_gradient_kernel(x);
         case RELIE:
             return relie_gradient_kernel(x);
         case RAMP:
diff --git a/src/activations.c b/src/activations.c
index 0cbb2f55..da1a17a8 100644
--- a/src/activations.c
+++ b/src/activations.c
@@ -16,6 +16,8 @@ char *get_activation_string(ACTIVATION a)
             return "relu";
         case ELU:
             return "elu";
+        case SELU:
+            return "selu";
         case RELIE:
             return "relie";
         case RAMP:
@@ -46,6 +48,7 @@ ACTIVATION get_activation(char *s)
     if (strcmp(s, "loggy")==0) return LOGGY;
     if (strcmp(s, "relu")==0) return RELU;
     if (strcmp(s, "elu")==0) return ELU;
+    if (strcmp(s, "selu")==0) return SELU;
     if (strcmp(s, "relie")==0) return RELIE;
     if (strcmp(s, "plse")==0) return PLSE;
     if (strcmp(s, "hardtan")==0) return HARDTAN;
@@ -72,6 +75,8 @@ float activate(float x, ACTIVATION a)
             return relu_activate(x);
         case ELU:
             return elu_activate(x);
+        case SELU:
+            return selu_activate(x);
         case RELIE:
             return relie_activate(x);
         case RAMP:
@@ -113,6 +118,8 @@ float gradient(float x, ACTIVATION a)
             return relu_gradient(x);
         case ELU:
             return elu_gradient(x);
+        case SELU:
+            return selu_gradient(x);
         case RELIE:
             return relie_gradient(x);
         case RAMP:
diff --git a/src/activations.h b/src/activations.h
index ff025682..65d85069 100644
--- a/src/activations.h
+++ b/src/activations.h
@@ -35,6 +35,7 @@ static inline float logistic_activate(float x){return 1./(1. + expf(-x));}
 static inline float loggy_activate(float x){return 2./(1. + expf(-x)) - 1;}
 static inline float relu_activate(float x){return x*(x>0);}
 static inline float elu_activate(float x){return (x >= 0)*x + (x < 0)*(expf(x)-1);}
+static inline float selu_activate(float x){return (x >= 0)*1.0507*x + (x < 0)*1.0507*1.6732*(expf(x)-1);}
 static inline float relie_activate(float x){return (x>0) ? x : .01*x;}
 static inline float ramp_activate(float x){return x*(x>0)+.1*x;}
 static inline float leaky_activate(float x){return (x>0) ? x : .1*x;}
@@ -77,6 +78,7 @@ static inline float stair_gradient(float x)
 }
 static inline float relu_gradient(float x){return (x>0);}
 static inline float elu_gradient(float x){return (x >= 0) + (x < 0)*(x + 1);}
+static inline float selu_gradient(float x){return (x >= 0)*1.0507 + (x < 0)*(x + 1.0507*1.6732);}
 static inline float relie_gradient(float x){return (x>0) ? 1 : .01;}
 static inline float ramp_gradient(float x){return (x>0)+.1;}
 static inline float leaky_gradient(float x){return (x>0) ? 1 : .1;}
diff --git a/src/blas_kernels.cl b/src/blas_kernels.cl
index 9c3dc51d..13e0079b 100644
--- a/src/blas_kernels.cl
+++ b/src/blas_kernels.cl
@@ -52,7 +52,7 @@ void scale_bias(float *output, float *biases, int batch, int n, int size)
 */
 __kernel void scale_bias_kernel(int N, __global float *output, __global float *biases, int batch, int n, int size)
 {
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
 
     if (id >= N) return;
 
@@ -82,7 +82,7 @@ void backward_scale_cpu(float *x_norm, float *delta, int batch, int n, int size,
 */
 __kernel void backward_scale_kernel(int N, __global float *x_norm, __global float *delta, int batch, int n, int size, __global float *scale_updates)
 {
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (id >= N) return;
 
     int f = (id / (batch*size));
@@ -123,7 +123,7 @@ void add_bias(float *output, float *biases, int batch, int n, int size)
 */
 __kernel void add_bias_kernel(int N, __global float *output, __global float *biases, int batch, int n, int size)
 {
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
 
     if (id >= N) return;
 
@@ -155,7 +155,7 @@ void backward_bias(float *bias_updates, float *delta, int batch, int n, int size
 */
 __kernel void backward_bias_kernel(int N, __global float *bias_updates, __global float *delta, int batch, int n, int size)
 {
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
 
     if (id >= N) return;
 
@@ -203,7 +203,7 @@ __kernel void  mean_kernel(int N, __global float *x, int batch, int filters, int
 {
     float scale = 1.f/(batch * spatial);
 
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (id >= N) return;
 
     int i = id;
@@ -240,7 +240,7 @@ __kernel void variance_kernel(int N, __global float *x, __global float *mean, in
 {
     float scale = 1.f/(batch * spatial - 1);
 
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (id >= N) return;
 
     int i = id;
@@ -327,7 +327,7 @@ __kernel void variance_delta_kernel(int N, __global float *x, __global float *de
 __kernel void accumulate_kernel(__global float *x, int n, int groups, __global float *sum)
 {
     int k;
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (i >= groups) return;
     sum[i] = 0;
     for(k = 0; k < n; ++k){
@@ -414,7 +414,7 @@ __kernel void fast_variance_delta_kernel(__const int threads, __global float *x,
 
 __kernel void adam_kernel(int N, __global float *x, __global float *m, __global float *v, float B1, float B2, float rate, float eps, int t)
 {
-    int index = get_global_id(0);
+    int index = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (index >= N) return;
 
     x[index] = x[index] + (rate * sqrt(1.f-pow(B2, t)) / (1.f-pow(B1, t)) * m[index] / (sqrt((v[index] + eps))));
@@ -437,7 +437,7 @@ void normalize_cpu(float *x, float *mean, float *variance, int batch, int filter
 */
 __kernel void normalize_kernel(int N, __global float *x, __global float *mean, __global float *variance, int batch, int filters, int spatial)
 {
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (id >= N) return;
 
     int b = (id / (filters*spatial));
@@ -465,7 +465,7 @@ void normalize_delta_cpu(float *x, float *mean, float *variance, float *mean_del
 */
 __kernel void normalize_delta_kernel(int N, __global float *x, __global float *mean, __global float *variance, __global float *mean_delta, __global float *variance_delta, int batch, int filters, int spatial, __global float *delta)
 {
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (id >= N) return;
 
     int j = (id / (filters*spatial));
@@ -500,7 +500,7 @@ void l2normalize_cpu(float *x, float *dx, int batch, int filters, int spatial)
 */
 __kernel void l2norm_kernel(int N, __global float *x, __global float *dx, int batch, int filters, int spatial)
 {
-    int index = get_global_id(0);
+    int index = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (index >= N) return;
     int b = index / spatial;
     int i = index % spatial;
@@ -522,7 +522,7 @@ __kernel void l2norm_kernel(int N, __global float *x, __global float *dx, int ba
 
 __kernel void reorg_kernel(int N, __global float *x, int w, int h, int c, int batch, int stride, int forward, __global float *out)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i >= N) return;
     int in_index = i;
     int in_w = i%w;
@@ -549,35 +549,35 @@ __kernel void reorg_kernel(int N, __global float *x, int w, int h, int c, int ba
 
 __kernel void axpy_kernel(int N, float ALPHA, __global float *X, int OFFX, int INCX,  __global float *Y, int OFFY, int INCY)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < N) Y[i*INCY+OFFY] += ALPHA*X[i*INCX+OFFX];
 }
 
 
 __kernel void pow_kernel(int N, float ALPHA, __global float *X, int OFFX, int INCX, __global float *Y, int OFFY, int INCY)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < N) Y[i*INCY + OFFY] = pow(X[i*INCX + OFFX], ALPHA);
 }
 
 
 __kernel void const_kernel(int N, float ALPHA, __global float *X, int OFFX, int INCX)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < N) X[i*INCX + OFFX] = ALPHA;
 }
 
 
 __kernel void constrain_kernel(int N, float ALPHA, __global float *X, int INCX)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < N) X[i*INCX] = min(ALPHA, max(-ALPHA, X[i*INCX]));
 }
 
 
 __kernel void supp_kernel(int N, float ALPHA, __global float *X, int INCX)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < N) {
         if((X[i*INCX] * X[i*INCX]) < (ALPHA * ALPHA)) X[i*INCX] = 0;
     }
@@ -586,49 +586,49 @@ __kernel void supp_kernel(int N, float ALPHA, __global float *X, int INCX)
 
 __kernel void add_kernel(int N, float ALPHA, __global float *X, int INCX)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < N) X[i*INCX] += ALPHA;
 }
 
 
 __kernel void scal_kernel(int N, float ALPHA, __global float *X, int INCX)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < N) X[i*INCX] *= ALPHA;
 }
 
 
 __kernel void fill_kernel(int N, float ALPHA, __global float *X, int OFFX, int INCX)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < N) X[i*INCX + OFFX] = ALPHA;
 }
 
 
 __kernel void mask_kernel(int n,  __global float *x, float mask_num, __global float *mask, float val)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n && mask[i] == mask_num) x[i] = val;
 }
 
 
 __kernel void copy_kernel(int N,  __global float *X, int OFFX, int INCX, __global float *Y, int OFFY, int INCY)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < N) Y[i*INCY + OFFY] = X[i*INCX + OFFX];
 }
 
 
 __kernel void mul_kernel(int N, __global float *X, int INCX, __global float *Y, int INCY)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < N) Y[i*INCY] *= X[i*INCX];
 }
 
 
 __kernel void flatten_kernel(int N, __global float *x, int spatial, int layers, int batch, int forward, __global float *out)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i >= N) return;
     int in_s = i%spatial;
     i = i/spatial;
@@ -646,7 +646,7 @@ __kernel void flatten_kernel(int N, __global float *x, int spatial, int layers,
 
 __kernel void shortcut_kernel(int size, int minw, int minh, int minc, int stride, int sample, int batch, int w1, int h1, int c1, __global float *add, int w2, int h2, int c2, float s1, float s2, __global float *out)
 {
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (id >= size) return;
     int i = id % minw;
     id /= minw;
@@ -665,7 +665,7 @@ __kernel void shortcut_kernel(int size, int minw, int minh, int minc, int stride
 
 __kernel void smooth_l1_kernel(int n, __global float *pred, __global float *truth, __global float *delta, __global float *error)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n){
         float diff = truth[i] - pred[i];
         float abs_val = fabs(diff);
@@ -683,7 +683,7 @@ __kernel void smooth_l1_kernel(int n, __global float *pred, __global float *trut
 
 __kernel void softmax_x_ent_kernel(int n, __global float *pred, __global float *truth, __global float *delta, __global float *error)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n) {
         float t = truth[i];
         float p = pred[i];
@@ -695,7 +695,7 @@ __kernel void softmax_x_ent_kernel(int n, __global float *pred, __global float *
 
 __kernel void logistic_x_ent_kernel(int n, __global float *pred, __global float *truth, __global float *delta, __global float *error)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n){
         float t = truth[i];
         float p = pred[i];
@@ -709,7 +709,7 @@ __kernel void logistic_x_ent_kernel(int n, __global float *pred, __global float
 
 __kernel void l2_kernel(int n, __global float *pred, __global float *truth, __global float *delta, __global float *error)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n){
         float t = truth[i];
         float p = pred[i];
@@ -722,7 +722,7 @@ __kernel void l2_kernel(int n, __global float *pred, __global float *truth, __gl
 
 __kernel void l1_kernel(int n, __global float *pred, __global float *truth, __global float *delta, __global float *error)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n){
         float diff = truth[i] - pred[i];
         error[i] = fabs(diff);
@@ -733,7 +733,7 @@ __kernel void l1_kernel(int n, __global float *pred, __global float *truth, __gl
 
 __kernel void wgan_kernel(int n, __global float *pred, __global float *truth, __global float *delta, __global float *error)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n){
         error[i] = (truth[i]!=0) ? -pred[i] : pred[i];
         delta[i] = (truth[i] > 0) ? 1 : -1;
@@ -743,7 +743,7 @@ __kernel void wgan_kernel(int n, __global float *pred, __global float *truth, __
 
 __kernel void weighted_sum_kernel(int n, __global float *a, __global float *b, __global float *s, __global float *c)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n){
         c[i] = s[i]*a[i] + (1-s[i])*(b ? b[i] : 0);
     }
@@ -752,7 +752,7 @@ __kernel void weighted_sum_kernel(int n, __global float *a, __global float *b, _
 
 __kernel void weighted_delta_kernel(int n, __global float *a, __global float *b, __global float *s, __global float *da, __global float *db, __global float *ds, __global float *dc)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n){
         if(da) da[i] += dc[i] * s[i];
         db[i] += dc[i] * (1-s[i]);
@@ -763,7 +763,7 @@ __kernel void weighted_delta_kernel(int n, __global float *a, __global float *b,
 
 __kernel void mult_add_into_kernel(int n, __global float *a, __global float *b, __global float *c)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n){
         c[i] += a[i]*b[i];
     }
@@ -772,7 +772,7 @@ __kernel void mult_add_into_kernel(int n, __global float *a, __global float *b,
 
 __kernel void deinter_kernel(int NX, __global float *X, int NY, __global float *Y, int B, __global float *OUT)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < (NX+NY)*B){
         int b = i / (NX+NY);
         int j = i % (NX+NY);
@@ -787,7 +787,7 @@ __kernel void deinter_kernel(int NX, __global float *X, int NY, __global float *
 
 __kernel void inter_kernel(int NX, __global float *X, int NY, __global float *Y, int B, __global float *OUT)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < (NX+NY)*B){
         int b = i / (NX+NY);
         int j = i % (NX+NY);
@@ -822,7 +822,7 @@ __kernel void softmax_device(__global float *input, int n, float temp, int strid
 
 __kernel void softmax_kernel(__global float *input, int offset, int n, int batch, int batch_offset, int groups, int group_offset, int stride, float temp, __global float *output)
 {
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (id >= batch*groups) return;
     int b = id / groups;
     int g = id % groups;
@@ -832,7 +832,7 @@ __kernel void softmax_kernel(__global float *input, int offset, int n, int batch
 
 __kernel void softmax_tree_kernel(__global float *input, int offset, int index, int spatial, int batch, int stride, float temp, __global float *output, int groups, __global float *group_size, __global float *group_offset)
 {
-    int id = get_global_id(0);
+    int id = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if (id >= spatial*batch*groups) return;
     int s = id % spatial;
     id = id / spatial;
@@ -846,14 +846,14 @@ __kernel void softmax_tree_kernel(__global float *input, int offset, int index,
 
 __kernel void scale_mask_kernel(int n, __global float *x, float mask_num, __global float *mask, float scale)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i < n && mask[i] == mask_num) x[i] *= scale;
 }
 
 
 __kernel void dot_kernel(__global float *output, float scale, int batch, int n, int size, __global float *delta)
 {
-    int index = get_global_id(0);
+    int index = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
 
     int f1 = index / n;
     int f2 = index % n;
@@ -887,6 +887,43 @@ __kernel void dot_kernel(__global float *output, float scale, int batch, int n,
 }
 
 //Source: https://www.sharcnet.ca/help/index.php/Porting_CUDA_to_OpenCL
+void atomic_add_global(volatile global float *source, const float operand);
+
+void atomic_add_global(volatile global float *source, const float operand) {
+    union {
+        unsigned int intVal;
+        float floatVal;
+    } newVal;
+    union {
+        unsigned int intVal;
+        float floatVal;
+    } prevVal;
+
+    do {
+        prevVal.floatVal = *source;
+        newVal.floatVal = prevVal.floatVal + operand;
+    } while (atomic_cmpxchg((volatile global unsigned int *)source, prevVal.intVal, newVal.intVal) != prevVal.intVal);
+}
+
+void atomic_add_local(volatile local float *source, const float operand);
+
+void atomic_add_local(volatile local float *source, const float operand) {
+    union {
+        unsigned int intVal;
+        float floatVal;
+    } newVal;
+
+    union {
+        unsigned int intVal;
+        float floatVal;
+    } prevVal;
+
+    do {
+        prevVal.floatVal = *source;
+        newVal.floatVal = prevVal.floatVal + operand;
+    } while (atomic_cmpxchg((volatile local unsigned int *)source, prevVal.intVal, newVal.intVal) != prevVal.intVal);
+}
+
 inline void atomicAdd_f(__global float* address, float value)
 {
     float old = value;
@@ -895,7 +932,7 @@ inline void atomicAdd_f(__global float* address, float value)
 
 __kernel void upsample_kernel(int N, __global float *x, int w, int h, int c, int batch, int stride, int forward, float scale, __global float *out)
 {
-    int i = get_global_id(0);
+    int i = (get_group_id(0) + get_group_id(1)*get_num_groups(0)) * get_local_size(0) + get_local_id(0);
     if(i >= N) return;
     int out_index = i;
     int out_w = i%(w*stride);
@@ -913,7 +950,7 @@ __kernel void upsample_kernel(int N, __global float *x, int w, int h, int c, int
     int in_index = b*w*h*c + in_c*w*h + in_h*w + in_w;
 
     if(forward) out[out_index] += scale * x[in_index];
-    else atomicAdd_f(x+in_index, scale * out[out_index]);
+    else atomic_add_global(x+in_index, scale * out[out_index]);
 }
 
 __kernel void gemm_kernel(
diff --git a/src/connected_layer.c b/src/connected_layer.c
index 849d6cfb..34f2fadc 100644
--- a/src/connected_layer.c
+++ b/src/connected_layer.c
@@ -44,7 +44,6 @@ layer make_connected_layer(int batch, int inputs, int outputs, ACTIVATION activa
 
     //float scale = 1./sqrt(inputs);
     float scale = sqrt(2./inputs);
-    //float scale = 1.0027f;
     for(i = 0; i < outputs*inputs; ++i){
         l.weights[i] = scale*rand_uniform(-1, 1);
     }
@@ -197,7 +196,7 @@ void denormalize_connected_layer(layer l)
 {
     int i, j;
     for(i = 0; i < l.outputs; ++i){
-        float scale = l.scales[i]/sqrt(l.rolling_variance[i]);
+        float scale = l.scales[i]/sqrt(l.rolling_variance[i] + .000001);
         for(j = 0; j < l.inputs; ++j){
             l.weights[i*l.inputs + j] *= scale;
         }
diff --git a/src/convolutional_kernels.c b/src/convolutional_kernels.c
index f4163a8f..20659216 100644
--- a/src/convolutional_kernels.c
+++ b/src/convolutional_kernels.c
@@ -285,4 +285,4 @@ void update_convolutional_layer_gpu(convolutional_layer l, update_args a)
     }
 }
 
-#endif // GPU
+#endif // GPU
\ No newline at end of file
diff --git a/src/convolutional_layer.c b/src/convolutional_layer.c
index 54f5151a..6d2b179f 100644
--- a/src/convolutional_layer.c
+++ b/src/convolutional_layer.c
@@ -33,7 +33,7 @@ void binarize_weights(float *weights, int n, int size, float *binary)
     for(f = 0; f < n; ++f){
         float mean = 0;
         for(i = 0; i < size; ++i){
-            mean += fabsf(weights[f*size + i]);
+            mean += fabs(weights[f*size + i]);
         }
         mean = mean / size;
         for(i = 0; i < size; ++i){
@@ -56,7 +56,7 @@ void binarize_input(float *input, int n, int size, float *binary)
     for(s = 0; s < size; ++s){
         float mean = 0;
         for(i = 0; i < n; ++i){
-            mean += fabsf(input[i*size + s]);
+            mean += fabs(input[i*size + s]);
         }
         mean = mean / n;
         for(i = 0; i < n; ++i){
@@ -86,8 +86,7 @@ image get_convolutional_delta(convolutional_layer l)
 }
 
 static size_t get_workspace_size(layer l){
-
-    return (size_t)l.out_h*l.out_w*l.size*l.size*l.c/l.groups;
+    return (size_t)l.out_h*l.out_w*l.size*l.size*l.c/l.groups*sizeof(float);
 }
 
 convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int n, int groups, int size, int stride, int padding, ACTIVATION activation, int batch_normalize, int binary, int xnor, int adam)
@@ -118,10 +117,10 @@ convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int
     l.nweights = c/groups*n*size*size;
     l.nbiases = n;
 
-    //float scale = 1.f/sqrt(size*size*c);
+    // float scale = 1./sqrt(size*size*c);
     float scale = sqrt(2./(size*size*c/l.groups));
-    //float scale = 1.0027f;
     //printf("convscale %f\n", scale);
+    //scale = .02;
     //for(i = 0; i < c*n*size*size; ++i) l.weights[i] = scale*rand_uniform(-1, 1);
     for(i = 0; i < l.nweights; ++i) l.weights[i] = scale*rand_normal();
     int out_w = convolutional_out_width(l);
@@ -152,7 +151,7 @@ convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int
         l.scales = calloc(n, sizeof(float));
         l.scale_updates = calloc(n, sizeof(float));
         for(i = 0; i < n; ++i){
-            l.scales[i] = 1.f;
+            l.scales[i] = 1;
         }
 
         l.mean = calloc(n, sizeof(float));
@@ -163,6 +162,7 @@ convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int
 
         l.rolling_mean = calloc(n, sizeof(float));
         l.rolling_variance = calloc(n, sizeof(float));
+
         l.x = calloc(l.batch*l.outputs, sizeof(float));
         l.x_norm = calloc(l.batch*l.outputs, sizeof(float));
     }
@@ -275,6 +275,29 @@ void test_convolutional_layer()
 
 void resize_convolutional_layer(convolutional_layer *l, int w, int h)
 {
+#ifdef GPU
+    if (gpu_index >= 0) {
+        if (l->delta_gpu.mem) opencl_free_gpu_only(l->delta_gpu);
+        if (l->output_gpu.mem) opencl_free_gpu_only(l->output_gpu);
+        if (l->batch_normalize) {
+            opencl_free_gpu_only(l->mean_gpu);
+            opencl_free_gpu_only(l->variance_gpu);
+            
+            opencl_free_gpu_only(l->rolling_mean_gpu);
+            opencl_free_gpu_only(l->rolling_variance_gpu);
+            
+            opencl_free_gpu_only(l->mean_delta_gpu);
+            opencl_free_gpu_only(l->variance_delta_gpu);
+
+            opencl_free_gpu_only(l->scales_gpu);
+            opencl_free_gpu_only(l->scale_updates_gpu);
+
+            opencl_free_gpu_only(l->x_gpu);
+            opencl_free_gpu_only(l->x_norm_gpu);
+        }
+    }
+#endif
+
     l->w = w;
     l->h = h;
     int out_w = convolutional_out_width(*l);
@@ -289,24 +312,41 @@ void resize_convolutional_layer(convolutional_layer *l, int w, int h)
     l->output = realloc(l->output, l->batch*l->outputs*sizeof(float));
     l->delta  = realloc(l->delta,  l->batch*l->outputs*sizeof(float));
     if(l->batch_normalize){
+        l->mean = realloc(l->mean, l->n*sizeof(float));
+        l->variance = realloc(l->variance, l->n*sizeof(float));
+
+        l->rolling_mean = realloc(l->rolling_mean, l->n*sizeof(float));
+        l->rolling_variance = realloc(l->rolling_variance, l->n*sizeof(float));
+
+        l->mean_delta = realloc(l->mean_delta, l->n*sizeof(float));
+        l->variance_delta = realloc(l->variance_delta, l->n*sizeof(float));
+
+        l->scales = realloc(l->scales, l->n*sizeof(float));
+        l->scale_updates = realloc(l->scale_updates, l->n*sizeof(float));
+
         l->x = realloc(l->x, l->batch*l->outputs*sizeof(float));
         l->x_norm  = realloc(l->x_norm, l->batch*l->outputs*sizeof(float));
     }
 
 #ifdef GPU
     if (gpu_index >= 0) {
-        if (l->delta_gpu.mem) opencl_free_gpu_only(l->delta_gpu);
-        if (l->output_gpu.mem) opencl_free_gpu_only(l->output_gpu);
-        if (l->batch_normalize) {
-            opencl_free_gpu_only(l->x_gpu);
-            opencl_free_gpu_only(l->x_norm_gpu);
-        }
-
         l->delta_gpu = opencl_make_array(l->delta, l->batch * l->outputs);
         l->output_gpu = opencl_make_array(l->output, l->batch * l->outputs);
         if (l->batch_normalize) {
-            l->x_gpu = opencl_make_array(l->output, l->batch * l->outputs);
-            l->x_norm_gpu = opencl_make_array(l->output, l->batch * l->outputs);
+            l->mean_gpu = opencl_make_array(l->mean, l->n);
+            l->variance_gpu = opencl_make_array(l->variance, l->n);
+
+            l->rolling_mean_gpu = opencl_make_array(l->rolling_mean, l->n);
+            l->rolling_variance_gpu = opencl_make_array(l->rolling_variance, l->n);
+
+            l->mean_delta_gpu = opencl_make_array(l->mean_delta, l->n);
+            l->variance_delta_gpu = opencl_make_array(l->variance_delta, l->n);
+
+            l->scales_gpu = opencl_make_array(l->scales, l->n);
+            l->scale_updates_gpu = opencl_make_array(l->scale_updates, l->n);
+
+            l->x_gpu = opencl_make_array(l->x, l->batch * l->outputs);
+            l->x_norm_gpu = opencl_make_array(l->x_norm, l->batch * l->outputs);
         }
     }
 #endif
@@ -325,13 +365,13 @@ void add_bias(float *output, float *biases, int batch, int n, int size)
     }
 }
 
-void scale_bias(float *output, float *biases, int batch, int n, int size)
+void scale_bias(float *output, float *scales, int batch, int n, int size)
 {
     int i,j,b;
     for(b = 0; b < batch; ++b){
         for(i = 0; i < n; ++i){
             for(j = 0; j < size; ++j){
-                output[(b*n + i)*size + j] *= biases[i];
+                output[(b*n + i)*size + j] *= scales[i];
             }
         }
     }
diff --git a/src/cost_layer.c b/src/cost_layer.c
index 660a05d5..d2481933 100644
--- a/src/cost_layer.c
+++ b/src/cost_layer.c
@@ -68,15 +68,19 @@ cost_layer make_cost_layer(int batch, int inputs, COST_TYPE cost_type, float sca
 
 void resize_cost_layer(cost_layer *l, int inputs)
 {
+#ifdef GPU
+    if (gpu_index >= 0) {
+        opencl_free_gpu_only(l->delta_gpu);
+        opencl_free_gpu_only(l->output_gpu);
+    }
+#endif
+
     l->inputs = inputs;
     l->outputs = inputs;
     l->delta = realloc(l->delta, inputs*l->batch*sizeof(float));
     l->output = realloc(l->output, inputs*l->batch*sizeof(float));
 #ifdef GPU
     if (gpu_index >= 0) {
-        opencl_free_gpu_only(l->delta_gpu);
-        opencl_free_gpu_only(l->output_gpu);
-
         l->delta_gpu = opencl_make_array(l->delta, inputs * l->batch);
         l->output_gpu = opencl_make_array(l->output, inputs * l->batch);
     }
diff --git a/src/crop_layer.c b/src/crop_layer.c
index 41836847..e22a6274 100644
--- a/src/crop_layer.c
+++ b/src/crop_layer.c
@@ -49,6 +49,12 @@ crop_layer make_crop_layer(int batch, int h, int w, int c, int crop_height, int
 
 void resize_crop_layer(layer *l, int w, int h)
 {
+#ifdef GPU
+    if (gpu_index >= 0) {
+        opencl_free_gpu_only(l->output_gpu);
+    }
+#endif
+
     l->w = w;
     l->h = h;
 
@@ -61,7 +67,6 @@ void resize_crop_layer(layer *l, int w, int h)
     l->output = realloc(l->output, l->batch*l->outputs*sizeof(float));
 #ifdef GPU
     if (gpu_index >= 0) {
-        opencl_free_gpu_only(l->output_gpu);
         l->output_gpu = opencl_make_array(l->output, l->outputs * l->batch);
     }
 #endif
diff --git a/src/data.c b/src/data.c
index 50acc516..03256ce3 100644
--- a/src/data.c
+++ b/src/data.c
@@ -430,7 +430,8 @@ void fill_truth_iseg(char *path, int num_boxes, float *truth, int classes, int w
     free_image(part);
 }
 
-void fill_truth_mask(char *path, int num_boxes, float *truth, int classes, int w, int h, augment_args aug, int flip, int mw, int mh){
+void fill_truth_mask(char *path, int num_boxes, float *truth, int classes, int w, int h, augment_args aug, int flip, int mw, int mh)
+{
     char labelpath[4096];
     find_replace(path, "images", "mask", labelpath);
     find_replace(labelpath, "JPEGImages", "mask", labelpath);
@@ -1084,8 +1085,7 @@ data load_data_detection(int n, char **paths, int m, int w, int h, int boxes, in
         float dh = jitter * orig.h;
 
         float new_ar = (orig.w + rand_uniform(-dw, dw)) / (orig.h + rand_uniform(-dh, dh));
-        //float scale = rand_uniform(.25, 2);
-        float scale = 1;
+        float scale = rand_uniform(.25, 2);
 
         float nw, nh;
 
@@ -1135,6 +1135,8 @@ void *load_thread(void *ptr)
         *a.d = load_data_super(a.paths, a.n, a.m, a.w, a.h, a.scale);
     } else if (a.type == WRITING_DATA){
         *a.d = load_data_writing(a.paths, a.n, a.m, a.w, a.h, a.out_w, a.out_h);
+    } else if (a.type == ISEG_DATA){
+        *a.d = load_data_iseg(a.n, a.paths, a.m, a.w, a.h, a.classes, a.num_boxes, a.scale, a.min, a.max, a.angle, a.aspect, a.hue, a.saturation, a.exposure);
     } else if (a.type == INSTANCE_DATA){
         *a.d = load_data_mask(a.n, a.paths, a.m, a.w, a.h, a.classes, a.num_boxes, a.coords, a.min, a.max, a.angle, a.aspect, a.hue, a.saturation, a.exposure);
     } else if (a.type == SEGMENTATION_DATA){
diff --git a/src/deconvolutional_layer.c b/src/deconvolutional_layer.c
index 52402697..788cc991 100644
--- a/src/deconvolutional_layer.c
+++ b/src/deconvolutional_layer.c
@@ -175,6 +175,13 @@ void denormalize_deconvolutional_layer(layer l)
 
 void resize_deconvolutional_layer(layer *l, int h, int w)
 {
+#ifdef GPU
+    if (gpu_index >= 0) {
+        opencl_free_gpu_only(l->delta_gpu);
+        opencl_free_gpu_only(l->output_gpu);
+    }
+#endif
+
     l->h = h;
     l->w = w;
     l->out_h = (l->h - 1) * l->stride + l->size - 2*l->pad;
@@ -192,9 +199,6 @@ void resize_deconvolutional_layer(layer *l, int h, int w)
 
 #ifdef GPU
     if (gpu_index >= 0) {
-        opencl_free_gpu_only(l->delta_gpu);
-        opencl_free_gpu_only(l->output_gpu);
-
         l->delta_gpu = opencl_make_array(l->delta, l->batch * l->outputs);
         l->output_gpu = opencl_make_array(l->output, l->batch * l->outputs);
 
@@ -202,8 +206,8 @@ void resize_deconvolutional_layer(layer *l, int h, int w)
             opencl_free_gpu_only(l->x_gpu);
             opencl_free_gpu_only(l->x_norm_gpu);
 
-            l->x_gpu = opencl_make_array(l->output, l->batch * l->outputs);
-            l->x_norm_gpu = opencl_make_array(l->output, l->batch * l->outputs);
+            l->x_gpu = opencl_make_array(l->x, l->batch * l->outputs);
+            l->x_norm_gpu = opencl_make_array(l->x_norm, l->batch * l->outputs);
         }
     }
 #endif
diff --git a/src/detection_layer.c b/src/detection_layer.c
index 845dc63c..90982b26 100644
--- a/src/detection_layer.c
+++ b/src/detection_layer.c
@@ -262,10 +262,10 @@ void forward_detection_layer_gpu(const detection_layer l, network net)
         return;
     }
 
-    opencl_pull_array(net.input_gpu, net.input, l.batch*l.inputs);
+    opencl_pull_array(net.output_gpu, net.input, l.batch*l.inputs);
     forward_detection_layer(l, net);
-
-    opencl_push_array(l.output_gpu, l.output, l.batch*l.outputs);
+    //opencl_push_array(l.output_gpu, l.output, l.batch*l.outputs);
+    if(!net.train) return;
     opencl_push_array(l.delta_gpu, l.delta, l.batch*l.inputs);
 }
 
diff --git a/src/dropout_layer.c b/src/dropout_layer.c
index 88a10fae..8e272594 100644
--- a/src/dropout_layer.c
+++ b/src/dropout_layer.c
@@ -29,10 +29,14 @@ dropout_layer make_dropout_layer(int batch, int inputs, float probability)
 
 void resize_dropout_layer(dropout_layer *l, int inputs)
 {
-    l->rand = realloc(l->rand, l->inputs*l->batch*sizeof(float));
 #ifdef GPU
     if (gpu_index >= 0) {
         opencl_free_gpu_only(l->rand_gpu);
+    }
+#endif
+    l->rand = realloc(l->rand, l->inputs*l->batch*sizeof(float));
+#ifdef GPU
+    if (gpu_index >= 0) {
         l->rand_gpu = opencl_make_array(l->rand, inputs * l->batch);
     }
 #endif
diff --git a/src/image_opencv.cpp b/src/image_opencv.cpp
new file mode 100644
index 00000000..5ba326d7
--- /dev/null
+++ b/src/image_opencv.cpp
@@ -0,0 +1,135 @@
+#ifdef OPENCV
+
+#include "stdio.h"
+#include "stdlib.h"
+#include "opencv2/opencv.hpp"
+#include "image.h"
+
+using namespace cv;
+
+extern "C" {
+
+IplImage *image_to_ipl(image im)
+{
+    int x,y,c;
+    IplImage *disp = cvCreateImage(cvSize(im.w,im.h), IPL_DEPTH_8U, im.c);
+    int step = disp->widthStep;
+    for(y = 0; y < im.h; ++y){
+        for(x = 0; x < im.w; ++x){
+            for(c= 0; c < im.c; ++c){
+                float val = im.data[c*im.h*im.w + y*im.w + x];
+                disp->imageData[y*step + x*im.c + c] = (unsigned char)(val*255);
+            }
+        }
+    }
+    return disp;
+}
+
+image ipl_to_image(IplImage* src)
+{
+    int h = src->height;
+    int w = src->width;
+    int c = src->nChannels;
+    image im = make_image(w, h, c);
+    unsigned char *data = (unsigned char *)src->imageData;
+    int step = src->widthStep;
+    int i, j, k;
+
+    for(i = 0; i < h; ++i){
+        for(k= 0; k < c; ++k){
+            for(j = 0; j < w; ++j){
+                im.data[k*w*h + i*w + j] = data[i*step + j*c + k]/255.;
+            }
+        }
+    }
+    return im;
+}
+
+Mat image_to_mat(image im)
+{
+    image copy = copy_image(im);
+    constrain_image(copy);
+    if(im.c == 3) rgbgr_image(copy);
+
+    IplImage *ipl = image_to_ipl(copy);
+    Mat m = cvarrToMat(ipl, true);
+    cvReleaseImage(&ipl);
+    free_image(copy);
+    return m;
+}
+
+image mat_to_image(Mat m)
+{
+    IplImage ipl = m;
+    image im = ipl_to_image(&ipl);
+    rgbgr_image(im);
+    return im;
+}
+
+void *open_video_stream(const char *f, int c, int w, int h, int fps)
+{
+    VideoCapture *cap;
+    if(f) cap = new VideoCapture(f);
+    else cap = new VideoCapture(c);
+    if(!cap->isOpened()) return 0;
+    if(w) cap->set(CV_CAP_PROP_FRAME_WIDTH, w);
+    if(h) cap->set(CV_CAP_PROP_FRAME_HEIGHT, w);
+    if(fps) cap->set(CV_CAP_PROP_FPS, w);
+    return (void *) cap;
+}
+
+image get_image_from_stream(void *p)
+{
+    VideoCapture *cap = (VideoCapture *)p;
+    Mat m;
+    *cap >> m;
+    if(m.empty()) return make_empty_image(0,0,0);
+    return mat_to_image(m);
+}
+
+image load_image_cv(char *filename, int channels)
+{
+    int flag = -1;
+    if (channels == 0) flag = -1;
+    else if (channels == 1) flag = 0;
+    else if (channels == 3) flag = 1;
+    else {
+        fprintf(stderr, "OpenCV can't force load with %d channels\n", channels);
+    }
+    Mat m;
+    m = imread(filename, flag);
+    if(!m.data){
+        fprintf(stderr, "Cannot load image \"%s\"\n", filename);
+        char buff[256];
+        sprintf(buff, "echo %s >> bad.list", filename);
+        system(buff);
+        return make_image(10,10,3);
+        //exit(0);
+    }
+    image im = mat_to_image(m);
+    return im;
+}
+
+int show_image_cv(image im, const char* name, int ms)
+{
+    Mat m = image_to_mat(im);
+    imshow(name, m);
+    int c = waitKey(ms);
+    if (c != -1) c = c%256;
+    return c;
+}
+
+void make_window(char *name, int w, int h, int fullscreen)
+{
+    namedWindow(name, WINDOW_NORMAL);
+    if (fullscreen) {
+        setWindowProperty(name, CV_WND_PROP_FULLSCREEN, CV_WINDOW_FULLSCREEN);
+    } else {
+        resizeWindow(name, w, h);
+        if(strcmp(name, "Demo") == 0) moveWindow(name, 0, 0);
+    }
+}
+
+}
+
+#endif
\ No newline at end of file
diff --git a/src/iseg_layer.c b/src/iseg_layer.c
index 3a6ce3f4..d6d5f27f 100644
--- a/src/iseg_layer.c
+++ b/src/iseg_layer.c
@@ -62,6 +62,13 @@ void resize_iseg_layer(layer *l, int w, int h)
     l->w = w;
     l->h = h;
 
+#ifdef GPU
+    if (gpu_index >= 0) {
+        opencl_free_gpu_only(l->delta_gpu);
+        opencl_free_gpu_only(l->output_gpu);
+    }
+#endif
+
     l->outputs = h*w*l->c;
     l->inputs = l->outputs;
 
@@ -70,11 +77,8 @@ void resize_iseg_layer(layer *l, int w, int h)
 
 #ifdef GPU
     if (gpu_index >= 0) {
-        opencl_free_gpu_only(l->delta_gpu);
-        opencl_free_gpu_only(l->output_gpu);
-
-        l->delta_gpu = opencl_make_array(l->delta, l->batch * l->outputs);
-        l->output_gpu = opencl_make_array(l->output, l->batch * l->outputs);
+        l->delta_gpu = opencl_make_array(l->delta, l->batch*l->outputs);
+        l->output_gpu = opencl_make_array(l->output, l->batch*l->outputs);
     }
 #endif
 }
@@ -113,11 +117,11 @@ void forward_iseg_layer(const layer l, network net)
             }
         }
 
+
         memset(l.counts, 0, 90*sizeof(int));
         for(i = 0; i < 90; ++i){
-            l.counts[i] = 0;
             fill_cpu(ids, 0, l.sums[i], 1);
-
+            
             int c = net.truth[b*l.truths + i*(l.w*l.h+1)];
             if(c < 0) break;
             // add up metric embeddings for each instance
@@ -217,9 +221,7 @@ void forward_iseg_layer_gpu(const layer l, network net)
         return;
     }
 
-    //opencl_pull_array(l.output_gpu, net.input, l.batch*l.inputs);
-    opencl_pull_array(l.output_gpu, l.output, l.batch*l.outputs);
-    memcpy(net.input, l.output, l.batch*l.outputs*sizeof(float));
+    opencl_pull_array(l.output_gpu, net.input, l.batch*l.inputs);
 
     forward_iseg_layer(l, net);
     opencl_push_array(l.delta_gpu, l.delta, l.batch*l.outputs);
diff --git a/src/layer.c b/src/layer.c
index 58874185..1400ba0f 100644
--- a/src/layer.c
+++ b/src/layer.c
@@ -6,7 +6,9 @@
 void free_layer(layer l)
 {
     if(l.type == DROPOUT){
-        if(l.rand)           free(l.rand);
+        if (gpu_index >= 0) {
+            if(l.rand)           free(l.rand);
+        }
 #ifdef GPU
         if (gpu_index >= 0) {
             if (l.rand_gpu.mem) opencl_free(l.rand_gpu);
@@ -14,47 +16,50 @@ void free_layer(layer l)
 #endif
         return;
     }
-    if(l.cweights)           free(l.cweights);
-    if(l.indexes)            free(l.indexes);
-    if(l.input_layers)       free(l.input_layers);
-    if(l.input_sizes)        free(l.input_sizes);
-    if(l.map)                free(l.map);
-    if(l.rand)               free(l.rand);
-    if(l.cost)               free(l.cost);
-    if(l.state)              free(l.state);
-    if(l.prev_state)         free(l.prev_state);
-    if(l.forgot_state)       free(l.forgot_state);
-    if(l.forgot_delta)       free(l.forgot_delta);
-    if(l.state_delta)        free(l.state_delta);
-    if(l.concat)             free(l.concat);
-    if(l.concat_delta)       free(l.concat_delta);
-    if(l.binary_weights)     free(l.binary_weights);
-    if(l.biases)             free(l.biases);
-    if(l.bias_updates)       free(l.bias_updates);
-    if(l.scales)             free(l.scales);
-    if(l.scale_updates)      free(l.scale_updates);
-    if(l.weights)            free(l.weights);
-    if(l.weight_updates)     free(l.weight_updates);
-    if(l.delta)              free(l.delta);
-    if(l.output)             free(l.output);
-    if(l.squared)            free(l.squared);
-    if(l.norms)              free(l.norms);
-    if(l.spatial_mean)       free(l.spatial_mean);
-    if(l.mean)               free(l.mean);
-    if(l.variance)           free(l.variance);
-    if(l.mean_delta)         free(l.mean_delta);
-    if(l.variance_delta)     free(l.variance_delta);
-    if(l.rolling_mean)       free(l.rolling_mean);
-    if(l.rolling_variance)   free(l.rolling_variance);
-    if(l.x)                  free(l.x);
-    if(l.x_norm)             free(l.x_norm);
-    if(l.m)                  free(l.m);
-    if(l.v)                  free(l.v);
-    if(l.z_cpu)              free(l.z_cpu);
-    if(l.r_cpu)              free(l.r_cpu);
-    if(l.h_cpu)              free(l.h_cpu);
-    if(l.binary_input)       free(l.binary_input);
-
+#ifndef GPU
+    if (gpu_index == -1) {
+        if(l.cweights)           free(l.cweights);
+        if(l.indexes)            free(l.indexes);
+        if(l.input_layers)       free(l.input_layers);
+        if(l.input_sizes)        free(l.input_sizes);
+        if(l.map)                free(l.map);
+        if(l.rand)               free(l.rand);
+        if(l.cost)               free(l.cost);
+        if(l.state)              free(l.state);
+        if(l.prev_state)         free(l.prev_state);
+        if(l.forgot_state)       free(l.forgot_state);
+        if(l.forgot_delta)       free(l.forgot_delta);
+        if(l.state_delta)        free(l.state_delta);
+        if(l.concat)             free(l.concat);
+        if(l.concat_delta)       free(l.concat_delta);
+        if(l.binary_weights)     free(l.binary_weights);
+        if(l.biases)             free(l.biases);
+        if(l.bias_updates)       free(l.bias_updates);
+        if(l.scales)             free(l.scales);
+        if(l.scale_updates)      free(l.scale_updates);
+        if(l.weights)            free(l.weights);
+        if(l.weight_updates)     free(l.weight_updates);
+        if(l.delta)              free(l.delta);
+        if(l.output)             free(l.output);
+        if(l.squared)            free(l.squared);
+        if(l.norms)              free(l.norms);
+        if(l.spatial_mean)       free(l.spatial_mean);
+        if(l.mean)               free(l.mean);
+        if(l.variance)           free(l.variance);
+        if(l.mean_delta)         free(l.mean_delta);
+        if(l.variance_delta)     free(l.variance_delta);
+        if(l.rolling_mean)       free(l.rolling_mean);
+        if(l.rolling_variance)   free(l.rolling_variance);
+        if(l.x)                  free(l.x);
+        if(l.x_norm)             free(l.x_norm);
+        if(l.m)                  free(l.m);
+        if(l.v)                  free(l.v);
+        if(l.z_cpu)              free(l.z_cpu);
+        if(l.r_cpu)              free(l.r_cpu);
+        if(l.h_cpu)              free(l.h_cpu);
+        if(l.binary_input)       free(l.binary_input);
+    }
+#endif
 #ifdef GPU
     if (gpu_index >= 0) {
         if (l.indexes_gpu.mem) opencl_free(l.indexes_gpu);
@@ -76,14 +81,16 @@ void free_layer(layer l)
         if (l.concat_delta_gpu.mem) opencl_free(l.concat_delta_gpu);
         if (l.binary_input_gpu.mem) opencl_free(l.binary_input_gpu);
         if (l.binary_weights_gpu.mem) opencl_free(l.binary_weights_gpu);
-        if (l.mean_gpu.mem) opencl_free(l.mean_gpu);
-        if (l.variance_gpu.mem) opencl_free(l.variance_gpu);
-        if (l.rolling_mean_gpu.mem) opencl_free(l.rolling_mean_gpu);
-        if (l.rolling_variance_gpu.mem) opencl_free(l.rolling_variance_gpu);
-        if (l.variance_delta_gpu.mem) opencl_free(l.variance_delta_gpu);
-        if (l.mean_delta_gpu.mem) opencl_free(l.mean_delta_gpu);
-        if (l.x_gpu.mem) opencl_free(l.x_gpu);
-        if (l.x_norm_gpu.mem) opencl_free(l.x_norm_gpu);
+        if (l.batch_normalize) {
+            if (l.mean_gpu.mem) opencl_free(l.mean_gpu);
+            if (l.variance_gpu.mem) opencl_free(l.variance_gpu);
+            if (l.rolling_mean_gpu.mem) opencl_free(l.rolling_mean_gpu);
+            if (l.rolling_variance_gpu.mem) opencl_free(l.rolling_variance_gpu);
+            if (l.variance_delta_gpu.mem) opencl_free(l.variance_delta_gpu);
+            if (l.mean_delta_gpu.mem) opencl_free(l.mean_delta_gpu);
+            if (l.x_gpu.mem) opencl_free(l.x_gpu);
+            if (l.x_norm_gpu.mem) opencl_free(l.x_norm_gpu);
+        }
         if (l.weights_gpu.mem) opencl_free(l.weights_gpu);
         if (l.weight_updates_gpu.mem) opencl_free(l.weight_updates_gpu);
         if (l.biases_gpu.mem) opencl_free(l.biases_gpu);
diff --git a/src/local_layer.c b/src/local_layer.c
index 7efc1641..8b927ed8 100644
--- a/src/local_layer.c
+++ b/src/local_layer.c
@@ -53,9 +53,8 @@ local_layer make_local_layer(int batch, int h, int w, int c, int n, int size, in
     l.biases = calloc(l.outputs, sizeof(float));
     l.bias_updates = calloc(l.outputs, sizeof(float));
 
-    //float scale = 1./sqrt(size*size*c);
+    // float scale = 1./sqrt(size*size*c);
     float scale = sqrt(2./(size*size*c));
-    //float scale = 1.0027f;
     for(i = 0; i < c*n*size*size; ++i) l.weights[i] = scale*rand_uniform(-1,1);
 
     l.output = calloc(l.batch*out_h * out_w * n, sizeof(float));
@@ -210,7 +209,7 @@ void forward_local_layer_gpu(const local_layer l, network net)
             int n = 1;
             int k = l.size*l.size*l.c;
             
-            gemm_offset_gpu(0,0,m,n,k,1,a,j*l.size*l.size*l.c,k,b,j,locations,1,c,j,locations);
+            gemm_offset_gpu(0,0,m,n,k,1,a,j*l.size*l.size*l.c*l.n,k,b,j,locations,1,c,i*l.outputs+j,locations);
         }
     }
     activate_array_gpu(l.output_gpu, l.outputs*l.batch, l.activation);
diff --git a/src/lstm_layer.c b/src/lstm_layer.c
index 7cff63cc..61964435 100644
--- a/src/lstm_layer.c
+++ b/src/lstm_layer.c
@@ -110,24 +110,24 @@ layer make_lstm_layer(int batch, int inputs, int outputs, int steps, int batch_n
         l.backward_gpu = backward_lstm_layer_gpu;
         l.update_gpu = update_lstm_layer_gpu;
 
-        l.output_gpu = opencl_make_array(l.output, batch * outputs * steps);
-        l.delta_gpu = opencl_make_array(l.delta, batch * l.outputs * steps);
-
-        l.prev_state_gpu = opencl_make_array(l.prev_state, batch * outputs);
-        l.prev_cell_gpu = opencl_make_array(l.prev_cell_cpu, batch * outputs);
-        l.cell_gpu = opencl_make_array(l.cell_cpu, batch * outputs * steps);
-
-        l.f_gpu = opencl_make_array(l.f_cpu, batch * outputs);
-        l.i_gpu = opencl_make_array(l.i_cpu, batch * outputs);
-        l.g_gpu = opencl_make_array(l.g_cpu, batch * outputs);
-        l.o_gpu = opencl_make_array(l.o_cpu, batch * outputs);
-        l.c_gpu = opencl_make_array(l.c_cpu, batch * outputs);
-        l.h_gpu = opencl_make_array(l.h_cpu, batch * outputs);
-        l.temp_gpu = opencl_make_array(l.temp_cpu, batch * outputs);
-        l.temp2_gpu = opencl_make_array(l.temp2_cpu, batch * outputs);
-        l.temp3_gpu = opencl_make_array(l.temp3_cpu, batch * outputs);
-        l.dc_gpu = opencl_make_array(l.dc_cpu, batch * outputs);
-        l.dh_gpu = opencl_make_array(l.dh_cpu, batch * outputs);
+        l.output_gpu = opencl_make_array(l.output, batch*outputs*steps);
+        l.delta_gpu = opencl_make_array(l.delta, batch*l.outputs*steps);
+
+        l.prev_state_gpu = opencl_make_array(l.prev_state, batch*outputs);
+        l.prev_cell_gpu = opencl_make_array(l.prev_cell_cpu, batch*outputs);
+        l.cell_gpu = opencl_make_array(l.cell_cpu, batch*outputs*steps);
+
+        l.f_gpu = opencl_make_array(l.f_cpu, batch*outputs);
+        l.i_gpu = opencl_make_array(l.i_cpu, batch*outputs);
+        l.g_gpu = opencl_make_array(l.g_cpu, batch*outputs);
+        l.o_gpu = opencl_make_array(l.o_cpu, batch*outputs);
+        l.c_gpu = opencl_make_array(l.c_cpu, batch*outputs);
+        l.h_gpu = opencl_make_array(l.h_cpu, batch*outputs);
+        l.temp_gpu = opencl_make_array(l.temp_cpu, batch*outputs);
+        l.temp2_gpu = opencl_make_array(l.temp2_cpu, batch*outputs);
+        l.temp3_gpu = opencl_make_array(l.temp3_cpu, batch*outputs);
+        l.dc_gpu = opencl_make_array(l.dc_cpu, batch*outputs);
+        l.dh_gpu = opencl_make_array(l.dh_cpu, batch*outputs);
     }
 #endif
 
diff --git a/src/maxpool_layer.c b/src/maxpool_layer.c
index c45c5dc2..9529b6e9 100644
--- a/src/maxpool_layer.c
+++ b/src/maxpool_layer.c
@@ -55,6 +55,14 @@ maxpool_layer make_maxpool_layer(int batch, int h, int w, int c, int size, int s
 
 void resize_maxpool_layer(maxpool_layer *l, int w, int h)
 {
+    #ifdef GPU
+    if (gpu_index >= 0) {
+        opencl_free_gpu_only(l->indexes_gpu);
+        opencl_free_gpu_only(l->output_gpu);
+        opencl_free_gpu_only(l->delta_gpu);
+    }
+    #endif
+
     l->h = h;
     l->w = w;
     l->inputs = h*w*l->c;
@@ -70,10 +78,6 @@ void resize_maxpool_layer(maxpool_layer *l, int w, int h)
 
     #ifdef GPU
     if (gpu_index >= 0) {
-        opencl_free_gpu_only(l->indexes_gpu);
-        opencl_free_gpu_only(l->output_gpu);
-        opencl_free_gpu_only(l->delta_gpu);
-
         l->indexes_gpu = opencl_make_int_array(l->indexes, output_size);
         l->output_gpu = opencl_make_array(l->output, output_size);
         l->delta_gpu = opencl_make_array(l->delta, output_size);
diff --git a/src/network.c b/src/network.c
index 87ee5739..8982d258 100644
--- a/src/network.c
+++ b/src/network.c
@@ -406,40 +406,35 @@ int resize_network(network *net, int w, int h)
     if(net->layers[net->n-1].truths) net->truths = net->layers[net->n-1].truths;
 #ifdef GPU
     if (gpu_index >= 0) {
-        //opencl_free_gpu_only(net->output_gpu);
         opencl_free_gpu_only(net->input_gpu);
         opencl_free_gpu_only(net->truth_gpu);
         opencl_free_gpu_only(net->delta_gpu);
     }
 #endif
-    //free(net->output);
     free(net->input);
     free(net->truth);
+    free(net->delta);
     net->output = out.output;
-    //net->output = calloc(net->outputs*net->batch, sizeof(float));
     net->input = calloc(net->inputs*net->batch, sizeof(float));
     net->truth = calloc(net->truths*net->batch, sizeof(float));
     net->delta = calloc(net->outputs*net->batch, sizeof(float));
 #ifdef GPU
-    if(gpu_index >= 0){
+    if (gpu_index >= 0) {
         net->output_gpu = out.output_gpu;
-        // net->output_gpu = opencl_make_array(net->output, net->outputs * net->batch);
         net->input_gpu = opencl_make_array(net->input, net->inputs * net->batch);
         net->truth_gpu = opencl_make_array(net->truth, net->truths * net->batch);
         net->delta_gpu = opencl_make_array(net->delta, net->outputs * net->batch);
-        if(workspace_size){
-            //printf("%ld\n", workspace_size*sizeof(float));
-            opencl_free(net->workspace_gpu);
-            net->workspace = calloc(workspace_size, sizeof(float));
-            net->workspace_gpu = opencl_make_array(net->workspace, workspace_size);
-        }
-    }else {
+        opencl_free(net->workspace_gpu);
+        net->workspace = calloc(1, workspace_size);
+        net->workspace_gpu = opencl_make_array(net->workspace, (workspace_size-1)/sizeof(float)+1);
+    }
+    else {
         free(net->workspace);
-        net->workspace = calloc(workspace_size, sizeof(float));
+        net->workspace = calloc(1, workspace_size);
     }
 #else
     free(net->workspace);
-    net->workspace = calloc(workspace_size, sizeof(float));
+    net->workspace = calloc(1, workspace_size);
 #endif
     //fprintf(stderr, " Done!\n");
     return 0;
@@ -463,7 +458,7 @@ image get_network_image_layer(network *net, int i)
     layer l = net->layers[i];
 #ifdef GPU
     if (gpu_index >= 0) {
-        //opencl_pull_array(l.output_gpu, l.output, l.outputs);
+        opencl_pull_array(l.output_gpu, l.output, l.outputs);
     }
 #endif
     if (l.out_w && l.out_h && l.out_c){
@@ -507,13 +502,7 @@ void top_predictions(network *net, int k, int *index)
 float *network_predict(network *net, float *input)
 {
     network orig = *net;
-    //net->input = input;
-#ifdef GPU
-    if (gpu_index >= 0) {
-        assert(net->input_gpu.ptr == (void *) net->input);
-    }
-#endif
-    memcpy(net->input, input, net->inputs*net->batch*sizeof(float));
+    net->input = input;
     net->truth = 0;
     net->train = 0;
     net->delta = 0;
@@ -594,11 +583,10 @@ void free_detections(detection *dets, int n)
 
 float *network_predict_image(network *net, image im)
 {
-    int resize = im.w != net->w || im.h != net->h;
-    image imr = resize ? letterbox_image(im, net->w, net->h) : im;
+    image imr = letterbox_image(im, net->w, net->h);
     set_batch_network(net, 1);
     float *p = network_predict(net, imr.data);
-    if(resize) free_image(imr);
+    free_image(imr);
     return p;
 }
 
@@ -737,12 +725,18 @@ void free_network(network *net)
         free_layer(net->layers[i]);
     }
     free(net->layers);
-    if(net->input) free(net->input);
-    if(net->truth) free(net->truth);
+#ifndef GPU
+if (gpu_index == -1) {
+    if (net->input) free(net->input);
+    if (net->truth) free(net->truth);
+    if (net->delta) free(net->delta);
+}
+#endif
 #ifdef GPU
     if (gpu_index >= 0) {
         if (net->input_gpu.mem) opencl_free(net->input_gpu);
         if (net->truth_gpu.mem) opencl_free(net->truth_gpu);
+        if (net->delta_gpu.mem) opencl_free(net->delta_gpu);
         if (net->workspace_gpu.mem) opencl_free(net->workspace_gpu);
     }
     else {
@@ -818,7 +812,7 @@ void forward_network_gpu(network *netp)
             net.truth = l.output;
         }
     }
-    if (!netp->train) pull_network_output(netp);
+    pull_network_output(netp);
     calc_network_cost(netp);
 }
 
@@ -1149,6 +1143,7 @@ void sync_nets(network **nets, int n, int interval)
 
 float train_networks(network **nets, int n, data d, int interval, int* gpus, int ngpus)
 {
+    assert(gpus);
     int i;
     int batch = nets[0]->batch;
     int subdivisions = nets[0]->subdivisions;
diff --git a/src/normalization_layer.c b/src/normalization_layer.c
index c47e8668..338fe030 100644
--- a/src/normalization_layer.c
+++ b/src/normalization_layer.c
@@ -41,7 +41,12 @@ layer make_normalization_layer(int batch, int w, int h, int c, int size, float a
 void resize_normalization_layer(layer *layer, int w, int h)
 {
 #ifdef GPU
-
+    if (gpu_index >= 0) {
+        opencl_free_gpu_only(layer->output_gpu);
+        opencl_free_gpu_only(layer->delta_gpu);
+        opencl_free_gpu_only(layer->squared_gpu);
+        opencl_free_gpu_only(layer->norms_gpu);
+    }
 #endif
     int c = layer->c;
     int batch = layer->batch;
@@ -57,11 +62,6 @@ void resize_normalization_layer(layer *layer, int w, int h)
     layer->norms = realloc(layer->norms, h * w * c * batch * sizeof(float));
 #ifdef GPU
     if (gpu_index >= 0) {
-        opencl_free_gpu_only(layer->output_gpu);
-        opencl_free_gpu_only(layer->delta_gpu);
-        opencl_free_gpu_only(layer->squared_gpu);
-        opencl_free_gpu_only(layer->norms_gpu);
-
         layer->output_gpu = opencl_make_array(layer->output, h * w * c * batch);
         layer->delta_gpu = opencl_make_array(layer->delta, h * w * c * batch);
         layer->squared_gpu = opencl_make_array(layer->squared, h * w * c * batch);
diff --git a/src/opencl.c b/src/opencl.c
index 88eac8c2..d936b8b2 100644
--- a/src/opencl.c
+++ b/src/opencl.c
@@ -424,6 +424,7 @@ char* concat(const char *s1, const char *s2)
 
 void opencl_load_buffer(const char *buffer, const size_t size, cl_program *output)
 {
+#ifdef RPI
     cl_int clErr;
 
     *output = clCreateProgramWithSource(opencl_context, CL_TRUE,
@@ -434,20 +435,13 @@ void opencl_load_buffer(const char *buffer, const size_t size, cl_program *outpu
         printf("opencl_load: could not create program. error: %s\n", clCheckError(clErr));
         return;
     }
-#ifdef RPI
+
     clErr = clBuildProgram(
             *output,
             1,
             &opencl_devices[opencl_device_id_t],
             NULL, NULL, NULL);
-#else
-    clErr = clBuildProgram(
-            *output,
-            1,
-            &opencl_devices[opencl_device_id_t],
-            "-cl-fp32-correctly-rounded-divide-sqrt "
-            "-Werror ", NULL, NULL);
-#endif
+
     if (clErr != CL_SUCCESS)
     {
         printf("opencl_load: could not compile. error: %s\n", clCheckError(clErr));
@@ -459,6 +453,89 @@ void opencl_load_buffer(const char *buffer, const size_t size, cl_program *outpu
         free(ebuffer);
         exit(1);
     }
+#else
+    cl_int clErr;
+
+    cl_program prhd = 0;
+
+    const char * code_header = "if CONFIG_USE_DOUBLE\n"
+                               "if defined(cl_khr_fp64)  // Khronos extension available?\n"
+                               "pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
+                               "define DOUBLE_SUPPORT_AVAILABLE\n"
+                               "elif defined(cl_amd_fp64)  // AMD extension available?\n"
+                               "pragma OPENCL EXTENSION cl_amd_fp64 : enable\n"
+                               "define DOUBLE_SUPPORT_AVAILABLE\n"
+                               "endif\n"
+                               "endif // CONFIG_USE_DOUBLE\n"
+                               "if defined(DOUBLE_SUPPORT_AVAILABLE)\n"
+                               "//#define float double\n"
+                               "else\n"
+                               "endif\n";
+
+    size_t headsz = strlen(code_header);
+
+    prhd = clCreateProgramWithSource(opencl_context, CL_TRUE,
+                                     (const char**)&code_header, &headsz, &clErr);
+
+    if (clErr != CL_SUCCESS)
+    {
+        printf("opencl_load: could not create header. error: %s\n", clCheckError(clErr));
+        return;
+    }
+
+    cl_program prog[1];
+
+    prog[0] = clCreateProgramWithSource(opencl_context, CL_TRUE,
+                                        (const char**)&buffer, &size, &clErr);
+
+    if (clErr != CL_SUCCESS)
+    {
+        printf("opencl_load: could not create program. error: %s\n", clCheckError(clErr));
+        return;
+    }
+
+    cl_program input_headers[1] = { prhd };
+    const char * input_header_names[1] = { "prhd" };
+
+    clErr = clCompileProgram(
+            prog[0],
+            opencl_device_ct_t,
+            opencl_devices,
+            "-cl-denorms-are-zero "
+            "-cl-fp32-correctly-rounded-divide-sqrt "
+            "-cl-std=CL1.2 "
+            "-Werror "
+            "-DCONFIG_USE_DOUBLE "
+            //"-cl-opt-disable "
+            ,
+            1, input_headers, input_header_names,
+            NULL, NULL);
+
+    if (clErr != CL_SUCCESS)
+    {
+        printf("opencl_load: could not compile. error: %s\n", clCheckError(clErr));
+        size_t len;
+        char *buffer = (char*)calloc(0x10000000, sizeof(char));
+        clGetProgramBuildInfo(prog[0], opencl_devices[opencl_device_id_t], CL_PROGRAM_BUILD_LOG, 0x10000000 * sizeof(char), buffer, &len);
+        printf("CL_PROGRAM_BUILD_LOG:\n%s\n", buffer);
+        free(buffer);
+    }
+
+    *output =
+            clLinkProgram(opencl_context, opencl_device_ct_t, opencl_devices,
+                          "-cl-denorms-are-zero "
+                          "-cl-kernel-arg-info", 1, prog, NULL, NULL, &clErr);
+
+    if (clErr != CL_SUCCESS)
+    {
+        printf("opencl_load: could not link. error: %s\n", clCheckError(clErr));
+        size_t len;
+        char *buffer = (char*)calloc(0x10000000, sizeof(char));
+        clGetProgramBuildInfo(*output, opencl_devices[opencl_device_id_t], CL_PROGRAM_BUILD_LOG, 0x10000000 * sizeof(char), buffer, &len);
+        printf("CL_PROGRAM_BUILD_LOG:\n%s\n", buffer);
+        free(buffer);
+    }
+#endif
 }
 
 void opencl_create_kernel(cl_program *program, const char *kernelName,
@@ -791,7 +868,7 @@ cl_mem_ext opencl_make_array(float *x, size_t n)
     buf.org = clCreateBuffer(opencl_context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR,
                              buf.len * buf.obs, buf.ptr, &clErr);
     if (clErr != CL_SUCCESS)
-        printf("could not push array to device. error: %s\n", clCheckError(clErr));
+        printf("could create buffer on device. error: %s\n", clCheckError(clErr));
 
     buf.mem = buf.org;
 
@@ -825,7 +902,7 @@ cl_mem_ext opencl_make_int_array(int *x, size_t n)
     buf.org = clCreateBuffer(opencl_context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR,
                              buf.len * buf.obs, buf.ptr, &clErr);
     if (clErr != CL_SUCCESS)
-        printf("could not push array to device. error: %s\n", clCheckError(clErr));
+        printf("could create buffer on device. error: %s\n", clCheckError(clErr));
 
     buf.mem = buf.org;
 
@@ -843,9 +920,9 @@ cl_mem_ext opencl_make_int_array(int *x, size_t n)
 
 void opencl_push_int_array(cl_mem_ext x_gpu, int *x, size_t n)
 {
-    if (ngpusg == 1) assert(x_gpu.ptr == (void*) x);
+    //if (ngpusg == 1) assert(x_gpu.ptr == (void*) x);
 
-    if (ngpusg == 1) assert(n != 0 && x_gpu.len - x_gpu.off == n);
+    //if (ngpusg == 1) assert(n != 0 && x_gpu.len - x_gpu.off == n);
 
     //x_gpu.ptr = x;
 
@@ -856,28 +933,28 @@ void opencl_push_int_array(cl_mem_ext x_gpu, int *x, size_t n)
     //cl_int clErr = clEnqueueWriteBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, 0,
     //                                    (n - x_gpu.off) * x_gpu.obs, x_gpu.ptr, 0, NULL, NULL);
     cl_int clErr;
-    void* map = clEnqueueMapBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, CL_MAP_WRITE,
+    int* map = clEnqueueMapBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE,
                                    0, x_gpu.len * x_gpu.obs, 0, NULL, NULL, &clErr);
     if (clErr != CL_SUCCESS)
         printf("could not map array to device. error: %s\n", clCheckError(clErr));
     memcpy(map, x, (n - x_gpu.off) * x_gpu.obs);
-    clErr = clEnqueueUnmapMemObject(opencl_queues[opencl_device_id_t], x_gpu.mem, map, 0, NULL, NULL);
-    if (clErr != CL_SUCCESS)
-        printf("could not unmap array from device. error: %s\n", clCheckError(clErr));
+    clErr = clEnqueueUnmapMemObject(opencl_queues[opencl_device_id_t], x_gpu.mem, (void*)map, 0, NULL, NULL);
 #ifdef BENCHMARK
     t = clock() - t;
     double time_taken = ((double)t);
     printf("%s\t%x\t%d\n", "opencl_push_int_array", x_gpu.ptr, (int)time_taken);
 #endif
+    if (clErr != CL_SUCCESS)
+        printf("could not unmap array from device. error: %s\n", clCheckError(clErr));
     //if (clErr != CL_SUCCESS)
     //    printf("could not push array to device. error: %s\n", clCheckError(clErr));
 }
 
 void opencl_pull_int_array(cl_mem_ext x_gpu, int *x, size_t n)
 {
-    if (ngpusg == 1) assert(x_gpu.ptr == (void*) x);
+    //if (ngpusg == 1) assert(x_gpu.ptr == (void*) x);
 
-    if (ngpusg == 1) assert(n != 0 && x_gpu.len - x_gpu.off == n);
+    //if (ngpusg == 1) assert(n != 0 && x_gpu.len - x_gpu.off == n);
 
     //x_gpu.ptr = x;
 
@@ -888,28 +965,28 @@ void opencl_pull_int_array(cl_mem_ext x_gpu, int *x, size_t n)
     //cl_int clErr = clEnqueueReadBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, 0,
     //                                   (n - x_gpu.off) * x_gpu.obs, x_gpu.ptr, 0, NULL, NULL);
     cl_int clErr;
-    void* map = clEnqueueMapBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, CL_MAP_READ,
+    int* map = clEnqueueMapBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE,
                                    0, x_gpu.len * x_gpu.obs, 0, NULL, NULL, &clErr);
     if (clErr != CL_SUCCESS)
         printf("could not map array to device. error: %s\n", clCheckError(clErr));
     memcpy(x, map, (n - x_gpu.off) * x_gpu.obs);
-    clErr = clEnqueueUnmapMemObject(opencl_queues[opencl_device_id_t], x_gpu.mem, map, 0, NULL, NULL);
-    if (clErr != CL_SUCCESS)
-        printf("could not unmap array from device. error: %s\n", clCheckError(clErr));
+    clErr = clEnqueueUnmapMemObject(opencl_queues[opencl_device_id_t], x_gpu.mem, (void*)map, 0, NULL, NULL);
 #ifdef BENCHMARK
     t = clock() - t;
     double time_taken = ((double)t);
     printf("%s\t%x\t%d\n", "opencl_pull_int_array", x_gpu.ptr, (int)time_taken);
 #endif
+    if (clErr != CL_SUCCESS)
+        printf("could not unmap array from device. error: %s\n", clCheckError(clErr));
     //if (clErr != CL_SUCCESS)
     //    printf("could not pull array to device. error: %s\n", clCheckError(clErr));
 }
 
 void opencl_push_array(cl_mem_ext x_gpu, float *x, size_t n)
 {
-    if (ngpusg == 1) assert(x_gpu.ptr == (void*) x);
+    //if (ngpusg == 1) assert(x_gpu.ptr == (void*) x);
 
-    if (ngpusg == 1) assert(n != 0 && x_gpu.len  - x_gpu.off == n);
+    //if (ngpusg == 1) assert(n != 0 && x_gpu.len  - x_gpu.off == n);
 
     //x_gpu.ptr = x;
 
@@ -920,28 +997,28 @@ void opencl_push_array(cl_mem_ext x_gpu, float *x, size_t n)
     //cl_int clErr = clEnqueueWriteBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, 0,
     //                                    (n - x_gpu.off) * x_gpu.obs, x_gpu.ptr, 0, NULL, NULL);
     cl_int clErr;
-    void* map = clEnqueueMapBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, CL_MAP_WRITE,
+    float* map = clEnqueueMapBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE,
                                    0, x_gpu.len * x_gpu.obs, 0, NULL, NULL, &clErr);
     if (clErr != CL_SUCCESS)
         printf("could not map array to device. error: %s\n", clCheckError(clErr));
     memcpy(map, x, (n - x_gpu.off) * x_gpu.obs);
-    clErr = clEnqueueUnmapMemObject(opencl_queues[opencl_device_id_t], x_gpu.mem, map, 0, NULL, NULL);
-    if (clErr != CL_SUCCESS)
-        printf("could not unmap array from device. error: %s\n", clCheckError(clErr));
+    clErr = clEnqueueUnmapMemObject(opencl_queues[opencl_device_id_t], x_gpu.mem, (void*)map, 0, NULL, NULL);
 #ifdef BENCHMARK
     t = clock() - t;
     double time_taken = ((double)t);
     printf("%s\t%x\t%d\n", "opencl_push_array", x_gpu.ptr, (int)time_taken);
 #endif
+    if (clErr != CL_SUCCESS)
+        printf("could not unmap array from device. error: %s\n", clCheckError(clErr));
     //if (clErr != CL_SUCCESS)
     //    printf("could not push array to device. error: %s\n", clCheckError(clErr));
 }
 
 void opencl_pull_array(cl_mem_ext x_gpu, float *x, size_t n)
 {
-    if (ngpusg == 1) assert(x_gpu.ptr == (void*) x);
+    //if (ngpusg == 1) assert(x_gpu.ptr == (void*) x);
 
-    if (ngpusg == 1) assert(n != 0 && x_gpu.len  - x_gpu.off == n);
+    //if (ngpusg == 1) assert(n != 0 && x_gpu.len  - x_gpu.off == n);
 
     //x_gpu.ptr = x;
 
@@ -952,19 +1029,19 @@ void opencl_pull_array(cl_mem_ext x_gpu, float *x, size_t n)
     //cl_int clErr = clEnqueueReadBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, 0,
     //                                   (n - x_gpu.off) * x_gpu.obs, x_gpu.ptr, 0, NULL, NULL);
     cl_int clErr;
-    void* map = clEnqueueMapBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, CL_MAP_READ,
+    float* map = clEnqueueMapBuffer(opencl_queues[opencl_device_id_t], x_gpu.mem, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE,
                                    0, x_gpu.len * x_gpu.obs, 0, NULL, NULL, &clErr);
     if (clErr != CL_SUCCESS)
         printf("could not map array to device. error: %s\n", clCheckError(clErr));
     memcpy(x, map, (n - x_gpu.off) * x_gpu.obs);
-    clErr = clEnqueueUnmapMemObject(opencl_queues[opencl_device_id_t], x_gpu.mem, map, 0, NULL, NULL);
-    if (clErr != CL_SUCCESS)
-        printf("could not unmap array from device. error: %s\n", clCheckError(clErr));
+    clErr = clEnqueueUnmapMemObject(opencl_queues[opencl_device_id_t], x_gpu.mem, (void*)map, 0, NULL, NULL);
 #ifdef BENCHMARK
     t = clock() - t;
     double time_taken = ((double)t);
     printf("%s\t%x\t%d\n", "opencl_pull_array", x_gpu.ptr, (int)time_taken);
 #endif
+    if (clErr != CL_SUCCESS)
+        printf("could not unmap array from device. error: %s\n", clCheckError(clErr));
     //if (clErr != CL_SUCCESS)
     //    printf("could not pull array to device. error: %s\n", clCheckError(clErr));
 }
@@ -980,9 +1057,12 @@ void opencl_free(cl_mem_ext x_gpu)
     x_gpu.dec = 0;
     x_gpu.add = 0;
     x_gpu.rem = 0;
-    if (!x_gpu.ptr)
-        free(x_gpu.ptr);
     clReleaseMemObject(x_gpu.org);
+    x_gpu.org = 0;
+    if (x_gpu.ptr) {
+        free(x_gpu.ptr);
+        x_gpu.ptr = 0;
+    }
 }
 
 void opencl_free_gpu_only(cl_mem_ext x_gpu){
@@ -996,6 +1076,7 @@ void opencl_free_gpu_only(cl_mem_ext x_gpu){
     x_gpu.add = 0;
     x_gpu.rem = 0;
     clReleaseMemObject(x_gpu.org);
+    x_gpu.org = 0;
 }
 
 cl_mem_ext inc(cl_mem_ext buf, int inc, size_t len) {
diff --git a/src/opencl.h b/src/opencl.h
index 59bb5487..471b7650 100644
--- a/src/opencl.h
+++ b/src/opencl.h
@@ -9,6 +9,7 @@
 
 #ifdef GPU
 #ifdef __APPLE__
+#define CL_SILENCE_DEPRECATION
 #include <OpenCL/cl.h>
 #include <OpenCL/cl_ext.h>
 #else
diff --git a/src/parser.c b/src/parser.c
index 474cae6f..027ec881 100644
--- a/src/parser.c
+++ b/src/parser.c
@@ -56,31 +56,31 @@ LAYER_TYPE string_to_layer_type(char * type)
     if (strcmp(type, "[iseg]")==0) return ISEG;
     if (strcmp(type, "[local]")==0) return LOCAL;
     if (strcmp(type, "[conv]")==0
-        || strcmp(type, "[convolutional]")==0) return CONVOLUTIONAL;
+            || strcmp(type, "[convolutional]")==0) return CONVOLUTIONAL;
     if (strcmp(type, "[deconv]")==0
-        || strcmp(type, "[deconvolutional]")==0) return DECONVOLUTIONAL;
+            || strcmp(type, "[deconvolutional]")==0) return DECONVOLUTIONAL;
     if (strcmp(type, "[activation]")==0) return ACTIVE;
     if (strcmp(type, "[logistic]")==0) return LOGXENT;
     if (strcmp(type, "[l2norm]")==0) return L2NORM;
     if (strcmp(type, "[net]")==0
-        || strcmp(type, "[network]")==0) return NETWORK;
+            || strcmp(type, "[network]")==0) return NETWORK;
     if (strcmp(type, "[crnn]")==0) return CRNN;
     if (strcmp(type, "[gru]")==0) return GRU;
     if (strcmp(type, "[lstm]") == 0) return LSTM;
     if (strcmp(type, "[rnn]")==0) return RNN;
     if (strcmp(type, "[conn]")==0
-        || strcmp(type, "[connected]")==0) return CONNECTED;
+            || strcmp(type, "[connected]")==0) return CONNECTED;
     if (strcmp(type, "[max]")==0
-        || strcmp(type, "[maxpool]")==0) return MAXPOOL;
+            || strcmp(type, "[maxpool]")==0) return MAXPOOL;
     if (strcmp(type, "[reorg]")==0) return REORG;
     if (strcmp(type, "[avg]")==0
-        || strcmp(type, "[avgpool]")==0) return AVGPOOL;
+            || strcmp(type, "[avgpool]")==0) return AVGPOOL;
     if (strcmp(type, "[dropout]")==0) return DROPOUT;
     if (strcmp(type, "[lrn]")==0
-        || strcmp(type, "[normalization]")==0) return NORMALIZATION;
+            || strcmp(type, "[normalization]")==0) return NORMALIZATION;
     if (strcmp(type, "[batchnorm]")==0) return BATCHNORM;
     if (strcmp(type, "[soft]")==0
-        || strcmp(type, "[softmax]")==0) return SOFTMAX;
+            || strcmp(type, "[softmax]")==0) return SOFTMAX;
     if (strcmp(type, "[route]")==0) return ROUTE;
     if (strcmp(type, "[upsample]")==0) return UPSAMPLE;
     return BLANK;
@@ -832,8 +832,8 @@ network *parse_network_cfg(char *filename)
             l.delta = net->layers[count-1].delta;
 #ifdef GPU
             if (gpu_index >= 0) {
-                l.output_gpu = net->layers[count - 1].output_gpu;
-                l.delta_gpu = net->layers[count - 1].delta_gpu;
+                l.output_gpu = net->layers[count-1].output_gpu;
+                l.delta_gpu = net->layers[count-1].delta_gpu;
             }
 #endif
         }else{
@@ -868,28 +868,26 @@ network *parse_network_cfg(char *filename)
     net->truths = out.outputs;
     if(net->layers[net->n-1].truths) net->truths = net->layers[net->n-1].truths;
     net->output = out.output;
-    //net->output = calloc(net->outputs*net->batch, sizeof(float));
     net->input = calloc(net->inputs*net->batch, sizeof(float));
     net->truth = calloc(net->truths*net->batch, sizeof(float));
 #ifdef GPU
     if (gpu_index >= 0) {
         net->output_gpu = out.output_gpu;
-        //net->output_gpu = opencl_make_array(net->output, net->outputs * net->batch);
         net->input_gpu = opencl_make_array(net->input, net->inputs * net->batch);
         net->truth_gpu = opencl_make_array(net->truth, net->truths * net->batch);
     }
 #endif
     if(workspace_size){
-        //printf("%ld\n", workspace_size*sizeof(float));
+        //printf("%ld\n", workspace_size);
 #ifdef GPU
         if(gpu_index >= 0){
-            net->workspace = calloc(workspace_size, sizeof(float));
-            net->workspace_gpu = opencl_make_array(net->workspace, workspace_size);
+            net->workspace = calloc(1, workspace_size);
+            net->workspace_gpu = opencl_make_array(net->workspace, (workspace_size-1)/sizeof(float)+1);
         }else {
-            net->workspace = calloc(workspace_size, sizeof(float));
+            net->workspace = calloc(1, workspace_size);
         }
 #else
-        net->workspace = calloc(workspace_size, sizeof(float));
+        net->workspace = calloc(1, workspace_size);
 #endif
     }
     return net;
@@ -1324,3 +1322,4 @@ void load_weights(network *net, char *filename)
 {
     load_weights_upto(net, filename, 0, net->n);
 }
+
diff --git a/src/region_layer.c b/src/region_layer.c
index 2ee79794..aa473070 100644
--- a/src/region_layer.c
+++ b/src/region_layer.c
@@ -44,8 +44,8 @@ layer make_region_layer(int batch, int w, int h, int n, int classes, int coords)
     if (gpu_index >= 0) {
         l.forward_gpu = forward_region_layer_gpu;
         l.backward_gpu = backward_region_layer_gpu;
-        l.output_gpu = opencl_make_array(l.output, batch * l.outputs);
-        l.delta_gpu = opencl_make_array(l.delta, batch * l.outputs);
+        l.output_gpu = opencl_make_array(l.output, batch*l.outputs);
+        l.delta_gpu = opencl_make_array(l.delta, batch*l.outputs);
     }
 #endif
 
@@ -60,6 +60,13 @@ void resize_region_layer(layer *l, int w, int h)
     l->w = w;
     l->h = h;
 
+#ifdef GPU
+    if (gpu_index >= 0) {
+        opencl_free_gpu_only(l->delta_gpu);
+        opencl_free_gpu_only(l->output_gpu);
+    }
+#endif
+
     l->outputs = h*w*l->n*(l->classes + l->coords + 1);
     l->inputs = l->outputs;
 
@@ -68,11 +75,8 @@ void resize_region_layer(layer *l, int w, int h)
 
 #ifdef GPU
     if (gpu_index >= 0) {
-        opencl_free_gpu_only(l->delta_gpu);
-        opencl_free_gpu_only(l->output_gpu);
-
-        l->delta_gpu = opencl_make_array(l->delta, l->batch * l->outputs);
-        l->output_gpu = opencl_make_array(l->output, l->batch * l->outputs);
+        l->delta_gpu = opencl_make_array(l->delta, l->batch*l->outputs);
+        l->output_gpu = opencl_make_array(l->output, l->batch*l->outputs);
     }
 #endif
 }
@@ -82,8 +86,8 @@ box get_region_box(float *x, float *biases, int n, int index, int i, int j, int
     box b;
     b.x = (i + x[index + 0*stride]) / w;
     b.y = (j + x[index + 1*stride]) / h;
-    b.w = expf(x[index + 2*stride]) * biases[2*n]   / w;
-    b.h = expf(x[index + 3*stride]) * biases[2*n+1] / h;
+    b.w = exp(x[index + 2*stride]) * biases[2*n]   / w;
+    b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h;
     return b;
 }
 
@@ -94,8 +98,8 @@ float delta_region_box(box truth, float *x, float *biases, int n, int index, int
 
     float tx = (truth.x*w - i);
     float ty = (truth.y*h - j);
-    float tw = logf(truth.w*w / biases[2*n]);
-    float th = logf(truth.h*h / biases[2*n + 1]);
+    float tw = log(truth.w*w / biases[2*n]);
+    float th = log(truth.h*h / biases[2*n + 1]);
 
     delta[index + 0*stride] = scale * (tx - x[index + 0*stride]);
     delta[index + 1*stride] = scale * (ty - x[index + 1*stride]);
@@ -162,35 +166,34 @@ int entry_index(layer l, int batch, int location, int entry)
 void forward_region_layer(const layer l, network net)
 {
     int i,j,b,t,n;
+    memcpy(l.output, net.input, l.outputs*l.batch*sizeof(float));
 
+#ifndef GPU
     if (gpu_index < 0) {
-        memcpy(l.output, net.input, l.outputs * l.batch * sizeof(float));
-
-        for (b = 0; b < l.batch; ++b) {
-            for (n = 0; n < l.n; ++n) {
-                int index = entry_index(l, b, n * l.w * l.h, 0);
-                activate_array(l.output + index, 2 * l.w * l.h, LOGISTIC);
-                index = entry_index(l, b, n * l.w * l.h, l.coords);
-                if (!l.background) activate_array(l.output + index, l.w * l.h, LOGISTIC);
-                index = entry_index(l, b, n * l.w * l.h, l.coords + 1);
-                if (!l.softmax && !l.softmax_tree) activate_array(l.output + index, l.classes * l.w * l.h, LOGISTIC);
-            }
+    for (b = 0; b < l.batch; ++b){
+        for(n = 0; n < l.n; ++n){
+            int index = entry_index(l, b, n*l.w*l.h, 0);
+            activate_array(l.output + index, 2*l.w*l.h, LOGISTIC);
+            index = entry_index(l, b, n*l.w*l.h, l.coords);
+            if(!l.background) activate_array(l.output + index,   l.w*l.h, LOGISTIC);
+            index = entry_index(l, b, n*l.w*l.h, l.coords + 1);
+            if(!l.softmax && !l.softmax_tree) activate_array(l.output + index, l.classes*l.w*l.h, LOGISTIC);
         }
-        if (l.softmax_tree) {
-            int i;
-            int count = l.coords + 1;
-            for (i = 0; i < l.softmax_tree->groups; ++i) {
-                int group_size = l.softmax_tree->group_size[i];
-                softmax_cpu(net.input + count, group_size, l.batch, l.inputs, l.n * l.w * l.h, 1, l.n * l.w * l.h,
-                            l.temperature, l.output + count);
-                count += group_size;
-            }
-        } else if (l.softmax) {
-            int index = entry_index(l, 0, 0, l.coords + !l.background);
-            softmax_cpu(net.input + index, l.classes + l.background, l.batch * l.n, l.inputs / l.n, l.w * l.h, 1,
-                        l.w * l.h, 1, l.output + index);
+    }
+    if (l.softmax_tree){
+        int i;
+        int count = l.coords + 1;
+        for (i = 0; i < l.softmax_tree->groups; ++i) {
+            int group_size = l.softmax_tree->group_size[i];
+            softmax_cpu(net.input + count, group_size, l.batch, l.inputs, l.n*l.w*l.h, 1, l.n*l.w*l.h, l.temperature, l.output + count);
+            count += group_size;
         }
+    } else if (l.softmax){
+        int index = entry_index(l, 0, 0, l.coords + !l.background);
+        softmax_cpu(net.input + index, l.classes + l.background, l.batch*l.n, l.inputs/l.n, l.w*l.h, 1, l.w*l.h, 1, l.output + index);
+    }
     }
+#endif
 
     memset(l.delta, 0, l.outputs * l.batch * sizeof(float));
     if(!net.train) return;
@@ -323,7 +326,7 @@ void forward_region_layer(const layer l, network net)
             ++class_count;
         }
     }
-    l.cost[0] = pow(mag_array(l.delta, l.outputs * l.batch), 2);
+    *(l.cost) = pow(mag_array(l.delta, l.outputs * l.batch), 2);
     printf("Region Avg IOU: %f, Class: %f, Obj: %f, No Obj: %f, Avg Recall: %f,  count: %d\n", avg_iou/count, avg_cat/class_count, avg_obj/count, avg_anyobj/(l.w*l.h*l.n*l.batch), recall/count, count);
 }
 
@@ -474,15 +477,11 @@ void forward_region_layer_gpu(const layer l, network net)
         return;
     }
 
-    //opencl_pull_array(l.output_gpu, net.input, l.batch*l.inputs);
-    opencl_pull_array(l.output_gpu, l.output, l.batch*l.outputs);
-    memcpy(net.input, l.output, l.batch*l.outputs*sizeof(float));
-
+    opencl_pull_array(l.output_gpu, net.input, l.batch*l.inputs);
     forward_region_layer(l, net);
-    // TODO: WHY NOT?
     //opencl_push_array(l.output_gpu, l.output, l.batch*l.outputs);
     if(!net.train) return;
-    opencl_pull_array(l.delta_gpu, l.delta, l.batch*l.outputs);
+    opencl_push_array(l.delta_gpu, l.delta, l.batch*l.outputs);
 }
 
 void backward_region_layer_gpu(const layer l, network net)
@@ -514,3 +513,4 @@ void zero_objectness(layer l)
         }
     }
 }
+
diff --git a/src/reorg_layer.c b/src/reorg_layer.c
index b808a1ff..ed78de0e 100644
--- a/src/reorg_layer.c
+++ b/src/reorg_layer.c
@@ -74,6 +74,13 @@ void resize_reorg_layer(layer *l, int w, int h)
         l->out_c = c*(stride*stride);
     }
 
+#ifdef GPU
+    if (gpu_index >= 0) {
+        opencl_free_gpu_only(l->output_gpu);
+        opencl_free_gpu_only(l->delta_gpu);
+    }
+#endif
+
     l->outputs = l->out_h * l->out_w * l->out_c;
     l->inputs = l->outputs;
     int output_size = l->outputs * l->batch;
@@ -83,9 +90,6 @@ void resize_reorg_layer(layer *l, int w, int h)
 
 #ifdef GPU
     if (gpu_index >= 0) {
-        opencl_free_gpu_only(l->output_gpu);
-        opencl_free_gpu_only(l->delta_gpu);
-
         l->output_gpu = opencl_make_array(l->output, output_size);
         l->delta_gpu = opencl_make_array(l->delta, output_size);
     }
diff --git a/src/rnn_layer.c b/src/rnn_layer.c
index e0d163b6..94351a68 100644
--- a/src/rnn_layer.c
+++ b/src/rnn_layer.c
@@ -69,8 +69,8 @@ layer make_rnn_layer(int batch, int inputs, int outputs, int steps, ACTIVATION a
         l.forward_gpu = forward_rnn_layer_gpu;
         l.backward_gpu = backward_rnn_layer_gpu;
         l.update_gpu = update_rnn_layer_gpu;
-        l.state_gpu = opencl_make_array(l.state, batch * outputs);
-        l.prev_state_gpu = opencl_make_array(l.prev_state, batch * outputs);
+        l.state_gpu = opencl_make_array(l.state, batch*outputs);
+        l.prev_state_gpu = opencl_make_array(l.prev_state, batch*outputs);
         l.output_gpu = l.output_layer->output_gpu;
         l.delta_gpu = l.output_layer->delta_gpu;
     }
diff --git a/src/route_layer.c b/src/route_layer.c
index bf2b7c63..62297470 100644
--- a/src/route_layer.c
+++ b/src/route_layer.c
@@ -31,8 +31,8 @@ route_layer make_route_layer(int batch, int n, int *input_layers, int *input_siz
     if (gpu_index >= 0) {
         l.forward_gpu = forward_route_layer_gpu;
         l.backward_gpu = backward_route_layer_gpu;
-        l.delta_gpu = opencl_make_array(l.delta, outputs * batch);
-        l.output_gpu = opencl_make_array(l.output, outputs * batch);
+        l.delta_gpu = opencl_make_array(l.delta, outputs*batch);
+        l.output_gpu = opencl_make_array(l.output, outputs*batch);
     }
 #endif
     return l;
diff --git a/src/shortcut_layer.c b/src/shortcut_layer.c
index 2f46e427..d6aaf6f6 100644
--- a/src/shortcut_layer.c
+++ b/src/shortcut_layer.c
@@ -32,8 +32,8 @@ layer make_shortcut_layer(int batch, int index, int w, int h, int c, int w2, int
     if (gpu_index >= 0) {
         l.forward_gpu = forward_shortcut_layer_gpu;
         l.backward_gpu = backward_shortcut_layer_gpu;
-        l.delta_gpu = opencl_make_array(l.delta, l.outputs * batch);
-        l.output_gpu = opencl_make_array(l.output, l.outputs * batch);
+        l.delta_gpu = opencl_make_array(l.delta, l.outputs*batch);
+        l.output_gpu = opencl_make_array(l.output, l.outputs*batch);
     }
 #endif
     return l;
@@ -45,6 +45,14 @@ void resize_shortcut_layer(layer *l, int w, int h)
     assert(l->h == l->out_h);
     l->w = l->out_w = w;
     l->h = l->out_h = h;
+
+#ifdef GPU
+    if (gpu_index >= 0) {
+        opencl_free_gpu_only(l->output_gpu);
+        opencl_free_gpu_only(l->delta_gpu);
+    }
+#endif
+
     l->outputs = w*h*l->out_c;
     l->inputs = l->outputs;
     l->delta =  realloc(l->delta, l->outputs*l->batch*sizeof(float));
@@ -52,11 +60,8 @@ void resize_shortcut_layer(layer *l, int w, int h)
 
 #ifdef GPU
     if (gpu_index >= 0) {
-        opencl_free_gpu_only(l->output_gpu);
-        opencl_free_gpu_only(l->delta_gpu);
-
-        l->output_gpu = opencl_make_array(l->output, l->outputs * l->batch);
-        l->delta_gpu = opencl_make_array(l->delta, l->outputs * l->batch);
+        l->output_gpu = opencl_make_array(l->output, l->outputs*l->batch);
+        l->delta_gpu = opencl_make_array(l->delta, l->outputs*l->batch);
     }
 #endif
     
diff --git a/src/softmax_layer.c b/src/softmax_layer.c
index 0d1f8c98..de57f51e 100644
--- a/src/softmax_layer.c
+++ b/src/softmax_layer.c
@@ -29,9 +29,9 @@ softmax_layer make_softmax_layer(int batch, int inputs, int groups)
     if (gpu_index >= 0) {
         l.forward_gpu = forward_softmax_layer_gpu;
         l.backward_gpu = backward_softmax_layer_gpu;
-        l.output_gpu = opencl_make_array(l.output, inputs * batch);
-        l.loss_gpu = opencl_make_array(l.loss, inputs * batch);
-        l.delta_gpu = opencl_make_array(l.delta, inputs * batch);
+        l.output_gpu = opencl_make_array(l.output, inputs*batch);
+        l.loss_gpu = opencl_make_array(l.loss, inputs*batch);
+        l.delta_gpu = opencl_make_array(l.delta, inputs*batch);
     }
 #endif
     return l;
@@ -89,13 +89,13 @@ void forward_softmax_layer_gpu(const softmax_layer l, network net)
         }
     }
     if(net.truth && !l.noloss) {
-        softmax_x_ent_gpu(l.batch * l.inputs, l.output_gpu, net.truth_gpu, l.delta_gpu, l.loss_gpu);
+        softmax_x_ent_gpu(l.batch*l.inputs, l.output_gpu, net.truth_gpu, l.delta_gpu, l.loss_gpu);
         if (l.softmax_tree) {
-            mask_gpu(l.batch * l.inputs, l.delta_gpu, SECRET_NUM, net.truth_gpu, 0);
-            mask_gpu(l.batch * l.inputs, l.loss_gpu, SECRET_NUM, net.truth_gpu, 0);
+            mask_gpu(l.batch*l.inputs, l.delta_gpu, SECRET_NUM, net.truth_gpu, 0);
+            mask_gpu(l.batch*l.inputs, l.loss_gpu, SECRET_NUM, net.truth_gpu, 0);
         }
-        opencl_pull_array(l.loss_gpu, l.loss, l.batch * l.inputs);
-        l.cost[0] = sum_array(l.loss, l.batch * l.inputs);
+        opencl_pull_array(l.loss_gpu, l.loss, l.batch*l.inputs);
+        l.cost[0] = sum_array(l.loss, l.batch*l.inputs);
     }
 }
 /*
diff --git a/src/upsample_layer.c b/src/upsample_layer.c
index f8445dae..8f940a4e 100644
--- a/src/upsample_layer.c
+++ b/src/upsample_layer.c
@@ -52,6 +52,14 @@ void resize_upsample_layer(layer *l, int w, int h)
         l->out_w = w/l->stride;
         l->out_h = h/l->stride;
     }
+
+#ifdef GPU
+    if (gpu_index >= 0) {
+        opencl_free_gpu_only(l->output_gpu);
+        opencl_free_gpu_only(l->delta_gpu);
+    }
+#endif
+
     l->outputs = l->out_w*l->out_h*l->out_c;
     l->inputs = l->h*l->w*l->c;
     l->delta =  realloc(l->delta, l->outputs*l->batch*sizeof(float));
@@ -59,11 +67,8 @@ void resize_upsample_layer(layer *l, int w, int h)
 
 #ifdef GPU
     if (gpu_index >= 0) {
-        opencl_free_gpu_only(l->output_gpu);
-        opencl_free_gpu_only(l->delta_gpu);
-
-        l->output_gpu = opencl_make_array(l->output, l->outputs * l->batch);
-        l->delta_gpu = opencl_make_array(l->delta, l->outputs * l->batch);
+        l->output_gpu = opencl_make_array(l->output, l->outputs*l->batch);
+        l->delta_gpu = opencl_make_array(l->delta, l->outputs*l->batch);
     }
 #endif
     
diff --git a/src/utils.c b/src/utils.c
index 609e199b..5332d191 100644
--- a/src/utils.c
+++ b/src/utils.c
@@ -188,6 +188,7 @@ char *find_char_arg(int argc, char **argv, char *arg, char *def)
     return def;
 }
 
+
 char *basecfg(char *cfgfile)
 {
     char *c = cfgfile;
@@ -697,15 +698,14 @@ float rand_normal()
 
 size_t rand_size_t()
 {
-    return  
-        ((size_t)(rand()&0xff) << 56) |
+    return  ((size_t)(rand()&0xff) << 56) | 
         ((size_t)(rand()&0xff) << 48) |
         ((size_t)(rand()&0xff) << 40) |
         ((size_t)(rand()&0xff) << 32) |
         ((size_t)(rand()&0xff) << 24) |
         ((size_t)(rand()&0xff) << 16) |
-        ((size_t)(rand()&0xff) <<  8) |
-        ((size_t)(rand()&0xff) <<  0);
+        ((size_t)(rand()&0xff) << 8) |
+        ((size_t)(rand()&0xff) << 0);
 }
 
 float rand_uniform(float min, float max)
diff --git a/src/yolo_layer.c b/src/yolo_layer.c
index 5fd12a2e..3ce4c3b2 100644
--- a/src/yolo_layer.c
+++ b/src/yolo_layer.c
@@ -89,8 +89,8 @@ box get_yolo_box(float *x, float *biases, int n, int index, int i, int j, int lw
     box b;
     b.x = (i + x[index + 0*stride]) / lw;
     b.y = (j + x[index + 1*stride]) / lh;
-    b.w = expf(x[index + 2*stride]) * biases[2*n]   / w;
-    b.h = expf(x[index + 3*stride]) * biases[2*n+1] / h;
+    b.w = exp(x[index + 2*stride]) * biases[2*n]   / w;
+    b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h;
     return b;
 }
 
@@ -101,8 +101,8 @@ float delta_yolo_box(box truth, float *x, float *biases, int n, int index, int i
 
     float tx = (truth.x*lw - i);
     float ty = (truth.y*lh - j);
-    float tw = logf(truth.w*w / biases[2*n]);
-    float th = logf(truth.h*h / biases[2*n + 1]);
+    float tw = log(truth.w*w / biases[2*n]);
+    float th = log(truth.h*h / biases[2*n + 1]);
 
     delta[index + 0*stride] = scale * (tx - x[index + 0*stride]);
     delta[index + 1*stride] = scale * (ty - x[index + 1*stride]);
@@ -136,19 +136,20 @@ static int entry_index(layer l, int batch, int location, int entry)
 void forward_yolo_layer(const layer l, network net)
 {
     int i,j,b,t,n;
+    memcpy(l.output, net.input, l.outputs*l.batch*sizeof(float));
 
+#ifndef GPU
     if (gpu_index < 0) {
-        memcpy(l.output, net.input, l.outputs * l.batch * sizeof(float));
-
-        for (b = 0; b < l.batch; ++b) {
-            for (n = 0; n < l.n; ++n) {
-                int index = entry_index(l, b, n * l.w * l.h, 0);
-                activate_array(l.output + index, 2 * l.w * l.h, LOGISTIC);
-                index = entry_index(l, b, n * l.w * l.h, 4);
-                activate_array(l.output + index, (1 + l.classes) * l.w * l.h, LOGISTIC);
+    for (b = 0; b < l.batch; ++b){
+        for(n = 0; n < l.n; ++n){
+            int index = entry_index(l, b, n*l.w*l.h, 0);
+            activate_array(l.output + index, 2*l.w*l.h, LOGISTIC);
+            index = entry_index(l, b, n*l.w*l.h, 4);
+            activate_array(l.output + index, (1+l.classes)*l.w*l.h, LOGISTIC);
             }
         }
     }
+#endif
 
     memset(l.delta, 0, l.outputs * l.batch * sizeof(float));
     if(!net.train) return;
@@ -201,7 +202,7 @@ void forward_yolo_layer(const layer l, network net)
             box truth = float_to_box(net.truth + t*(4 + 1) + b*l.truths, 1);
 
             if(!truth.x) break;
-            float best_iou = -1.f;
+            float best_iou = 0;
             int best_n = 0;
             i = (truth.x * l.w);
             j = (truth.y * l.h);
@@ -240,7 +241,7 @@ void forward_yolo_layer(const layer l, network net)
             }
         }
     }
-    l.cost[0] = pow(mag_array(l.delta, l.outputs * l.batch), 2);
+    *(l.cost) = pow(mag_array(l.delta, l.outputs * l.batch), 2);
     printf("Region %d Avg IOU: %f, Class: %f, Obj: %f, No Obj: %f, .5R: %f, .75R: %f,  count: %d\n", net.index, avg_iou/count, avg_cat/class_count, avg_obj/count, avg_anyobj/(l.w*l.h*l.n*l.batch), recall/count, recall75/count, count);
 }
 
@@ -366,15 +367,9 @@ void forward_yolo_layer_gpu(const layer l, network net)
         return;
     }
 
-    //opencl_pull_array(l.output_gpu, net.input, l.batch*l.inputs);
-    opencl_pull_array(l.output_gpu, l.output, l.batch*l.outputs);
-    memcpy(net.input, l.output, l.batch*l.outputs*sizeof(float));
-
+    opencl_pull_array(l.output_gpu, net.input, l.batch*l.inputs);
     forward_yolo_layer(l, net);
-    // TODO: WHY NOT?
-    //opencl_push_array(l.output_gpu, l.output, l.batch*l.outputs);
-    if(!net.train) return;
-    opencl_pull_array(l.delta_gpu, l.delta, l.batch*l.outputs);
+    opencl_push_array(l.delta_gpu, l.delta, l.batch*l.outputs);
 }
 
 void backward_yolo_layer_gpu(const layer l, network net)