arm_compute v17.09

Change-Id: I4bf8f4e6e5f84ce0d5b6f5ba570d276879f42a81

Author: Kaizen

LICENSE

@@ -1,19 +1,23 @@
 
 Please report issues here: https://github.com/ARM-software/ComputeLibrary/issues
+Make sure you are using the latest version of the library before opening an issue. Thanks
 
 Related projects:
 
 - [Caffe on Compute Library](https://github.com/OAID/caffeOnACL)
+- [Tutorial: Cartoonifying Images on Raspberry Pi with the Compute Library](https://community.arm.com/graphics/b/blog/posts/cartoonifying-images-on-raspberry-pi-with-the-compute-library)
 
 Documentation available here:
 
+- [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/)
 - [v17.04](https://arm-software.github.io/ComputeLibrary/v17.04/)
 - [v17.03.1](https://arm-software.github.io/ComputeLibrary/v17.03.1/)
 
 Binaries available here:
 
+- [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)
 - [v17.04](https://github.com/ARM-software/ComputeLibrary/releases/download/v17.04/arm_compute-v17.04-bin.tar.gz)

README.md

@@ -24,8 +24,8 @@ import os.path
 import re
 import subprocess
 
-VERSION = "v17.06"
-SONAME_VERSION="3.0.0"
+VERSION = "v17.09"
+SONAME_VERSION="4.0.0"
 
 Import('env')
 Import('vars')
@@ -138,22 +138,22 @@ core_files = Glob('src/core/*.cpp')
 core_files += Glob('src/core/CPP/*.cpp')
 core_files += Glob('src/core/CPP/kernels/*.cpp')
 
-files = Glob('src/runtime/*.cpp')
+runtime_files = Glob('src/runtime/*.cpp')
 # CLHarrisCorners uses the Scheduler to run CPP kernels
-files += Glob('src/runtime/CPP/SingleThreadScheduler.cpp')
+runtime_files += Glob('src/runtime/CPP/SingleThreadScheduler.cpp')
 
 if env['cppthreads']:
-     files += Glob('src/runtime/CPP/CPPScheduler.cpp')
+     runtime_files += Glob('src/runtime/CPP/CPPScheduler.cpp')
 
 if env['openmp']:
-     files += Glob('src/runtime/OMP/OMPScheduler.cpp')
+     runtime_files += Glob('src/runtime/OMP/OMPScheduler.cpp')
 
 if env['opencl']:
     core_files += Glob('src/core/CL/*.cpp')
     core_files += Glob('src/core/CL/kernels/*.cpp')
 
-    files += Glob('src/runtime/CL/*.cpp')
-    files += Glob('src/runtime/CL/functions/*.cpp')
+    runtime_files += Glob('src/runtime/CL/*.cpp')
+    runtime_files += Glob('src/runtime/CL/functions/*.cpp')
 
     # Generate embed files
     if env['embed_kernels']:
@@ -169,31 +169,64 @@ if env['neon']:
     core_files += Glob('src/core/NEON/*.cpp')
     core_files += Glob('src/core/NEON/kernels/*.cpp')
 
-    files += Glob('src/runtime/NEON/*.cpp')
-    files += Glob('src/runtime/NEON/functions/*.cpp')
+    if env['arch'] == "armv7a":
+        core_files += Glob('src/core/NEON/kernels/arm32/*.cpp')
+
+    if "arm64-v8" in env['arch']:
+        core_files += Glob('src/core/NEON/kernels/arm64/*.cpp')
+
+    runtime_files += Glob('src/runtime/NEON/*.cpp')
+    runtime_files += Glob('src/runtime/NEON/functions/*.cpp')
 
 static_core_objects = [arm_compute_env.StaticObject(f) for f in core_files]
 shared_core_objects = [arm_compute_env.SharedObject(f) for f in core_files]
 
 arm_compute_core_a = build_library('arm_compute_core-static', static_core_objects, static=True)
 Export('arm_compute_core_a')
 
-if env['os'] != 'bare_metal':
+if env['os'] != 'bare_metal' and not env['standalone']:
     arm_compute_core_so = build_library('arm_compute_core', shared_core_objects, static=False)
     Export('arm_compute_core_so')
 
-shared_objects = [arm_compute_env.SharedObject(f) for f in files]
-static_objects = [arm_compute_env.StaticObject(f) for f in files]
+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_objects, static=True)
+arm_compute_a = build_library('arm_compute-static', static_core_objects + static_runtime_objects, static=True)
 Export('arm_compute_a')
 
-if env['os'] != 'bare_metal':
-    arm_compute_so = build_library('arm_compute', shared_core_objects + shared_objects, static=False)
+if env['os'] != 'bare_metal' and not env['standalone']:
+    arm_compute_so = build_library('arm_compute', shared_core_objects + shared_runtime_objects, static=False)
     Export('arm_compute_so')
 
-alias = arm_compute_env.Alias("arm_compute", [arm_compute_a, arm_compute_so])
+if env['neon'] and env['opencl']:
+    graph_files = Glob('src/graph/*.cpp')
+    graph_files += Glob('src/graph/nodes/*.cpp')
+
+    graph_files += Glob('src/graph/CL/*.cpp')
+    graph_files += Glob('src/graph/NEON/*.cpp')
+
+    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)
+    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)
+    Export('arm_compute_graph_so')
+
+    graph_alias = arm_compute_env.Alias("arm_compute_graph", [arm_compute_graph_a, arm_compute_graph_so])
+    Default(graph_alias)
+
+if env['standalone']:
+    alias = arm_compute_env.Alias("arm_compute", [arm_compute_a])
+else:
+    alias = arm_compute_env.Alias("arm_compute", [arm_compute_a, arm_compute_so])
+
 Default(alias)
 
 Default(generate_embed)
-Depends([alias,arm_compute_core_so, arm_compute_core_a], generate_embed)
+
+if env['standalone']:
+    Depends([alias,arm_compute_core_a], generate_embed)
+else:
+    Depends([alias,arm_compute_core_so, arm_compute_core_a], generate_embed)

SConscript

@@ -44,6 +44,7 @@ vars.AddVariables(
     EnumVariable("build", "Build type", "cross_compile", allowed_values=("native", "cross_compile")),
     BoolVariable("examples", "Build example programs", True),
     BoolVariable("Werror", "Enable/disable the -Werror compilation flag", True),
+    BoolVariable("standalone", "Builds the tests as standalone executables, links statically with libgcc, libstdc++ and libarm_compute", False),
     BoolVariable("opencl", "Enable OpenCL support", True),
     BoolVariable("neon", "Enable Neon support", False),
     BoolVariable("embed_kernels", "Embed OpenCL kernels in library binary", False),
@@ -102,25 +103,31 @@ prefix = ""
 if env['arch'] == 'armv7a':
     env.Append(CXXFLAGS = ['-march=armv7-a', '-mthumb', '-mfpu=neon'])
 
-    if env['os'] in ['linux', 'bare_metal']:
+    if env['os'] == 'linux':
         prefix = "arm-linux-gnueabihf-"
         env.Append(CXXFLAGS = ['-mfloat-abi=hard'])
+    elif env['os'] == 'bare_metal':
+        prefix = "arm-none-eabi-"
+        env.Append(CXXFLAGS = ['-mfloat-abi=hard'])
     elif env['os'] == 'android':
         prefix = "arm-linux-androideabi-"
         env.Append(CXXFLAGS = ['-mfloat-abi=softfp'])
 elif env['arch'] == 'arm64-v8a':
     env.Append(CXXFLAGS = ['-march=armv8-a'])
 
-    if env['os'] in ['linux', 'bare_metal']:
+    if env['os'] == 'linux':
         prefix = "aarch64-linux-gnu-"
+    elif env['os'] == 'bare_metal':
+        prefix = "aarch64-none-elf-"
     elif env['os'] == 'android':
         prefix = "aarch64-linux-android-"
 elif env['arch'] == 'arm64-v8.2-a':
     env.Append(CXXFLAGS = ['-march=armv8.2-a+fp16+simd'])
     env.Append(CPPDEFINES = ['ARM_COMPUTE_ENABLE_FP16'])
-
-    if env['os'] in ['linux', 'bare_metal']:
+    if env['os'] == 'linux':
         prefix = "aarch64-linux-gnu-"
+    elif env['os'] == 'bare_metal':
+        prefix = "aarch64-elf-"
     elif env['os'] == 'android':
         prefix = "aarch64-linux-android-"
 elif env['arch'] == 'x86_32':
@@ -161,6 +168,10 @@ if not GetOption("help"):
         if compiler_ver == '4.8.3':
             env.Append(CXXFLAGS = ['-Wno-array-bounds'])
 
+if env['standalone']:
+    env.Append(CXXFLAGS = ['-fPIC'])
+    env.Append(LINKFLAGS = ['-static-libgcc','-static-libstdc++'])
+
 if env['Werror']:
     env.Append(CXXFLAGS = ['-Werror'])
 
@@ -169,8 +180,10 @@ if env['os'] == 'android':
     env.Append(LINKFLAGS = ['-pie', '-static-libstdc++'])
 elif env['os'] == 'bare_metal':
     env.Append(LINKFLAGS = ['-static'])
+    env.Append(LINKFLAGS = ['-specs=rdimon.specs'])
     env.Append(CXXFLAGS = ['-fPIC'])
     env.Append(CPPDEFINES = ['NO_MULTI_THREADING'])
+    env.Append(CPPDEFINES = ['BARE_METAL'])
 
 if env['opencl']:
     if env['os'] == 'bare_metal':
@@ -189,6 +202,7 @@ else:
 
 if env['asserts']:
     env.Append(CPPDEFINES = ['ARM_COMPUTE_ASSERTS_ENABLED'])
+    env.Append(CXXFLAGS = ['-fstack-protector-strong'])
 
 env.Append(CPPPATH = ['#/include', "#"])
 env.Append(CXXFLAGS = env['extra_cxx_flags'])
@@ -205,4 +219,5 @@ if env['opencl']:
 if env['examples']:
     SConscript('./examples/SConscript', variant_dir='#build/%s/examples' % env['build_dir'], duplicate=0)
 
-SConscript('./tests/SConscript', variant_dir='#build/%s/tests' % env['build_dir'], duplicate=0)
+if env['os'] != 'bare_metal':
+    SConscript('./tests/SConscript', variant_dir='#build/%s/tests' % env['build_dir'], duplicate=0)

SConstruct

@@ -26,6 +26,7 @@
 
 #include "arm_compute/core/CL/OpenCL.h"
 #include "arm_compute/core/Helpers.h"
+#include "support/ToolchainSupport.h"
 
 #include <string>
 
@@ -52,6 +53,22 @@ static constexpr const unsigned int max_cl_vector_width = 16;
  */
 std::string get_cl_type_from_data_type(const DataType &dt);
 
+/** Get the size of a data type in number of bits.
+ *
+ * @param[in] dt @ref DataType.
+ *
+ * @return Number of bits in the data type specified.
+ */
+std::string get_data_size_from_data_type(const DataType &dt);
+
+/** Translates fixed point tensor data type to the underlying OpenCL type.
+ *
+ * @param[in] dt @ref DataType to be translated to OpenCL type.
+ *
+ * @return The string specifying the underlying OpenCL type to be used.
+ */
+std::string get_underlying_cl_type_from_data_type(const DataType &dt);
+
 /** Translates a given gpu device target to string.
  *
  * @param[in] target Given gpu target.
@@ -70,7 +87,7 @@ const std::string &string_from_target(GPUTarget target);
 template <typename Kernel, typename... T>
 std::unique_ptr<Kernel> create_configure_kernel(T &&... args)
 {
-    std::unique_ptr<Kernel> k = arm_compute::cpp14::make_unique<Kernel>();
+    std::unique_ptr<Kernel> k = arm_compute::support::cpp14::make_unique<Kernel>();
     k->configure(std::forward<T>(args)...);
     return k;
 }
@@ -82,7 +99,7 @@ std::unique_ptr<Kernel> create_configure_kernel(T &&... args)
 template <typename Kernel>
 std::unique_ptr<Kernel> create_kernel()
 {
-    std::unique_ptr<Kernel> k = arm_compute::cpp14::make_unique<Kernel>();
+    std::unique_ptr<Kernel> k = arm_compute::support::cpp14::make_unique<Kernel>();
     return k;
 }
 
@@ -101,5 +118,20 @@ GPUTarget get_target_from_device(cl::Device &device);
  * @return the GPU target which shows the arch
  */
 GPUTarget get_arch_from_target(GPUTarget target);
+
+/** Helper function to get the highest OpenCL version supported
+ *
+ * @param[in] device A CL device
+ *
+ * @return the highest OpenCL version supported
+ */
+CLVersion get_cl_version(const cl::Device &device);
+/** Helper function to check whether the arm_non_uniform_work_group_size extension is supported
+ *
+ * @param[in] device A CL device
+ *
+ * @return True if the extension is supported
+ */
+bool non_uniform_workgroup_support(const cl::Device &device);
 }
-#endif
+#endif /* __ARM_COMPUTE_CLHELPERS_H__ */

arm_compute/core/CL/CLHelpers.h

@@ -162,9 +162,11 @@ 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);
+        _kernel_path        = std::move(kernel_path);
+        _context            = std::move(context);
+        _device             = std::move(device);
+        _max_workgroup_size = 0;
+        max_local_workgroup_size();
     }
     /** Sets the path that the kernels reside in.
      *
@@ -174,6 +176,17 @@ class CLKernelLibrary
     {
         _kernel_path = kernel_path;
     };
+    /** Gets the path that the kernels reside in.
+     */
+    std::string get_kernel_path()
+    {
+        return _kernel_path;
+    };
+    /** Gets the source of the selected program
+      *
+      * @param[in] program_name Program name.
+     */
+    std::string get_program_source(const std::string &program_name);
     /** Sets the CL context used to create programs.
      *
      * @note Setting the context also resets the device to the
@@ -195,14 +208,19 @@ 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);
+        _device             = std::move(device);
+        _max_workgroup_size = 0;
+        max_local_workgroup_size();
     };
     /** Creates a kernel from the kernel library.
      *
@@ -220,6 +238,15 @@ class CLKernelLibrary
      *
      */
     void load_binary();
+    /** Find the maximum number of local work items in a workgroup can be supported by the device
+     *
+     */
+    size_t max_local_workgroup_size();
+
+    /** Return the default NDRange that is suitable for the device.
+     *
+     */
+    cl::NDRange default_ndrange();
 
 private:
     /** Load program and its dependencies.
@@ -243,6 +270,7 @@ 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__ */