[Android] Enable WebCL on devices which only have OpenCL 1.1 support.

WebCL 1.0 is based on OpenCL 1.1, so this feature should be available as
long as the device supports OpenCL 1.1 or above. This commit eliminates
the hard dependencies on several OpenCL 1.2 APIs.

WebCLDevice::release() is deliberately removed since clReleaseDevice() is only
available since OpenCL 1.2 to release a sub device, while it is actually not
required by the WebCL 1.0 implementation.

BUG=XWALK-3975

Author: hujiajie

Source/modules/webcl/WebCL.cpp

@@ -60,9 +60,6 @@ PassRefPtr<WebCL> WebCL::create()
 WebCL::~WebCL()
 {
     releaseAll();
-
-    for (auto platform : m_platforms)
-        platform->releaseAll();
 }
 
 Vector<RefPtr<WebCLPlatform>> WebCL::getPlatforms(ExceptionState& es)

Source/modules/webcl/WebCLCommandQueue.cpp

@@ -106,7 +106,7 @@ void WebCLCommandQueue::finish(WebCLCallback* whenFinished, ExceptionState& es)
 void WebCLCommandQueue::finishCommandQueues(SyncMethod method)
 {
     if (method == ASYNC) {
-        cl_int err = clEnqueueMarkerWithWaitList(m_clCommandQueue, 0, nullptr, &m_eventForCallback);
+        cl_int err = clEnqueueMarker(m_clCommandQueue, &m_eventForCallback);
         if (err != CL_SUCCESS || !m_eventForCallback)
             return;
         WebCLCommandQueueHolder* holder = new WebCLCommandQueueHolder;
@@ -154,7 +154,7 @@ void WebCLCommandQueue::enqueueBarrier(ExceptionState& es)
         return;
     }
 
-    cl_int err = clEnqueueBarrierWithWaitList(m_clCommandQueue, 0, nullptr, nullptr);
+    cl_int err = clEnqueueBarrier(m_clCommandQueue);
     if (err != CL_SUCCESS)
         WebCLException::throwException(err, es);
 }
@@ -175,7 +175,7 @@ void WebCLCommandQueue::enqueueMarker(WebCLEvent* event, ExceptionState& es)
     if (event && !clEventId)
         return;
 
-    cl_int err = clEnqueueMarkerWithWaitList(m_clCommandQueue, 0, nullptr, clEventId);
+    cl_int err = clEnqueueMarker(m_clCommandQueue, clEventId);
     if (err != CL_SUCCESS)
         WebCLException::throwException(err, es);
 }
@@ -196,7 +196,7 @@ void WebCLCommandQueue::enqueueWaitForEvents(const Vector<RefPtr<WebCLEvent>>& e
     if (clEvents.size() != events.size())
         return;
 
-    cl_int err = clEnqueueBarrierWithWaitList(m_clCommandQueue, clEvents.size(), clEvents.data(), nullptr);
+    cl_int err = clEnqueueWaitForEvents(m_clCommandQueue, clEvents.size(), clEvents.data());
     if (err != CL_SUCCESS)
         WebCLException::throwException(err, es);
 }

Source/modules/webcl/WebCLContext.cpp

@@ -390,10 +390,9 @@ PassRefPtr<WebCLImage> WebCLContext::createImage2DBase(unsigned flags, unsigned
     imageDescriptor.setChannelOrder(channelOrder);
     imageDescriptor.setChannelType(channelType);
     cl_image_format image_format = {channelOrder, channelType};
-    cl_image_desc desc = {CL_MEM_OBJECT_IMAGE2D, static_cast<size_t>(width), static_cast<size_t>(height), 0, 0, static_cast<size_t>(rowPitch), 0, 0, 0, 0};
 
     cl_int err = CL_SUCCESS;
-    cl_mem clMemId = clCreateImage(m_clContext, flags, &image_format, &desc, data, &err);
+    cl_mem clMemId = clCreateImage2D(m_clContext, flags, &image_format, width, height, rowPitch, data, &err);
     if (err != CL_SUCCESS) {
         WebCLException::throwException(err, es);
         return nullptr;

Source/modules/webcl/WebCLDevice.cpp

@@ -17,8 +17,13 @@ namespace blink {
 
 WebCLDevice::~WebCLDevice()
 {
-    release();
-    ASSERT(!m_clDeviceId);
+    // Unlike WebCLContext / WebCLCommandQueue / WebCLProgram / ...,
+    // WebCLDevice does not need to call a release() method here:
+    // 1) OpenCL 1.1 runtime has no clReleaseDevice() or an alternative, so
+    //    there's no need to release the device.
+    // 2) The OpenCL 1.2 (or above) spec implies that clReleaseDevice() is only
+    //    meaningful for sub devices, but no sub device is created in our WebCL
+    //    1.0 implementation.
 }
 
 PassRefPtr<WebCLDevice> WebCLDevice::create(cl_device_id deviceId)
@@ -455,18 +460,6 @@ void WebCLDevice::getEnabledExtensions(HashSet<String>& extensions)
     m_extension.getEnabledExtensions(extensions);
 }
 
-void WebCLDevice::release()
-{
-    if (isReleased())
-        return;
-
-    cl_int err = clReleaseDevice(m_clDeviceId);
-    if (err != CL_SUCCESS)
-        ASSERT_NOT_REACHED();
-
-    m_clDeviceId = 0;
-}
-
 WebCLDevice::WebCLDevice(cl_device_id device, WebCLPlatform* platform)
     : m_platform(platform)
     , m_clDeviceId(device)

Source/modules/webcl/WebCLDevice.h

@@ -40,13 +40,11 @@ class WebCLDevice : public RefCounted<WebCLDevice>, public ScriptWrappable {
     unsigned getImage2DMaxHeight();
     unsigned getMaxWorkGroup();
     Vector<unsigned> getMaxWorkItem();
-    void release();
     PassRefPtr<WebCLPlatform> getPlatform() const { return m_platform; }
     cl_device_id getDeviceId() { return m_clDeviceId; }
 
 private:
     WebCLDevice(cl_device_id, WebCLPlatform* platform);
-    bool isReleased() { return !m_clDeviceId; }
 
     WebCLPlatform* m_platform;
     WebCLExtension m_extension;

Source/modules/webcl/WebCLException.cpp

@@ -215,9 +215,6 @@ void WebCLException::throwException(int& code, ExceptionState& es)
     case CL_INVALID_PROPERTY:
         es.throwWebCLException(WebCLException::INVALID_PROPERTY, WebCLException::invalidPropertyMessage);
         break;
-    case CL_INVALID_IMAGE_DESCRIPTOR:
-        es.throwWebCLException(WebCLException::INVALID_IMAGE_FORMAT_DESCRIPTOR, WebCLException::invalidImageFormatDescriptorMessage);
-        break;
     default:
         es.throwWebCLException(WebCLException::FAILURE, WebCLException::failureMessage);
         break;

Source/modules/webcl/WebCLOpenCL.cpp

@@ -43,6 +43,8 @@ cl_int (CL_API_CALL *web_clGetPlatformIDs)(cl_uint num_entries, cl_platform_id*
 
 cl_int (CL_API_CALL *web_clGetPlatformInfo)(cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret);
 
+cl_int (CL_API_CALL *web_clUnloadCompiler)(cl_platform_id platform);
+
 cl_int (CL_API_CALL *web_clUnloadPlatformCompiler)(cl_platform_id platform);
 
 /* Device APIs */
@@ -79,6 +81,8 @@ cl_int (CL_API_CALL *web_clReleaseMemObject)(cl_mem memobj);
 
 cl_int (CL_API_CALL *web_clGetImageInfo)(cl_mem image, cl_image_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret);
 
+cl_mem (CL_API_CALL *web_clCreateImage2D)(cl_context context, cl_mem_flags flags, const cl_image_format* image_format, size_t image_width, size_t image_height, size_t image_row_pitch, void* host_ptr, cl_int* errcode_ret);
+
 cl_mem (CL_API_CALL *web_clCreateImage)(cl_context context, cl_mem_flags flags, const cl_image_format* image_format, const cl_image_desc* image_desc, void* host_ptr, cl_int* errcode_ret);
 
 cl_int (CL_API_CALL *web_clGetSupportedImageFormats)(cl_context context, cl_mem_flags flags, cl_mem_object_type image_type, cl_uint num_entries, cl_image_format* image_formats, cl_uint* num_image_formats);
@@ -145,8 +149,12 @@ cl_int (CL_API_CALL *web_clEnqueueWriteImage)(cl_command_queue command_queue, cl
 
 cl_int (CL_API_CALL *web_clEnqueueCopyBuffer)(cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_buffer, size_t src_offset, size_t dst_offset, size_t size, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event);
 
+cl_int (CL_API_CALL *web_clEnqueueBarrier)(cl_command_queue command_queue);
+
 cl_int (CL_API_CALL *web_clEnqueueBarrierWithWaitList)(cl_command_queue command_queue, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event);
 
+cl_int (CL_API_CALL *web_clEnqueueMarker)(cl_command_queue command_queue, cl_event* event);
+
 cl_int (CL_API_CALL *web_clEnqueueMarkerWithWaitList)(cl_command_queue command_queue, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event);
 
 cl_int (CL_API_CALL *web_clEnqueueTask)(cl_command_queue command_queue, cl_kernel kernel, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event);
@@ -165,6 +173,8 @@ cl_int (CL_API_CALL *web_clEnqueueCopyImageToBuffer)(cl_command_queue command_qu
 
 cl_int (CL_API_CALL *web_clEnqueueCopyBufferToImage)(cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_image, size_t src_offset, const size_t* dst_origin, const size_t* region, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event);
 
+cl_int (CL_API_CALL *web_clEnqueueWaitForEvents)(cl_command_queue command_queue, cl_uint num_events, const cl_event* event_list);
+
 /* OpenCL Extention */
 cl_int (CL_API_CALL *web_clEnqueueAcquireGLObjects)(cl_command_queue command_queue, cl_uint num_objects, const cl_mem* mem_objects, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event);
 
@@ -174,13 +184,27 @@ cl_mem (CL_API_CALL *web_clCreateFromGLBuffer)(cl_context context, cl_mem_flags
 
 cl_mem (CL_API_CALL *web_clCreateFromGLRenderbuffer)(cl_context context, cl_mem_flags flags, GLuint renderbuffer, cl_int* errcode_ret);
 
+cl_mem (CL_API_CALL *web_clCreateFromGLTexture2D)(cl_context context, cl_mem_flags flags, GLenum texture_target, GLint miplevel, GLuint texture, cl_int* errcode_ret);
+
 cl_mem (CL_API_CALL *web_clCreateFromGLTexture)(cl_context context, cl_mem_flags flags, GLenum texture_target, GLint miplevel, GLuint texture, cl_int* errcode_ret);
 
 cl_int (CL_API_CALL *web_clGetGLTextureInfo)(cl_mem, cl_gl_texture_info, size_t, void *, size_t *);
 
+// These aliases are missing from WebCLOpenCL.h. Put them here for internal use only.
+#define clReleaseDevice web_clReleaseDevice
+#define clCreateImage web_clCreateImage
+#define clUnloadPlatformCompiler web_clUnloadPlatformCompiler
+#define clEnqueueMarkerWithWaitList web_clEnqueueMarkerWithWaitList
+#define clEnqueueBarrierWithWaitList web_clEnqueueBarrierWithWaitList
+#define clCreateFromGLTexture web_clCreateFromGLTexture
+
 #if defined(WTF_OS_LINUX) || OS(ANDROID)
 #define MAP_FUNC(fn)  { *(void**)(&fn) = dlsym(handle, #fn); }
 #define MAP_FUNC_OR_BAIL(fn)  { *(void**)(&fn) = dlsym(handle, #fn); if(!fn) return false; }
+// In case `fn' is not defined or deprecated in the OpenCL spec tagged by
+// `major' and `minor', map `fn' to a wrapper implemented with APIs defined
+// by this spec.
+#define MAP_FUNC_TO_WRAPPER(fn, major, minor) { *(void**)(&fn) = (void*)fn##Impl##major##minor; }
 
 static const char* DEFAULT_SO[] = LIBS;
 static const int DEFAULT_SO_LEN = SO_LEN;
@@ -198,6 +222,49 @@ static bool getCLHandle(const char** libs, int length)
 }
 #endif // defined(WTF_OS_LINUX) || OS(ANDROID)
 
+// In OpenCL 1.1 spec, no release opertion is needed for device.
+static cl_int CL_API_CALL clReleaseDeviceImpl11(cl_device_id device)
+{
+    return CL_SUCCESS;
+}
+
+static cl_mem CL_API_CALL clCreateImage2DImpl12(cl_context context, cl_mem_flags flags, const cl_image_format* format, size_t width, size_t height, size_t rowPitch, void* hostPtr, cl_int* err)
+{
+    cl_image_desc desc = {CL_MEM_OBJECT_IMAGE2D, static_cast<size_t>(width), static_cast<size_t>(height), 0, 0, static_cast<size_t>(rowPitch), 0, 0, 0, 0};
+    ASSERT(clCreateImage);
+    return clCreateImage(context, flags, format, &desc, hostPtr, err);
+}
+
+static cl_int CL_API_CALL clUnloadCompilerImpl12(cl_platform_id platform)
+{
+    ASSERT(clUnloadPlatformCompiler);
+    return clUnloadPlatformCompiler(platform);
+}
+
+static cl_int CL_API_CALL clEnqueueMarkerImpl12(cl_command_queue queue, cl_event* event)
+{
+    ASSERT(clEnqueueMarkerWithWaitList);
+    return clEnqueueMarkerWithWaitList(queue, 0, nullptr, event);
+}
+
+static cl_int CL_API_CALL clEnqueueBarrierImpl12(cl_command_queue queue)
+{
+    ASSERT(clEnqueueBarrierWithWaitList);
+    return clEnqueueBarrierWithWaitList(queue, 0, nullptr, nullptr);
+}
+
+static cl_int CL_API_CALL clEnqueueWaitForEventsImpl12(cl_command_queue queue, cl_uint numEvents, const cl_event* eventList)
+{
+    ASSERT(clEnqueueBarrierWithWaitList);
+    return clEnqueueBarrierWithWaitList(queue, numEvents, eventList, nullptr);
+}
+
+static cl_mem CL_API_CALL clCreateFromGLTexture2DImpl12(cl_context context, cl_mem_flags flags, GLenum textureTarget, GLint miplevel, GLuint texture, cl_int* err)
+{
+    ASSERT(clCreateFromGLTexture);
+    return clCreateFromGLTexture(context, flags, textureTarget, miplevel, texture, err);
+}
+
 bool init(const char** libs, int length)
 {
     const char** mLibs = (libs == 0 ? DEFAULT_SO : libs);
@@ -211,7 +278,6 @@ bool init(const char** libs, int length)
     MAP_FUNC_OR_BAIL(clCreateCommandQueue);
     MAP_FUNC_OR_BAIL(clCreateContext);
     MAP_FUNC_OR_BAIL(clCreateContextFromType);
-    MAP_FUNC_OR_BAIL(clCreateImage);
     MAP_FUNC_OR_BAIL(clCreateKernel);
     MAP_FUNC_OR_BAIL(clCreateKernelsInProgram);
     MAP_FUNC_OR_BAIL(clCreateProgramWithSource);
@@ -230,8 +296,6 @@ bool init(const char** libs, int length)
     MAP_FUNC_OR_BAIL(clEnqueueCopyImageToBuffer);
     MAP_FUNC_OR_BAIL(clEnqueueReadBufferRect);
     MAP_FUNC_OR_BAIL(clEnqueueWriteBufferRect);
-    MAP_FUNC_OR_BAIL(clEnqueueBarrierWithWaitList);
-    MAP_FUNC_OR_BAIL(clEnqueueMarkerWithWaitList);
     MAP_FUNC_OR_BAIL(clEnqueueNDRangeKernel);
     MAP_FUNC_OR_BAIL(clEnqueueTask);
 
@@ -258,7 +322,6 @@ bool init(const char** libs, int length)
 
     MAP_FUNC_OR_BAIL(clReleaseCommandQueue);
     MAP_FUNC_OR_BAIL(clReleaseContext);
-    MAP_FUNC_OR_BAIL(clReleaseDevice);
     MAP_FUNC_OR_BAIL(clReleaseEvent);
     MAP_FUNC_OR_BAIL(clReleaseKernel);
     MAP_FUNC_OR_BAIL(clReleaseMemObject);
@@ -269,7 +332,6 @@ bool init(const char** libs, int length)
     MAP_FUNC_OR_BAIL(clSetKernelArg);
     MAP_FUNC_OR_BAIL(clSetUserEventStatus);
 
-    MAP_FUNC_OR_BAIL(clUnloadPlatformCompiler);
     MAP_FUNC_OR_BAIL(clWaitForEvents);
 
     // They depends on whether OpenCL library support gl_sharing extension.
@@ -278,9 +340,49 @@ bool init(const char** libs, int length)
     MAP_FUNC(clEnqueueReleaseGLObjects);
     MAP_FUNC(clCreateFromGLBuffer);
     MAP_FUNC(clCreateFromGLRenderbuffer);
-    MAP_FUNC(clCreateFromGLTexture);
     MAP_FUNC(clGetGLTextureInfo);
 
+    // The following APIs are not available in all versions of the OpenCL
+    // spec, so wrappers may be needed if they are not exported by the OpenCL
+    // runtime library.
+    MAP_FUNC(clReleaseDevice)
+    if (!clReleaseDevice)
+        MAP_FUNC_TO_WRAPPER(clReleaseDevice, 1, 1)
+
+    MAP_FUNC(clCreateImage)
+    if (clCreateImage)
+        MAP_FUNC_TO_WRAPPER(clCreateImage2D, 1, 2)
+    else
+        MAP_FUNC_OR_BAIL(clCreateImage2D)
+
+    MAP_FUNC(clUnloadPlatformCompiler)
+    if (clUnloadPlatformCompiler)
+        MAP_FUNC_TO_WRAPPER(clUnloadCompiler, 1, 2)
+    else
+        MAP_FUNC_OR_BAIL(clUnloadCompiler)
+
+    MAP_FUNC(clEnqueueMarkerWithWaitList)
+    if (clEnqueueMarkerWithWaitList)
+        MAP_FUNC_TO_WRAPPER(clEnqueueMarker, 1, 2)
+    else
+        MAP_FUNC_OR_BAIL(clEnqueueMarker)
+
+    MAP_FUNC(clEnqueueBarrierWithWaitList)
+    if (clEnqueueBarrierWithWaitList)
+        MAP_FUNC_TO_WRAPPER(clEnqueueBarrier, 1, 2)
+    else
+        MAP_FUNC_OR_BAIL(clEnqueueBarrier)
+
+    MAP_FUNC(clEnqueueWaitForEvents)
+    if (!clEnqueueWaitForEvents)
+        MAP_FUNC_TO_WRAPPER(clEnqueueWaitForEvents, 1, 2)
+
+    MAP_FUNC(clCreateFromGLTexture)
+    if (clCreateFromGLTexture)
+        MAP_FUNC_TO_WRAPPER(clCreateFromGLTexture2D, 1, 2)
+    else
+        MAP_FUNC(clCreateFromGLTexture2D)
+
     return true;
 }