arm_compute v17.10

Change-Id: If1489af40eccd0219ede8946577afbf04db31b29

Author: Kaizen

README.md

@@ -9,6 +9,7 @@ Related projects:
 
 Documentation available here:
 
+- [v17.10](https://arm-software.github.io/ComputeLibrary/v17.10/)
 - [v17.09](https://arm-software.github.io/ComputeLibrary/v17.09/)
 - [v17.06](https://arm-software.github.io/ComputeLibrary/v17.06/)
 - [v17.05](https://arm-software.github.io/ComputeLibrary/v17.05/)
@@ -17,6 +18,7 @@ Documentation available here:
 
 Binaries available here:
 
+- [v17.10](https://github.com/ARM-software/ComputeLibrary/releases/download/v17.10/arm_compute-v17.10-bin.tar.gz)
 - [v17.09](https://github.com/ARM-software/ComputeLibrary/releases/download/v17.09/arm_compute-v17.09-bin.tar.gz)
 - [v17.06](https://github.com/ARM-software/ComputeLibrary/releases/download/v17.06/arm_compute-v17.06-bin.tar.gz)
 - [v17.05](https://github.com/ARM-software/ComputeLibrary/releases/download/v17.05/arm_compute-v17.05-bin.tar.gz)

SConscript

@@ -24,18 +24,18 @@ import os.path
 import re
 import subprocess
 
-VERSION = "v17.09"
-SONAME_VERSION="4.0.0"
+VERSION = "v17.10"
+SONAME_VERSION="5.0.0"
 
 Import('env')
 Import('vars')
 
-def build_library(name, sources, static=False):
+def build_library(name, sources, static=False, libs=[]):
     if static:
-        obj = arm_compute_env.StaticLibrary(name, source=sources)
+        obj = arm_compute_env.StaticLibrary(name, source=sources, LIBS = arm_compute_env["LIBS"] + libs)
     else:
         if env['set_soname']:
-            obj = arm_compute_env.SharedLibrary(name, source=sources, SHLIBVERSION = SONAME_VERSION)
+            obj = arm_compute_env.SharedLibrary(name, source=sources, SHLIBVERSION = SONAME_VERSION, LIBS = arm_compute_env["LIBS"] + libs)
 
             symlinks = []
             # Manually delete symlinks or SCons will get confused:
@@ -51,7 +51,7 @@ def build_library(name, sources, static=False):
             Default(clean)
             Depends(obj, clean)
         else:
-            obj = arm_compute_env.SharedLibrary(name, source=sources)
+            obj = arm_compute_env.SharedLibrary(name, source=sources, LIBS = arm_compute_env["LIBS"] + libs)
 
     Default(obj)
     return obj
@@ -191,11 +191,12 @@ if env['os'] != 'bare_metal' and not env['standalone']:
 shared_runtime_objects = [arm_compute_env.SharedObject(f) for f in runtime_files]
 static_runtime_objects = [arm_compute_env.StaticObject(f) for f in runtime_files]
 
-arm_compute_a = build_library('arm_compute-static', static_core_objects + static_runtime_objects, static=True)
+arm_compute_a = build_library('arm_compute-static', static_runtime_objects, static=True, libs = [ arm_compute_core_a ])
 Export('arm_compute_a')
 
 if env['os'] != 'bare_metal' and not env['standalone']:
-    arm_compute_so = build_library('arm_compute', shared_core_objects + shared_runtime_objects, static=False)
+    arm_compute_so = build_library('arm_compute', shared_runtime_objects, static=False, libs = [ "arm_compute_core" ])
+    Depends(arm_compute_so, arm_compute_core_so)
     Export('arm_compute_so')
 
 if env['neon'] and env['opencl']:
@@ -208,10 +209,11 @@ if env['neon'] and env['opencl']:
     shared_graph_objects = [arm_compute_env.SharedObject(f) for f in graph_files]
     static_graph_objects = [arm_compute_env.StaticObject(f) for f in graph_files]
 
-    arm_compute_graph_a = build_library('arm_compute_graph-static', static_core_objects + static_runtime_objects + static_graph_objects, static=True)
+    arm_compute_graph_a = build_library('arm_compute_graph-static', static_graph_objects, static=True, libs = [ arm_compute_a ])
     Export('arm_compute_graph_a')
 
-    arm_compute_graph_so = build_library('arm_compute_graph', shared_core_objects + shared_runtime_objects + shared_graph_objects, static=False)
+    arm_compute_graph_so = build_library('arm_compute_graph', shared_graph_objects, static=False, libs = [ "arm_compute", "arm_compute_core" ])
+    Depends( arm_compute_graph_so, arm_compute_so)
     Export('arm_compute_graph_so')
 
     graph_alias = arm_compute_env.Alias("arm_compute_graph", [arm_compute_graph_a, arm_compute_graph_so])

SConstruct

@@ -56,6 +56,7 @@ vars.AddVariables(
 )
 
 env = Environment(platform="posix", variables=vars, ENV = os.environ)
+env.Append(LIBPATH = ["#build/%s" % env['build_dir']])
 
 SConsignFile('build/.%s' % env['build_dir'])
 
@@ -107,7 +108,7 @@ if env['arch'] == 'armv7a':
         prefix = "arm-linux-gnueabihf-"
         env.Append(CXXFLAGS = ['-mfloat-abi=hard'])
     elif env['os'] == 'bare_metal':
-        prefix = "arm-none-eabi-"
+        prefix = "arm-eabi-"
         env.Append(CXXFLAGS = ['-mfloat-abi=hard'])
     elif env['os'] == 'android':
         prefix = "arm-linux-androideabi-"
@@ -118,7 +119,7 @@ elif env['arch'] == 'arm64-v8a':
     if env['os'] == 'linux':
         prefix = "aarch64-linux-gnu-"
     elif env['os'] == 'bare_metal':
-        prefix = "aarch64-none-elf-"
+        prefix = "aarch64-elf-"
     elif env['os'] == 'android':
         prefix = "aarch64-linux-android-"
 elif env['arch'] == 'arm64-v8.2-a':
@@ -216,7 +217,7 @@ SConscript('./SConscript', variant_dir='#build/%s' % env['build_dir'], duplicate
 if env['opencl']:
     SConscript("./opencl-1.2-stubs/SConscript", variant_dir="build/%s/opencl-1.2-stubs" % env['build_dir'], duplicate=0)
 
-if env['examples']:
+if env['examples'] and env['os'] != 'bare_metal':
     SConscript('./examples/SConscript', variant_dir='#build/%s/examples' % env['build_dir'], duplicate=0)
 
 if env['os'] != 'bare_metal':

arm_compute/core/CL/CLKernelLibrary.h

@@ -162,11 +162,9 @@ class CLKernelLibrary
      */
     void init(std::string kernel_path = ".", cl::Context context = cl::Context::getDefault(), cl::Device device = cl::Device::getDefault())
     {
-        _kernel_path        = std::move(kernel_path);
-        _context            = std::move(context);
-        _device             = std::move(device);
-        _max_workgroup_size = 0;
-        max_local_workgroup_size();
+        _kernel_path = std::move(kernel_path);
+        _context     = std::move(context);
+        _device      = std::move(device);
     }
     /** Sets the path that the kernels reside in.
      *
@@ -208,20 +206,15 @@ class CLKernelLibrary
         {
             _device = cl_devices[0];
         }
-
-        _max_workgroup_size = 0;
-        max_local_workgroup_size();
-    };
+    }
     /** Sets the CL device for which the programs are created.
      *
      * @param[in] device A CL device.
      */
     void set_device(cl::Device device)
     {
-        _device             = std::move(device);
-        _max_workgroup_size = 0;
-        max_local_workgroup_size();
-    };
+        _device = std::move(device);
+    }
     /** Creates a kernel from the kernel library.
      *
      * @param[in] kernel_name       Kernel name.
@@ -238,15 +231,14 @@ class CLKernelLibrary
      *
      */
     void load_binary();
-    /** Find the maximum number of local work items in a workgroup can be supported by the device
+    /** Find the maximum number of local work items in a workgroup can be supported for the kernel.
      *
      */
-    size_t max_local_workgroup_size();
-
-    /** Return the default NDRange that is suitable for the device.
+    size_t max_local_workgroup_size(const cl::Kernel &kernel) const;
+    /** Return the default NDRange for the device.
      *
      */
-    cl::NDRange default_ndrange();
+    cl::NDRange default_ndrange() const;
 
 private:
     /** Load program and its dependencies.
@@ -270,7 +262,6 @@ class CLKernelLibrary
     static const std::map<std::string, std::string> _kernel_program_map; /**< Map that associates kernel names with programs. */
     static const std::map<std::string, std::string> _program_source_map; /**< Contains sources for all programs.
                                                                               Used for compile-time kernel inclusion. >*/
-    size_t _max_workgroup_size;                                          /** Maximum local workgroup size supported on the device */
 };
 }
 #endif /* __ARM_COMPUTE_CLKERNELLIBRARY_H__ */

arm_compute/core/CL/ICLKernel.h

@@ -175,6 +175,12 @@ class ICLKernel : public IKernel
      */
     GPUTarget get_target() const;
 
+    /** Get the maximum workgroup size for the device the CLKernelLibrary uses.
+     *
+     * @return The maximum workgroup size value.
+     */
+    size_t get_max_workgroup_size();
+
 private:
     /** Add the passed array's parameters to the object's kernel's arguments starting from the index idx.
      *
@@ -208,10 +214,11 @@ class ICLKernel : public IKernel
     unsigned int           num_arguments_per_tensor() const;
 
 protected:
-    cl::Kernel  _kernel;    /**< OpenCL kernel to run */
-    cl::NDRange _lws_hint;  /**< Local workgroup size hint for the OpenCL kernel */
-    GPUTarget   _target;    /**< The targeted GPU */
-    std::string _config_id; /**< Configuration ID */
+    cl::Kernel  _kernel;             /**< OpenCL kernel to run */
+    cl::NDRange _lws_hint;           /**< Local workgroup size hint for the OpenCL kernel */
+    GPUTarget   _target;             /**< The targeted GPU */
+    std::string _config_id;          /**< Configuration ID */
+    size_t      _max_workgroup_size; /**< The maximum workgroup size for this kernel */
 };
 
 /** Add the kernel to the command queue with the given window.
@@ -223,7 +230,7 @@ class ICLKernel : public IKernel
  * @param[in,out] queue    OpenCL command queue.
  * @param[in]     kernel   Kernel to enqueue
  * @param[in]     window   Window the kernel has to process.
- * @param[in]     lws_hint Local workgroup size requested, by default (128,1)
+ * @param[in]     lws_hint Local workgroup size requested, by default (128,1).
  *
  * @note If any dimension of the lws is greater than the global workgroup size then no lws will be passed.
  */

arm_compute/core/CL/OpenCL.h

@@ -83,6 +83,8 @@ class CLSymbols final
     using clGetDeviceInfo_func           = cl_int (*)(cl_device_id, cl_device_info, size_t, void *, size_t *);
     using clGetDeviceIDs_func            = cl_int (*)(cl_platform_id, cl_device_type, cl_uint, cl_device_id *, cl_uint *);
     using clRetainEvent_func             = cl_int (*)(cl_event);
+    using clGetPlatformIDs_func          = cl_int (*)(cl_uint, cl_platform_id *, cl_uint *);
+    using clGetKernelWorkGroupInfo_func  = cl_int (*)(cl_kernel, cl_device_id, cl_kernel_work_group_info, size_t, void *, size_t *);
 
     clBuildProgram_func            clBuildProgram            = nullptr;
     clEnqueueNDRangeKernel_func    clEnqueueNDRangeKernel    = nullptr;
@@ -113,6 +115,8 @@ class CLSymbols final
     clGetDeviceInfo_func           clGetDeviceInfo           = nullptr;
     clGetDeviceIDs_func            clGetDeviceIDs            = nullptr;
     clRetainEvent_func             clRetainEvent             = nullptr;
+    clGetPlatformIDs_func          clGetPlatformIDs          = nullptr;
+    clGetKernelWorkGroupInfo_func  clGetKernelWorkGroupInfo  = nullptr;
 
 private:
     std::pair<bool, bool> _loaded{ false, false };

arm_compute/core/Logger.h

@@ -0,0 +1,71 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+
+#ifndef __ARM_COMPUTE_LOGGER_H__
+#define __ARM_COMPUTE_LOGGER_H__
+
+#include <iostream>
+#include <memory>
+
+#ifdef ARM_COMPUTE_DEBUG_ENABLED
+#define ARM_COMPUTE_LOG(x) (arm_compute::Logger::get().log_info() << x)
+#else /* ARM_COMPUTE_DEBUG_ENABLED */
+#define ARM_COMPUTE_LOG(...)
+#endif /* ARM_COMPUTE_DEBUG_ENABLED */
+
+namespace arm_compute
+{
+/**< Verbosity of the logger */
+enum class LoggerVerbosity
+{
+    NONE, /**< No info */
+    INFO  /**< Log info */
+};
+
+/** Logger singleton class */
+class Logger
+{
+public:
+    static Logger &get();
+    void set_logger(std::ostream &ostream, LoggerVerbosity verbosity);
+    std::ostream &log_info();
+
+private:
+    /** Default constructor */
+    Logger();
+    /** Allow instances of this class to be moved */
+    Logger(Logger &&) = default;
+    /** Prevent instances of this class from being copied (As this class contains pointers) */
+    Logger(const Logger &) = delete;
+    /** Prevent instances of this class from being copied (As this class contains pointers) */
+    Logger &operator=(const Logger &) = delete;
+    /** Allow instances of this class to be moved */
+    Logger &operator=(Logger &&) = default;
+
+    std::ostream   *_ostream;
+    std::ostream    _nullstream;
+    LoggerVerbosity _verbosity;
+};
+} // arm_compute
+#endif /* __ARM_COMPUTE_LOGGER_H__ */

arm_compute/core/NEON/NEFixedPoint.inl

@@ -21,6 +21,7 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  * SOFTWARE.
  */
+#include <array>
 #include <limits>
 
 namespace arm_compute

arm_compute/core/NEON/kernels/NELKTrackerKernel.h

@@ -109,7 +109,7 @@ class NELKTrackerKernel : public INEKernel
      *
      * @return Values A11, A12, A22
      */
-    std::tuple<int, int, int> compute_spatial_gradient_matrix(const NELKInternalKeypoint &keypoint, int *bilinear_ix, int *bilinear_iy);
+    std::tuple<int, int, int> compute_spatial_gradient_matrix(const NELKInternalKeypoint &keypoint, int32_t *bilinear_ix, int32_t *bilinear_iy);
     /** Compute the vector A^T * b, i.e. -sum(I_d * I_t) for d in {x,y}
      *
      * @param[in] old_keypoint Old keypoint for which gradient is computed
@@ -119,7 +119,7 @@ class NELKTrackerKernel : public INEKernel
      *
      * @return Values b1, b2
      */
-    std::pair<int, int> compute_image_mismatch_vector(const NELKInternalKeypoint &old_keypoint, const NELKInternalKeypoint &new_keypoint, const int *bilinear_ix, const int *bilinear_iy);
+    std::pair<int, int> compute_image_mismatch_vector(const NELKInternalKeypoint &old_keypoint, const NELKInternalKeypoint &new_keypoint, const int32_t *bilinear_ix, const int32_t *bilinear_iy);
 
     const ITensor              *_input_old;
     const ITensor              *_input_new;

arm_compute/core/NEON/kernels/NEPoolingLayerKernel.h

@@ -47,6 +47,8 @@ class NEPoolingLayerKernel : public INEKernel
     /** Default destructor */
     ~NEPoolingLayerKernel() = default;
     /** Set the input and output tensors.
+     *
+     * @note QS8, QS16 and F16 are supported for pool sizes 2 and 3 only
      *
      * @param[in]  input     Source tensor. Data types supported: QS8/QS16/F16/F32.
      * @param[out] output    Destination tensor. Data types supported: Same as @p input.
@@ -123,6 +125,13 @@ class NEPoolingLayerKernel : public INEKernel
      */
     template <PoolingType pooling_type>
     void pooling7_f32(const Window &window_input, const Window &window);
+    /** Function to perform NxN pooling.
+     *
+     * @param[in] window_input Input region on which to execute the kernel.
+     * @param[in] window       Output region on which to execute the kernel.
+     */
+    template <PoolingType pooling_type>
+    void poolingN_f32(const Window &window_input, const Window &window);
     /** Common signature for all the specialised Pooling functions
      *
      * @param[in] window_input Input region on which to execute the kernel.