Add support for OpenCL ICD Loader

BUILD.gn

@@ -345,6 +345,9 @@ if (angle_has_build) {
 
 angle_static_library("angle_common") {
   sources = libangle_common_sources
+  if (angle_enable_cl) {
+    sources += libangle_common_cl_sources
+  }
 
   configs += [
     ":angle_common_config",

gni/angle.gni

@@ -140,6 +140,11 @@ declare_args() {
   angle_enable_cl_passthrough = angle_enable_cl
 }
 
+# OpenCL is not supported on Windows UWP, because the CL headers include DXD9, which is not
+# supported by UWP. A workaround might be possible if CL support on UWP is required.
+assert(!angle_is_winuwp || !angle_enable_cl,
+       "OpenCL is not supported on Windows UWP")
+
 if (!angle_enable_cl) {
   angle_enable_cl_passthrough = false
 }

include/angle_cl.h

@@ -3,9 +3,7 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-// angle_cl.h:
-//   Includes all necessary CL headers and definitions for ANGLE.
-//
+// angle_cl.h: Includes all necessary CL headers and definitions for ANGLE.
 
 #ifndef ANGLECL_H_
 #define ANGLECL_H_
@@ -18,6 +16,91 @@
 #define CL_USE_DEPRECATED_OPENCL_2_1_APIS
 #define CL_USE_DEPRECATED_OPENCL_2_2_APIS
 
-#include "CL/opencl.h"
+#include "CL/cl_icd.h"
+
+#include <cstddef>
+#include <type_traits>
+
+namespace cl
+{
+
+template <typename CLObjectType>
+struct Dispatch
+{
+    constexpr Dispatch(const cl_icd_dispatch &dispatch) : mDispatch(&dispatch)
+    {
+        static_assert(
+            std::is_standard_layout<CLObjectType>::value && offsetof(CLObjectType, mDispatch) == 0u,
+            "Not ICD compatible");
+    }
+    ~Dispatch() = default;
+
+    constexpr const cl_icd_dispatch &getDispatch() { return *mDispatch; }
+
+  private:
+    // This has to be the first member to be OpenCL ICD compatible
+    const cl_icd_dispatch *const mDispatch;
+};
+
+}  // namespace cl
+
+struct _cl_platform_id : public cl::Dispatch<_cl_platform_id>
+{
+    constexpr _cl_platform_id(const cl_icd_dispatch &dispatch)
+        : cl::Dispatch<_cl_platform_id>(dispatch)
+    {}
+    ~_cl_platform_id() = default;
+};
+
+struct _cl_device_id : public cl::Dispatch<_cl_device_id>
+{
+    constexpr _cl_device_id(const cl_icd_dispatch &dispatch) : cl::Dispatch<_cl_device_id>(dispatch)
+    {}
+    ~_cl_device_id() = default;
+};
+
+struct _cl_context : public cl::Dispatch<_cl_context>
+{
+    constexpr _cl_context(const cl_icd_dispatch &dispatch) : cl::Dispatch<_cl_context>(dispatch) {}
+    ~_cl_context() = default;
+};
+
+struct _cl_command_queue : public cl::Dispatch<_cl_command_queue>
+{
+    constexpr _cl_command_queue(const cl_icd_dispatch &dispatch)
+        : cl::Dispatch<_cl_command_queue>(dispatch)
+    {}
+    ~_cl_command_queue() = default;
+};
+
+struct _cl_mem : public cl::Dispatch<_cl_mem>
+{
+    constexpr _cl_mem(const cl_icd_dispatch &dispatch) : cl::Dispatch<_cl_mem>(dispatch) {}
+    ~_cl_mem() = default;
+};
+
+struct _cl_program : public cl::Dispatch<_cl_program>
+{
+    constexpr _cl_program(const cl_icd_dispatch &dispatch) : cl::Dispatch<_cl_program>(dispatch) {}
+    ~_cl_program() = default;
+};
+
+struct _cl_kernel : public cl::Dispatch<_cl_kernel>
+{
+    constexpr _cl_kernel(const cl_icd_dispatch &dispatch) : cl::Dispatch<_cl_kernel>(dispatch) {}
+    ~_cl_kernel() = default;
+};
+
+struct _cl_event : public cl::Dispatch<_cl_event>
+{
+    constexpr _cl_event(const cl_icd_dispatch &dispatch) : cl::Dispatch<_cl_event>(dispatch) {}
+    ~_cl_event() = default;
+};
+
+struct _cl_sampler : public cl::Dispatch<_cl_sampler>
+{
+    constexpr _cl_sampler(const cl_icd_dispatch &dispatch) : cl::Dispatch<_cl_sampler>(dispatch) {}
+    ~_cl_sampler() = default;
+};
 
 #endif  // ANGLECL_H_

scripts/code_generation_hashes/GL_EGL_WGL_loader.json

@@ -6,7 +6,7 @@
   "scripts/egl_angle_ext.xml":
     "5bcc01462b355d933cf3ada15198fb68",
   "scripts/generate_loader.py":
-    "e24af68efd9f8149b53225e4b07a20aa",
+    "5592f86f5841d24afa660c67330426d6",
   "scripts/gl.xml":
     "2a73a58a7e26d8676a2c0af6d528cae6",
   "scripts/gl_angle_ext.xml":
@@ -19,8 +19,6 @@
     "3740eb7bd4928f17c4239ab294930469",
   "src/libEGL/egl_loader_autogen.h":
     "9cbf4d491497058a32642865eb032276",
-  "src/libOpenCL/cl_loader_autogen.cpp":
-    "1251dfd7f095459ff076abb02a5bbf79",
   "src/tests/restricted_traces/trace_egl_loader_autogen.cpp":
     "ab1ce9e72e1e248b13302349f2228a89",
   "src/tests/restricted_traces/trace_egl_loader_autogen.h":

scripts/code_generation_hashes/GL_EGL_entry_points.json

@@ -10,7 +10,7 @@
   "scripts/entry_point_packed_gl_enums.json":
     "4f7b43863a5e61991bba4010db463679",
   "scripts/generate_entry_points.py":
-    "ce49f151aac3d4a163686d7e45bcb03e",
+    "ef46ad67466c14b792ca7e9ac1cea8d7",
   "scripts/gl.xml":
     "2a73a58a7e26d8676a2c0af6d528cae6",
   "scripts/gl_angle_ext.xml":
@@ -130,9 +130,9 @@
   "src/libGLESv2/egl_stubs_autogen.h":
     "6439daa350c1663e71dd0af37dcc91df",
   "src/libGLESv2/entry_points_cl_autogen.cpp":
-    "2c43ce51701c4499003638fa30de340e",
+    "2de1fb4947c632a11803a459b3fc2b25",
   "src/libGLESv2/entry_points_cl_autogen.h":
-    "129fa7936c8bc6a20de9269da0c3c7d3",
+    "c53bd6b7025be98440ae9ca90161d550",
   "src/libGLESv2/entry_points_egl_autogen.cpp":
     "e7b708af1c8de435532058eb165d421e",
   "src/libGLESv2/entry_points_egl_autogen.h":
@@ -174,5 +174,5 @@
   "src/libGLESv2/libGLESv2_with_capture_autogen.def":
     "6b895f17f1a745f626a7534f14971fcd",
   "src/libOpenCL/libOpenCL_autogen.cpp":
-    "be80b03e4d121921f4a27aba9ad16aa1"
+    "06d55739c92ece6b669074451634b106"
 }
\ No newline at end of file

scripts/generate_loader.py

@@ -102,32 +102,6 @@ def write_source(data_source_name,
         out.close()
 
 
-def gen_libcl_loader():
-
-    xml = registry_xml.RegistryXML("cl.xml")
-    for major_version, minor_version in registry_xml.CL_VERSIONS:
-        name_prefix = "CL_VERSION_"
-        annotation = "%d_%d" % (major_version, minor_version)
-        feature_name = "%s%s" % (name_prefix, annotation)
-        xml.AddCommands(feature_name, annotation)
-    all_cmds = xml.all_cmd_names.get_all_commands()
-
-    path = os.path.join("..", "src", "libOpenCL")
-    source_path = registry_xml.path_to(path, "cl_loader_autogen.cpp")
-
-    with open(source_path, "w") as out:
-        setter = "    cl_loader.%s = reinterpret_cast<cl_api_%s>(loadProc(\"CL_%s\"));"
-        setters = [setter % (cmd, cmd, cmd[2:]) for cmd in all_cmds]
-
-        loader_source = template_cl_loader_cpp.format(
-            script_name=os.path.basename(sys.argv[0]),
-            data_source_name="cl.xml",
-            set_pointers="\n".join(setters))
-
-        out.write(loader_source)
-        out.close()
-
-
 def gen_libegl_loader():
 
     data_source_name = "egl.xml and egl_angle_ext.xml"
@@ -294,7 +268,6 @@ def main():
     if len(sys.argv) > 1:
         inputs = registry_xml.xml_inputs
         outputs = [
-            '../src/libOpenCL/cl_loader_autogen.cpp',
             '../src/libEGL/egl_loader_autogen.cpp',
             '../src/libEGL/egl_loader_autogen.h',
             '../util/egl_loader_autogen.cpp',
@@ -318,7 +291,6 @@ def main():
             return 1
         return 0
 
-    gen_libcl_loader()
     gen_libegl_loader()
     gen_util_gles_and_egl_loaders()
     gen_util_wgl_loader()
@@ -418,28 +390,5 @@ void {load_fn_name}(LoadProc loadProc)
 }}  // namespace angle
 """
 
-template_cl_loader_cpp = """// GENERATED FILE - DO NOT EDIT.
-// Generated by {script_name} using data from {data_source_name}.
-//
-// Copyright 2021 The ANGLE Project Authors. All rights reserved.
-// Use of this source code is governed by a BSD-style license that can be
-// found in the LICENSE file.
-//
-// cl_loader_autogen.cpp:
-//   Simple CL function loader.
-
-#include "cl_loader.h"
-
-cl_icd_dispatch cl_loader;
-
-namespace angle
-{{
-void LoadCL(LoadProc loadProc)
-{{
-{set_pointers}
-}}
-}}  // namespace angle
-"""
-
 if __name__ == '__main__':
     sys.exit(main())

src/libANGLE/CLCommandQueue.cpp

@@ -3,12 +3,13 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLCommandQueue.cpp: Implements the cl::CommandQueue class.
 
 #include "libANGLE/CLCommandQueue.h"
 
 namespace cl
 {
-// TODO
+
+CommandQueue::CommandQueue(const cl_icd_dispatch &dispatch) : _cl_command_queue(dispatch) {}
+
 }  // namespace cl

src/libANGLE/CLCommandQueue.h

@@ -3,21 +3,22 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLCommandQueue.h: Defines the cl::CommandQueue class, which can be used to queue a set of OpenCL
 // operations.
 
 #ifndef LIBANGLE_CLCOMMANDQUEUE_H_
 #define LIBANGLE_CLCOMMANDQUEUE_H_
 
-#include "libANGLE/CLtypes.h"
+#include "libANGLE/CLObject.h"
 
 namespace cl
 {
-class CommandQueue final
+
+class CommandQueue final : public _cl_command_queue, public Object
 {
   public:
-    using IsCLObjectType = std::true_type;
+    CommandQueue(const cl_icd_dispatch &dispatch);
+    ~CommandQueue() = default;
 };
 
 }  // namespace cl

src/libANGLE/CLContext.cpp

@@ -3,12 +3,13 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLContext.cpp: Implements the cl::Context class.
 
 #include "libANGLE/CLContext.h"
 
 namespace cl
 {
-// TODO
+
+Context::Context(const cl_icd_dispatch &dispatch) : _cl_context(dispatch) {}
+
 }  // namespace cl

src/libANGLE/CLContext.h

@@ -3,21 +3,22 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLContext.h: Defines the cl::Context class, which manages OpenCL objects such as command-queues,
 // memory, program and kernel objects and for executing kernels on one or more devices.
 
 #ifndef LIBANGLE_CLCONTEXT_H_
 #define LIBANGLE_CLCONTEXT_H_
 
-#include "libANGLE/CLtypes.h"
+#include "libANGLE/CLObject.h"
 
 namespace cl
 {
-class Context final
+
+class Context final : public _cl_context, public Object
 {
   public:
-    using IsCLObjectType = std::true_type;
+    Context(const cl_icd_dispatch &dispatch);
+    ~Context() = default;
 };
 
 }  // namespace cl

src/libANGLE/CLDevice.cpp

@@ -3,12 +3,13 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLDevice.cpp: Implements the cl::Device class.
 
 #include "libANGLE/CLDevice.h"
 
 namespace cl
 {
-// TODO
+
+Device::Device(const cl_icd_dispatch &dispatch) : _cl_device_id(dispatch) {}
+
 }  // namespace cl

src/libANGLE/CLDevice.h

@@ -3,21 +3,22 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLDevice.h: Defines the cl::Device class, which provides information about OpenCL device
 // configurations.
 
 #ifndef LIBANGLE_CLDEVICE_H_
 #define LIBANGLE_CLDEVICE_H_
 
-#include "libANGLE/CLtypes.h"
+#include "libANGLE/CLObject.h"
 
 namespace cl
 {
-class Device final
+
+class Device final : public _cl_device_id, public Object
 {
   public:
-    using IsCLObjectType = std::true_type;
+    Device(const cl_icd_dispatch &dispatch);
+    ~Device() = default;
 };
 
 }  // namespace cl

src/libANGLE/CLEvent.cpp

@@ -3,12 +3,13 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLEvent.cpp: Implements the cl::Event class.
 
 #include "libANGLE/CLEvent.h"
 
 namespace cl
 {
-// TODO
+
+Event::Event(const cl_icd_dispatch &dispatch) : _cl_event(dispatch) {}
+
 }  // namespace cl

src/libANGLE/CLEvent.h

@@ -3,21 +3,22 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLEvent.h: Defines the cl::Event class, which can be used to track the execution status of an
 // OpenCL command.
 
 #ifndef LIBANGLE_CLEVENT_H_
 #define LIBANGLE_CLEVENT_H_
 
-#include "libANGLE/CLtypes.h"
+#include "libANGLE/CLObject.h"
 
 namespace cl
 {
-class Event final
+
+class Event final : public _cl_event, public Object
 {
   public:
-    using IsCLObjectType = std::true_type;
+    Event(const cl_icd_dispatch &dispatch);
+    ~Event() = default;
 };
 
 }  // namespace cl

src/libANGLE/CLKernel.cpp

@@ -3,12 +3,13 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLKernel.cpp: Implements the cl::Kernel class.
 
 #include "libANGLE/CLKernel.h"
 
 namespace cl
 {
-// TODO
+
+Kernel::Kernel(const cl_icd_dispatch &dispatch) : _cl_kernel(dispatch) {}
+
 }  // namespace cl

src/libANGLE/CLKernel.h

@@ -3,20 +3,21 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLKernel.h: Defines the cl::Kernel class, which is a function declared in an OpenCL program.
 
 #ifndef LIBANGLE_CLKERNEL_H_
 #define LIBANGLE_CLKERNEL_H_
 
-#include "libANGLE/CLtypes.h"
+#include "libANGLE/CLObject.h"
 
 namespace cl
 {
-class Kernel final
+
+class Kernel final : public _cl_kernel, public Object
 {
   public:
-    using IsCLObjectType = std::true_type;
+    Kernel(const cl_icd_dispatch &dispatch);
+    ~Kernel() = default;
 };
 
 }  // namespace cl

src/libANGLE/CLMemory.cpp

@@ -3,12 +3,13 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLMemory.cpp: Implements the cl::Memory class.
 
 #include "libANGLE/CLMemory.h"
 
 namespace cl
 {
-// TODO
+
+Memory::Memory(const cl_icd_dispatch &dispatch) : _cl_mem(dispatch) {}
+
 }  // namespace cl

src/libANGLE/CLMemory.h

@@ -3,21 +3,22 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLMemory.h: Defines the cl::Memory class, which is a memory object and represents OpenCL objects
 // such as buffers, images and pipes.
 
 #ifndef LIBANGLE_CLMEMORY_H_
 #define LIBANGLE_CLMEMORY_H_
 
-#include "libANGLE/CLtypes.h"
+#include "libANGLE/CLObject.h"
 
 namespace cl
 {
-class Memory final
+
+class Memory final : public _cl_mem, public Object
 {
   public:
-    using IsCLObjectType = std::true_type;
+    Memory(const cl_icd_dispatch &dispatch);
+    ~Memory() = default;
 };
 
 }  // namespace cl

src/libANGLE/CLObject.h

+//
+// Copyright 2021 The ANGLE Project Authors. All rights reserved.
+// Use of this source code is governed by a BSD-style license that can be
+// found in the LICENSE file.
+//
+// CLObject.h: Defines the cl::Object class, which is the base class of all ANGLE CL objects.
+
+#ifndef LIBANGLE_CLOBJECT_H_
+#define LIBANGLE_CLOBJECT_H_
+
+#include "libANGLE/CLtypes.h"
+
+namespace cl
+{
+
+class Object
+{
+  public:
+    constexpr Object() {}
+    ~Object() = default;
+};
+
+}  // namespace cl
+
+#endif  // LIBANGLE_CLCONTEXT_H_

src/libANGLE/CLPlatform.cpp

@@ -3,12 +3,13 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLPlatform.cpp: Implements the cl::Platform class.
 
 #include "libANGLE/CLPlatform.h"
 
 namespace cl
 {
-// TODO
+
+Platform::Platform(const cl_icd_dispatch &dispatch) : _cl_platform_id(dispatch) {}
+
 }  // namespace cl

src/libANGLE/CLPlatform.h

@@ -3,21 +3,22 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLPlatform.h: Defines the cl::Platform class, which provides information about platform-specific
 // OpenCL features.
 
 #ifndef LIBANGLE_CLPLATFORM_H_
 #define LIBANGLE_CLPLATFORM_H_
 
-#include "libANGLE/CLtypes.h"
+#include "libANGLE/CLObject.h"
 
 namespace cl
 {
-class Platform final
+
+class Platform final : public _cl_platform_id, public Object
 {
   public:
-    using IsCLObjectType = std::true_type;
+    Platform(const cl_icd_dispatch &dispatch);
+    ~Platform() = default;
 };
 
 }  // namespace cl

src/libANGLE/CLProgram.cpp

@@ -3,12 +3,13 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLProgram.cpp: Implements the cl::Program class.
 
 #include "libANGLE/CLProgram.h"
 
 namespace cl
 {
-// TODO
+
+Program::Program(const cl_icd_dispatch &dispatch) : _cl_program(dispatch) {}
+
 }  // namespace cl

src/libANGLE/CLProgram.h

@@ -3,20 +3,21 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLProgram.h: Defines the cl::Program class, which consists of a set of OpenCL kernels.
 
 #ifndef LIBANGLE_CLPROGRAM_H_
 #define LIBANGLE_CLPROGRAM_H_
 
-#include "libANGLE/CLtypes.h"
+#include "libANGLE/CLObject.h"
 
 namespace cl
 {
-class Program final
+
+class Program final : public _cl_program, public Object
 {
   public:
-    using IsCLObjectType = std::true_type;
+    Program(const cl_icd_dispatch &dispatch);
+    ~Program() = default;
 };
 
 }  // namespace cl

src/libANGLE/CLSampler.cpp

@@ -3,12 +3,13 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLSampler.cpp: Implements the cl::Sampler class.
 
 #include "libANGLE/CLSampler.h"
 
 namespace cl
 {
-// TODO
+
+Sampler::Sampler(const cl_icd_dispatch &dispatch) : _cl_sampler(dispatch) {}
+
 }  // namespace cl

src/libANGLE/CLSampler.h

@@ -3,20 +3,21 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLSampler.h: Defines the cl::Sampler class, which describes how to sample an OpenCL Image.
 
 #ifndef LIBANGLE_CLSAMPLER_H_
 #define LIBANGLE_CLSAMPLER_H_
 
-#include "libANGLE/CLtypes.h"
+#include "libANGLE/CLObject.h"
 
 namespace cl
 {
-class Sampler final
+
+class Sampler final : public _cl_sampler, public Object
 {
   public:
-    using IsCLObjectType = std::true_type;
+    Sampler(const cl_icd_dispatch &dispatch);
+    ~Sampler() = default;
 };
 
 }  // namespace cl

src/libANGLE/CLtypes.h

@@ -3,7 +3,6 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-
 // CLtypes.h: Defines common types for the OpenCL support in ANGLE.
 
 #ifndef LIBANGLE_CLTYPES_H_
@@ -11,7 +10,7 @@
 
 #include "angle_cl.h"
 
-#include <type_traits>
+#include <memory>
 
 namespace cl
 {
@@ -21,6 +20,7 @@ class Device;
 class Event;
 class Kernel;
 class Memory;
+class Object;
 class Platform;
 class Program;
 class Sampler;

src/libGLESv2.gni

@@ -11,8 +11,6 @@ libangle_common_sources = [
   "src/common/MemoryBuffer.cpp",
   "src/common/MemoryBuffer.h",
   "src/common/Optional.h",
-  "src/common/PackedCLEnums_autogen.cpp",
-  "src/common/PackedCLEnums_autogen.h",
   "src/common/PackedEGLEnums_autogen.cpp",
   "src/common/PackedEGLEnums_autogen.h",
   "src/common/PackedEnums.cpp",
@@ -67,6 +65,11 @@ libangle_common_sources = [
   "src/common/vector_utils.h",
 ]
 
+libangle_common_cl_sources = [
+  "src/common/PackedCLEnums_autogen.cpp",
+  "src/common/PackedCLEnums_autogen.h",
+]
+
 xxhash_sources = [
   "src/common/third_party/xxhash/xxhash.c",
   "src/common/third_party/xxhash/xxhash.h",
@@ -456,6 +459,7 @@ libangle_cl_headers = [
   "src/libANGLE/CLEvent.h",
   "src/libANGLE/CLKernel.h",
   "src/libANGLE/CLMemory.h",
+  "src/libANGLE/CLObject.h",
   "src/libANGLE/CLPlatform.h",
   "src/libANGLE/CLProgram.h",
   "src/libANGLE/CLSampler.h",
@@ -577,6 +581,8 @@ libglesv2_sources = [
 ]
 
 libglesv2_cl_sources = [
+  "src/libGLESv2/cl_dispatch_table.cpp",
+  "src/libGLESv2/cl_dispatch_table.h",
   "src/libGLESv2/cl_stubs.cpp",
   "src/libGLESv2/cl_stubs_autogen.h",
   "src/libGLESv2/entry_points_cl_autogen.cpp",

src/libGLESv2/cl_dispatch_table.cpp

+//
+// Copyright 2021 The ANGLE Project Authors. All rights reserved.
+// Use of this source code is governed by a BSD-style license that can be
+// found in the LICENSE file.
+//
+// cl_dispatch_table_autogen.cpp: Dispatch table for CL ICD Loader.
+
+#include "libGLESv2/cl_dispatch_table.h"
+
+#include "libGLESv2/entry_points_cl_autogen.h"
+
+// clang-format off
+
+// The correct order is required as defined in 'include/CL/cl_icd.h'.
+cl_icd_dispatch gCLIcdDispatchTable = {
+
+    // OpenCL 1.0
+    cl::clGetPlatformIDs,
+    cl::clGetPlatformInfo,
+    cl::clGetDeviceIDs,
+    cl::clGetDeviceInfo,
+    cl::clCreateContext,
+    cl::clCreateContextFromType,
+    cl::clRetainContext,
+    cl::clReleaseContext,
+    cl::clGetContextInfo,
+    cl::clCreateCommandQueue,
+    cl::clRetainCommandQueue,
+    cl::clReleaseCommandQueue,
+    cl::clGetCommandQueueInfo,
+    cl::clSetCommandQueueProperty,
+    cl::clCreateBuffer,
+    cl::clCreateImage2D,
+    cl::clCreateImage3D,
+    cl::clRetainMemObject,
+    cl::clReleaseMemObject,
+    cl::clGetSupportedImageFormats,
+    cl::clGetMemObjectInfo,
+    cl::clGetImageInfo,
+    cl::clCreateSampler,
+    cl::clRetainSampler,
+    cl::clReleaseSampler,
+    cl::clGetSamplerInfo,
+    cl::clCreateProgramWithSource,
+    cl::clCreateProgramWithBinary,
+    cl::clRetainProgram,
+    cl::clReleaseProgram,
+    cl::clBuildProgram,
+    cl::clUnloadCompiler,
+    cl::clGetProgramInfo,
+    cl::clGetProgramBuildInfo,
+    cl::clCreateKernel,
+    cl::clCreateKernelsInProgram,
+    cl::clRetainKernel,
+    cl::clReleaseKernel,
+    cl::clSetKernelArg,
+    cl::clGetKernelInfo,
+    cl::clGetKernelWorkGroupInfo,
+    cl::clWaitForEvents,
+    cl::clGetEventInfo,
+    cl::clRetainEvent,
+    cl::clReleaseEvent,
+    cl::clGetEventProfilingInfo,
+    cl::clFlush,
+    cl::clFinish,
+    cl::clEnqueueReadBuffer,
+    cl::clEnqueueWriteBuffer,
+    cl::clEnqueueCopyBuffer,
+    cl::clEnqueueReadImage,
+    cl::clEnqueueWriteImage,
+    cl::clEnqueueCopyImage,
+    cl::clEnqueueCopyImageToBuffer,
+    cl::clEnqueueCopyBufferToImage,
+    cl::clEnqueueMapBuffer,
+    cl::clEnqueueMapImage,
+    cl::clEnqueueUnmapMemObject,
+    cl::clEnqueueNDRangeKernel,
+    cl::clEnqueueTask,
+    cl::clEnqueueNativeKernel,
+    cl::clEnqueueMarker,
+    cl::clEnqueueWaitForEvents,
+    cl::clEnqueueBarrier,
+    cl::clGetExtensionFunctionAddress,
+    nullptr, // clCreateFromGLBuffer,
+    nullptr, // clCreateFromGLTexture2D,
+    nullptr, // clCreateFromGLTexture3D,
+    nullptr, // clCreateFromGLRenderbuffer,
+    nullptr, // clGetGLObjectInfo,
+    nullptr, // clGetGLTextureInfo,
+    nullptr, // clEnqueueAcquireGLObjects,
+    nullptr, // clEnqueueReleaseGLObjects,
+    nullptr, // clGetGLContextInfoKHR,
+
+    // cl_khr_d3d10_sharing
+    nullptr, // clGetDeviceIDsFromD3D10KHR,
+    nullptr, // clCreateFromD3D10BufferKHR,
+    nullptr, // clCreateFromD3D10Texture2DKHR,
+    nullptr, // clCreateFromD3D10Texture3DKHR,
+    nullptr, // clEnqueueAcquireD3D10ObjectsKHR,
+    nullptr, // clEnqueueReleaseD3D10ObjectsKHR,
+
+    // OpenCL 1.1
+    cl::clSetEventCallback,
+    cl::clCreateSubBuffer,
+    cl::clSetMemObjectDestructorCallback,
+    cl::clCreateUserEvent,
+    cl::clSetUserEventStatus,
+    cl::clEnqueueReadBufferRect,
+    cl::clEnqueueWriteBufferRect,
+    cl::clEnqueueCopyBufferRect,
+
+    // cl_ext_device_fission
+    nullptr, // clCreateSubDevicesEXT,
+    nullptr, // clRetainDeviceEXT,
+    nullptr, // clReleaseDeviceEXT,
+
+    // cl_khr_gl_event
+    nullptr, // clCreateEventFromGLsyncKHR,
+
+    // OpenCL 1.2
+    cl::clCreateSubDevices,
+    cl::clRetainDevice,
+    cl::clReleaseDevice,
+    cl::clCreateImage,
+    cl::clCreateProgramWithBuiltInKernels,
+    cl::clCompileProgram,
+    cl::clLinkProgram,
+    cl::clUnloadPlatformCompiler,
+    cl::clGetKernelArgInfo,
+    cl::clEnqueueFillBuffer,
+    cl::clEnqueueFillImage,
+    cl::clEnqueueMigrateMemObjects,
+    cl::clEnqueueMarkerWithWaitList,
+    cl::clEnqueueBarrierWithWaitList,
+    cl::clGetExtensionFunctionAddressForPlatform,
+    nullptr, // clCreateFromGLTexture,
+
+    // cl_khr_d3d11_sharing
+    nullptr, // clGetDeviceIDsFromD3D11KHR,
+    nullptr, // clCreateFromD3D11BufferKHR,
+    nullptr, // clCreateFromD3D11Texture2DKHR,
+    nullptr, // clCreateFromD3D11Texture3DKHR,
+    nullptr, // clCreateFromDX9MediaSurfaceKHR,
+    nullptr, // clEnqueueAcquireD3D11ObjectsKHR,
+    nullptr, // clEnqueueReleaseD3D11ObjectsKHR,
+
+    // cl_khr_dx9_media_sharing
+    nullptr, // clGetDeviceIDsFromDX9MediaAdapterKHR,
+    nullptr, // clEnqueueAcquireDX9MediaSurfacesKHR,
+    nullptr, // clEnqueueReleaseDX9MediaSurfacesKHR,
+
+    // cl_khr_egl_image
+    nullptr, // clCreateFromEGLImageKHR,
+    nullptr, // clEnqueueAcquireEGLObjectsKHR,
+    nullptr, // clEnqueueReleaseEGLObjectsKHR,
+
+    // cl_khr_egl_event
+    nullptr, // clCreateEventFromEGLSyncKHR,
+
+    // OpenCL 2.0
+    cl::clCreateCommandQueueWithProperties,
+    cl::clCreatePipe,
+    cl::clGetPipeInfo,
+    cl::clSVMAlloc,
+    cl::clSVMFree,
+    cl::clEnqueueSVMFree,
+    cl::clEnqueueSVMMemcpy,
+    cl::clEnqueueSVMMemFill,
+    cl::clEnqueueSVMMap,
+    cl::clEnqueueSVMUnmap,
+    cl::clCreateSamplerWithProperties,
+    cl::clSetKernelArgSVMPointer,
+    cl::clSetKernelExecInfo,
+
+    // cl_khr_sub_groups
+    nullptr, // clGetKernelSubGroupInfoKHR,
+
+    // OpenCL 2.1
+    cl::clCloneKernel,
+    cl::clCreateProgramWithIL,
+    cl::clEnqueueSVMMigrateMem,
+    cl::clGetDeviceAndHostTimer,
+    cl::clGetHostTimer,
+    cl::clGetKernelSubGroupInfo,
+    cl::clSetDefaultDeviceCommandQueue,
+
+    // OpenCL 2.2
+    cl::clSetProgramReleaseCallback,
+    cl::clSetProgramSpecializationConstant,
+
+    // OpenCL 3.0
+    cl::clCreateBufferWithProperties,
+    cl::clCreateImageWithProperties,
+    cl::clSetContextDestructorCallback
+};
+
+// clang-format on

src/libGLESv2/cl_dispatch_table.h

+//
+// Copyright 2021 The ANGLE Project Authors. All rights reserved.
+// Use of this source code is governed by a BSD-style license that can be
+// found in the LICENSE file.
+//
+// cl_dispatch_table.h: Declares dispatch table for CL ICD Loader.
+
+#ifndef LIBGLESV2_CL_DISPATCH_TABLE_H_
+#define LIBGLESV2_CL_DISPATCH_TABLE_H_
+
+#include "angle_cl.h"
+#include "export.h"
+
+extern "C" {
+
+ANGLE_EXPORT extern cl_icd_dispatch gCLIcdDispatchTable;
+
+}  // extern "C"
+
+#endif  // LIBGLESV2_CL_DISPATCH_TABLE_H_

src/libGLESv2/entry_points_cl_utils.h

@@ -23,6 +23,7 @@
 
 namespace cl
 {
+
 // First case: handling packed enums.
 template <typename PackedT, typename FromT>
 typename std::enable_if_t<std::is_enum<PackedT>::value, PackedT> PackParam(FromT from)
@@ -30,30 +31,16 @@ typename std::enable_if_t<std::is_enum<PackedT>::value, PackedT> PackParam(FromT
     return FromCLenum<PackedT>(from);
 }
 
-// Cast CL object types to ANGLE types marked with 'using IsCLObjectType = std::true_type;'
+// Cast CL object types to ANGLE CL object types
 template <typename PackedT, typename FromT>
 inline std::enable_if_t<
-    std::remove_pointer_t<std::remove_pointer_t<PackedT>>::IsCLObjectType::value,
+    std::is_base_of<cl::Object, std::remove_pointer_t<std::remove_pointer_t<PackedT>>>::value,
     PackedT>
 PackParam(FromT from)
 {
     return reinterpret_cast<PackedT>(from);
 }
 
-// First case: handling packed enums.
-template <typename UnpackedT, typename FromT>
-typename std::enable_if_t<std::is_enum<FromT>::value, UnpackedT> UnpackParam(FromT from)
-{
-    return ToCLenum(from);
-}
-
-// Cast ANGLE types marked with 'using IsCLObjectType = std::true_type;' to CL object types
-template <typename UnpackedT, typename FromT>
-inline typename std::enable_if_t<std::remove_pointer_t<FromT>::IsCLObjectType::value, UnpackedT>
-UnpackParam(FromT from)
-{
-    return reinterpret_cast<UnpackedT>(from);
-}
 }  // namespace cl
 
 #endif  // LIBGLESV2_ENTRY_POINTS_CL_UTILS_H_

src/libOpenCL/BUILD.gn

@@ -8,12 +8,12 @@ import("../../gni/angle.gni")
 
 assert(angle_enable_cl)
 
-angle_shared_library("OpenCL") {
+angle_shared_library("OpenCL_ANGLE") {
   defines = [ "LIBCL_IMPLEMENTATION" ]
 
   sources = [
-    "cl_loader.h",
-    "cl_loader_autogen.cpp",
+    "dispatch.cpp",
+    "dispatch.h",
     "libOpenCL_autogen.cpp",
   ]
 
@@ -28,6 +28,29 @@ angle_shared_library("OpenCL") {
   ]
 }
 
+angle_shared_library("OpenCL_ICD_ANGLE") {
+  defines = [ "LIBCL_IMPLEMENTATION" ]
+
+  sources = [
+    "dispatch.cpp",
+    "dispatch.h",
+    "libOpenCL_ICD.cpp",
+  ]
+
+  configs += [
+    "$angle_root:debug_annotations_config",
+    "$angle_root:library_name_config",
+  ]
+
+  deps = [
+    "$angle_root:angle_common",
+    "$angle_root:cl_includes",
+  ]
+}
+
 group("angle_cl") {
-  data_deps = [ ":OpenCL" ]
+  data_deps = [
+    ":OpenCL_ANGLE",
+    ":OpenCL_ICD_ANGLE",
+  ]
 }

src/libOpenCL/dispatch.cpp

+//
+// Copyright 2021 The ANGLE Project Authors. All rights reserved.
+// Use of this source code is governed by a BSD-style license that can be
+// found in the LICENSE file.
+//
+// dispatch.cpp: Implements a function to fetch the ANGLE OpenCL dispatch table.
+
+#include "libOpenCL/dispatch.h"
+
+#include "anglebase/no_destructor.h"
+#include "common/system_utils.h"
+
+#include <iostream>
+#include <memory>
+
+namespace cl
+{
+
+namespace
+{
+std::unique_ptr<angle::Library> &EntryPointsLib()
+{
+    static angle::base::NoDestructor<std::unique_ptr<angle::Library>> sEntryPointsLib;
+    return *sEntryPointsLib;
+}
+}  // anonymous namespace
+
+cl_icd_dispatch &GetDispatch()
+{
+    static cl_icd_dispatch *sDispatch = nullptr;
+
+    if (sDispatch == nullptr)
+    {
+        EntryPointsLib().reset(
+            angle::OpenSharedLibrary(ANGLE_GLESV2_LIBRARY_NAME, angle::SearchType::ApplicationDir));
+        if (EntryPointsLib())
+        {
+            sDispatch = reinterpret_cast<cl_icd_dispatch *>(
+                EntryPointsLib()->getSymbol("gCLIcdDispatchTable"));
+            if (sDispatch == nullptr)
+            {
+                std::cerr << "Error loading CL dispatch table." << std::endl;
+            }
+        }
+        else
+        {
+            std::cerr << "Error opening GLESv2 library." << std::endl;
+        }
+    }
+
+    return *sDispatch;
+}
+
+}  // namespace cl

src/libOpenCL/cl_loader.h → src/libOpenCL/dispatch.h

@@ -3,11 +3,10 @@
 // Use of this source code is governed by a BSD-style license that can be
 // found in the LICENSE file.
 //
-// cl_loader.h:
-//   Simple CL function loader.
+// dispatch.h: Declares a function to fetch the ANGLE OpenCL dispatch table.
 
-#ifndef LIBCL_CL_LOADER_H_
-#define LIBCL_CL_LOADER_H_
+#ifndef LIBOPENCL_DISPATCH_H_
+#define LIBOPENCL_DISPATCH_H_
 
 #include "export.h"
 
@@ -16,16 +15,11 @@
 #endif
 #include "angle_cl.h"
 
-// 'angle_cl.h' has to be included before this to enable CL defines
-#include "CL/cl_icd.h"
+namespace cl
+{
 
-ANGLE_NO_EXPORT extern cl_icd_dispatch cl_loader;
+cl_icd_dispatch &GetDispatch();
 
-namespace angle
-{
-using GenericProc = void (*)();
-using LoadProc    = GenericProc(CL_API_CALL *)(const char *);
-ANGLE_NO_EXPORT void LoadCL(LoadProc loadProc);
-}  // namespace angle
+}  // namespace cl
 
-#endif  // LIBCL_CL_LOADER_H_
+#endif  // LIBOPENCL_DISPATCH_H_

src/libOpenCL/libOpenCL_ICD.cpp

+//
+// Copyright 2021 The ANGLE Project Authors. All rights reserved.
+// Use of this source code is governed by a BSD-style license that can be
+// found in the LICENSE file.
+//
+// libOpenCL_ICD.cpp: Implements the CL entry points required for extension cl_khr_icd.
+
+#include "libOpenCL/dispatch.h"
+
+extern "C" {
+
+cl_int CL_API_CALL clIcdGetPlatformIDsKHR(cl_uint num_entries,
+                                          cl_platform_id *platforms,
+                                          cl_uint *num_platforms)
+{
+    return cl::GetDispatch().clGetPlatformIDs(num_entries, platforms, num_platforms);
+}
+
+cl_int CL_API_CALL clGetPlatformInfo(cl_platform_id platform,
+                                     cl_platform_info param_name,
+                                     size_t param_value_size,
+                                     void *param_value,
+                                     size_t *param_value_size_ret)
+{
+    return cl::GetDispatch().clGetPlatformInfo(platform, param_name, param_value_size, param_value,
+                                               param_value_size_ret);
+}
+
+void *CL_API_CALL clGetExtensionFunctionAddress(const char *func_name)
+{
+    return cl::GetDispatch().clGetExtensionFunctionAddress(func_name);
+}
+
+}  // extern "C"