From f01f9de88c8bced98387e0c735de3c1fb27a031f Mon Sep 17 00:00:00 2001
From: steniu01 <steven.niu@arm.com>
Date: Wed, 27 Sep 2017 17:00:11 +0100
Subject: COMPMID-545 add CL printf support

Change-Id: I685a68e7bc8d2cdff19851d839f244206b3d5790
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/89391
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
---
 arm_compute/core/CL/OpenCL.h         |  2 ++
 arm_compute/runtime/CL/CLScheduler.h | 26 ++++++++++++++++++++++++++
 src/core/CL/OpenCL.cpp               | 15 +++++++++++++++
 src/core/CL/cl_kernels/helpers.h     |  3 +++
 4 files changed, 46 insertions(+)

diff --git a/arm_compute/core/CL/OpenCL.h b/arm_compute/core/CL/OpenCL.h
index 897e9368f1..151cc9b53d 100644
--- a/arm_compute/core/CL/OpenCL.h
+++ b/arm_compute/core/CL/OpenCL.h
@@ -83,6 +83,7 @@ public:
     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 *);
 
     clBuildProgram_func            clBuildProgram            = nullptr;
     clEnqueueNDRangeKernel_func    clEnqueueNDRangeKernel    = nullptr;
@@ -113,6 +114,7 @@ public:
     clGetDeviceInfo_func           clGetDeviceInfo           = nullptr;
     clGetDeviceIDs_func            clGetDeviceIDs            = nullptr;
     clRetainEvent_func             clRetainEvent             = nullptr;
+    clGetPlatformIDs_func          clGetPlatformIDs          = nullptr;
 
 private:
     std::pair<bool, bool> _loaded{ false, false };
diff --git a/arm_compute/runtime/CL/CLScheduler.h b/arm_compute/runtime/CL/CLScheduler.h
index 11affebc48..1a7befc046 100644
--- a/arm_compute/runtime/CL/CLScheduler.h
+++ b/arm_compute/runtime/CL/CLScheduler.h
@@ -32,6 +32,28 @@
 #include "arm_compute/core/Types.h"
 #include "arm_compute/runtime/CL/CLTuner.h"
 
+#if defined(ARM_COMPUTE_DEBUG_ENABLED)
+namespace
+{
+void printf_callback(const char *buffer, unsigned int len, size_t complete, void *user_data)
+{
+    printf("%.*s", len, buffer);
+}
+
+// Create a cl_context with a printf_callback and user specified buffer size.
+cl_context_properties properties[] =
+{
+    // Enable a printf callback function for this context.
+    CL_PRINTF_CALLBACK_ARM, reinterpret_cast<cl_context_properties>(printf_callback),
+    // Request a minimum printf buffer size of 4MB for devices in the
+    // context that support this extension.
+    CL_PRINTF_BUFFERSIZE_ARM, static_cast<cl_context_properties>(0x100000),
+    CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(cl::Platform::get()()),
+    0
+};
+}
+#endif /* defined(ARM_COMPUTE_DEBUG_ENABLED) */
+
 namespace arm_compute
 {
 class ICLKernel;
@@ -60,6 +82,10 @@ public:
      */
     void default_init(ICLTuner *cl_tuner = nullptr)
     {
+#if defined(ARM_COMPUTE_DEBUG_ENABLED)
+        cl::Context::setDefault(cl::Context(CL_DEVICE_TYPE_DEFAULT, properties));
+#endif // defined(ARM_COMPUTE_DEBUG_ENABLED)
+
         CLKernelLibrary::get().init("./cl_kernels/", cl::Context::getDefault(), cl::Device::getDefault());
         init(cl::Context::getDefault(), cl::CommandQueue::getDefault(), cl::Device::getDefault(), cl_tuner);
     }
diff --git a/src/core/CL/OpenCL.cpp b/src/core/CL/OpenCL.cpp
index 1d04f397fd..cc2391977b 100644
--- a/src/core/CL/OpenCL.cpp
+++ b/src/core/CL/OpenCL.cpp
@@ -100,6 +100,7 @@ bool CLSymbols::load(const std::string &library)
     clGetDeviceInfo           = reinterpret_cast<clGetDeviceInfo_func>(dlsym(handle, "clGetDeviceInfo"));
     clGetDeviceIDs            = reinterpret_cast<clGetDeviceIDs_func>(dlsym(handle, "clGetDeviceIDs"));
     clRetainEvent             = reinterpret_cast<clRetainEvent_func>(dlsym(handle, "clRetainEvent"));
+    clGetPlatformIDs          = reinterpret_cast<clGetPlatformIDs_func>(dlsym(handle, "clGetPlatformIDs"));
 
     dlclose(handle);
 
@@ -632,3 +633,17 @@ cl_int clRetainEvent(cl_event event)
         return CL_OUT_OF_RESOURCES;
     }
 }
+
+cl_int clGetPlatformIDs(cl_uint num_entries, cl_platform_id *platforms, cl_uint *num_platforms)
+{
+    arm_compute::CLSymbols::get().load_default();
+    auto func = arm_compute::CLSymbols::get().clGetPlatformIDs;
+    if(func != nullptr)
+    {
+        return func(num_entries, platforms, num_platforms);
+    }
+    else
+    {
+        return CL_OUT_OF_RESOURCES;
+    }
+}
diff --git a/src/core/CL/cl_kernels/helpers.h b/src/core/CL/cl_kernels/helpers.h
index 68af64e344..4421e74816 100644
--- a/src/core/CL/cl_kernels/helpers.h
+++ b/src/core/CL/cl_kernels/helpers.h
@@ -25,6 +25,9 @@
 #define ARM_COMPUTE_HELPER_H
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#if defined(ARM_COMPUTE_DEBUG_ENABLED)
+#pragma OPENCL EXTENSION cl_arm_printf : enable
+#endif // defined(ARM_COMPUTE_DEBUG_ENABLED)
 
 #define EXPAND(x) x
 
-- 
cgit v1.2.1