Add loader for CL pass-through back end

.gitignore

@@ -55,6 +55,7 @@
 /third_party/llvm-build
 /third_party/markupsafe
 /third_party/nasm
+/third_party/OpenCL-ICD-Loader/src
 /third_party/proguard/lib
 /third_party/protobuf
 /third_party/Python-Markdown

BUILD.gn

@@ -268,6 +268,7 @@ config("angle_uwp_env") {
 
 angle_source_set("cl_includes") {
   sources = cl_includes
+  public_configs = [ ":external_config" ]
 }
 
 angle_source_set("includes") {

DEPS

@@ -405,6 +405,11 @@ deps = {
     'condition': 'not build_with_chromium',
   },
 
+  'third_party/OpenCL-ICD-Loader/src': {
+    'url': '{chromium_git}/external/github.com/KhronosGroup/OpenCL-ICD-Loader@9b5e3849b49a1448996c8b96ba086cd774d987db',
+    'condition': 'not build_with_chromium',
+  },
+
   'third_party/proguard': {
       'packages': [
           {

gni/angle.gni

@@ -140,7 +140,7 @@ 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
+# OpenCL is not supported on Windows UWP, because the CL headers include DX9, 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")

scripts/code_generation_hashes/GL_EGL_WGL_loader.json

@@ -12,7 +12,7 @@
   "scripts/gl_angle_ext.xml":
     "08f74b35d908b7c02b45fdf45572c434",
   "scripts/registry_xml.py":
-    "47013d2963204bdbb5097f077f452f0c",
+    "a5e473c991c18c71b07ae7b7c7509f4e",
   "scripts/wgl.xml":
     "c36001431919e1c435f1215a85f7e1db",
   "src/libEGL/egl_loader_autogen.cpp":

scripts/code_generation_hashes/GL_EGL_entry_points.json

@@ -10,19 +10,19 @@
   "scripts/entry_point_packed_gl_enums.json":
     "4f7b43863a5e61991bba4010db463679",
   "scripts/generate_entry_points.py":
-    "ef46ad67466c14b792ca7e9ac1cea8d7",
+    "de7a2201b2a550bf1e890b9941e9ad57",
   "scripts/gl.xml":
     "2a73a58a7e26d8676a2c0af6d528cae6",
   "scripts/gl_angle_ext.xml":
     "08f74b35d908b7c02b45fdf45572c434",
   "scripts/registry_xml.py":
-    "47013d2963204bdbb5097f077f452f0c",
+    "a5e473c991c18c71b07ae7b7c7509f4e",
   "scripts/wgl.xml":
     "c36001431919e1c435f1215a85f7e1db",
   "src/common/entry_points_enum_autogen.cpp":
-    "f44e96c3297ed6787434ae7ee02596af",
+    "c12d59977a450452ee72ddf8d767d654",
   "src/common/entry_points_enum_autogen.h":
-    "9db537e91358777907a2bba88ee98be5",
+    "a194a0bddfd0a239d7090c47a58a0672",
   "src/libANGLE/Context_gl_1_autogen.h":
     "6be1391ee21b3754d9e9c512255d4c5d",
   "src/libANGLE/Context_gl_2_autogen.h":
@@ -74,7 +74,7 @@
   "src/libANGLE/capture/frame_capture_utils_autogen.h":
     "4405cac61b4f79b893bd40c7fc452397",
   "src/libANGLE/validationCL_autogen.h":
-    "9b030d9bb634fdd7f38b28b2dc25e0f0",
+    "7fd03346f5cb8859f45cdd380af4bc64",
   "src/libANGLE/validationEGL_autogen.h":
     "3927fa260ad183fd9193d65b3f8d82c5",
   "src/libANGLE/validationES1_autogen.h":
@@ -122,7 +122,7 @@
   "src/libGL/libGL_autogen.def":
     "2789d87b05eea9f53d52e2aff499b785",
   "src/libGLESv2/cl_stubs_autogen.h":
-    "93e4a07e582d3ee0955d1e82bc43dd9b",
+    "15437a0236c79d37f2ee4655fc99409f",
   "src/libGLESv2/egl_ext_stubs_autogen.h":
     "2ef3b8d087f2a97f7270b96077c93856",
   "src/libGLESv2/egl_get_labeled_object_data.json":
@@ -130,9 +130,9 @@
   "src/libGLESv2/egl_stubs_autogen.h":
     "6439daa350c1663e71dd0af37dcc91df",
   "src/libGLESv2/entry_points_cl_autogen.cpp":
-    "2de1fb4947c632a11803a459b3fc2b25",
+    "4e7af65ebb7f126992adcf932a5b3060",
   "src/libGLESv2/entry_points_cl_autogen.h":
-    "c53bd6b7025be98440ae9ca90161d550",
+    "dde2f94c3004874a7da995dae69da811",
   "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":
-    "06d55739c92ece6b669074451634b106"
+    "10849978c910dc1af5dd4f0c815d1581"
 }
\ No newline at end of file

scripts/code_generation_hashes/GLenum_value_to_string_map.json

@@ -6,7 +6,7 @@
   "scripts/gl_angle_ext.xml":
     "08f74b35d908b7c02b45fdf45572c434",
   "scripts/registry_xml.py":
-    "47013d2963204bdbb5097f077f452f0c",
+    "a5e473c991c18c71b07ae7b7c7509f4e",
   "src/libANGLE/capture/gl_enum_utils_autogen.cpp":
     "8a310f105f1e120203d724776c4a33fb",
   "src/libANGLE/capture/gl_enum_utils_autogen.h":

scripts/code_generation_hashes/packed_enum.json

 {
   "src/common/PackedCLEnums_autogen.cpp":
-    "399217e4821b02c00fb96788de683c55",
+    "91ef59eed1815db6947faae73a8d3e2d",
   "src/common/PackedCLEnums_autogen.h":
-    "76e2457971997ca26ff3761ac21573e4",
+    "4ca5b1eb69c7b690c402a91644d564db",
   "src/common/PackedEGLEnums_autogen.cpp":
     "1e5cf84847c9880d32d1f638b366ed08",
   "src/common/PackedEGLEnums_autogen.h":
@@ -14,7 +14,7 @@
   "src/common/gen_packed_gl_enums.py":
     "7453287e3f184427043604d172a6765b",
   "src/common/packed_cl_enums.json":
-    "a30fcf0740b85c0a32518fe4f7d1e138",
+    "b38037f250994e7d275ac2185afea820",
   "src/common/packed_egl_enums.json":
     "76e1b814421e121164d60a0d89cb16c1",
   "src/common/packed_gl_enums.json":

scripts/code_generation_hashes/proc_table.json

@@ -6,17 +6,19 @@
   "scripts/egl_angle_ext.xml":
     "5bcc01462b355d933cf3ada15198fb68",
   "scripts/gen_proc_table.py":
-    "e4b824a2f13d8b00745158b95b000dde",
+    "1d92e72e9b98865366d1e56f5e31954c",
   "scripts/gl.xml":
     "2a73a58a7e26d8676a2c0af6d528cae6",
   "scripts/gl_angle_ext.xml":
     "08f74b35d908b7c02b45fdf45572c434",
   "scripts/registry_xml.py":
-    "47013d2963204bdbb5097f077f452f0c",
+    "a5e473c991c18c71b07ae7b7c7509f4e",
   "scripts/wgl.xml":
     "c36001431919e1c435f1215a85f7e1db",
   "src/libGL/proc_table_wgl_autogen.cpp":
     "2174a66447df2d577184fbe4b96c77e9",
+  "src/libGLESv2/proc_table_cl_autogen.cpp":
+    "ed003b0f041aaaa35b67d3fe07e61f91",
   "src/libGLESv2/proc_table_egl_autogen.cpp":
     "53d085f3f46ea5d841d93e7063925ae4"
 }
\ No newline at end of file

scripts/gen_proc_table.py

@@ -12,6 +12,7 @@ import registry_xml
 
 out_file_name_gles = "../src/libGLESv2/proc_table_egl_autogen.cpp"
 out_file_name_gl = "../src/libGL/proc_table_wgl_autogen.cpp"
+out_file_name_cl = "../src/libGLESv2/proc_table_cl_autogen.cpp"
 
 # The EGL_ANGLE_explicit_context extension is generated differently from other extensions.
 # Toggle generation here.
@@ -44,6 +45,33 @@ const size_t g_numProcs = {num_procs};
 }}  // namespace {namespace}
 """
 
+template_map_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.
+//
+// proc_table:
+//   Mapping from a string entry point name to function address.
+//
+
+{includes}
+#define P(FUNC) reinterpret_cast<{cast}>(FUNC)
+
+namespace {namespace}
+{{
+
+const ProcTable &GetProcTable()
+{{
+    static angle::base::NoDestructor<ProcTable> sProcTable(
+        {{{proc_data}}});
+    return *sProcTable;
+}}
+
+}}  // namespace {namespace}
+"""
+
 includes_gles = """#include "libGLESv2/proc_table_egl.h"
 
 #include "libGLESv2/entry_points_egl_autogen.h"
@@ -67,6 +95,16 @@ includes_gl = """#include "libGL/proc_table_wgl.h"
 #include "platform/PlatformMethods.h"
 """
 
+includes_cl = """#include "libGLESv2/proc_table_cl.h"
+
+#include "libGLESv2/entry_points_cl_autogen.h"
+
+#include "anglebase/no_destructor.h"
+
+// Using fully qualified entry point identifiers to make sure that missing entry points would not
+// pick up the global declarations of OpenCL
+"""
+
 sys.path.append('../src/libANGLE/renderer')
 import angle_format
 
@@ -80,7 +118,7 @@ def main():
     # auto_script parameters.
     if len(sys.argv) > 1:
         inputs = [source for source in registry_xml.xml_inputs]
-        outputs = [out_file_name_gles, out_file_name_gl]
+        outputs = [out_file_name_gles, out_file_name_gl, out_file_name_cl]
         if sys.argv[1] == 'inputs':
             print(','.join(inputs))
         elif sys.argv[1] == 'outputs':
@@ -197,6 +235,31 @@ def main():
             num_procs=len(proc_data))
         out_file.write(output_cpp)
         out_file.close()
+
+    # libCL proc table
+    clxml = 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)
+        clxml.AddCommands(feature_name, annotation)
+
+    clxml.AddExtensionCommands(registry_xml.supported_cl_extensions, ['cl'])
+    cl_commands = clxml.all_cmd_names.get_all_commands()
+    proc_data = ['{"%s", P(::cl::%s)}' % (cmd, cmd) for cmd in cl_commands]
+
+    with open(out_file_name_cl, 'w') as out_file:
+        output_cpp = template_map_cpp.format(
+            script_name=sys.argv[0],
+            data_source_name="cl.xml",
+            includes=includes_cl,
+            cast="void *",
+            namespace="cl",
+            proc_data=",\n         ".join(proc_data))
+        out_file.write(output_cpp)
+        out_file.close()
+
     return 0
 
 

scripts/generate_entry_points.py

@@ -245,6 +245,7 @@ TEMPLATE_EGL_ENTRY_POINT_WITH_RETURN = """\
 TEMPLATE_CL_ENTRY_POINT_NO_RETURN = """\
 void CL_API_CALL cl{name}({params})
 {{
+    ANGLE_SCOPED_GLOBAL_LOCK();
     CL_EVENT({name}, "{format_params}"{comma_if_needed}{pass_params});
 
     {packed_gl_enum_conversions}
@@ -258,6 +259,7 @@ void CL_API_CALL cl{name}({params})
 TEMPLATE_CL_ENTRY_POINT_WITH_RETURN_ERROR = """\
 cl_int CL_API_CALL cl{name}({params})
 {{
+    ANGLE_SCOPED_GLOBAL_LOCK();
     CL_EVENT({name}, "{format_params}"{comma_if_needed}{pass_params});
 
     {packed_gl_enum_conversions}
@@ -271,6 +273,7 @@ cl_int CL_API_CALL cl{name}({params})
 TEMPLATE_CL_ENTRY_POINT_WITH_RETURN_POINTER = """\
 {return_type} CL_API_CALL cl{name}({params})
 {{
+    ANGLE_SCOPED_GLOBAL_LOCK();
     CL_EVENT({name}, "{format_params}"{comma_if_needed}{pass_params});
 
     {packed_gl_enum_conversions}
@@ -296,8 +299,6 @@ TEMPLATE_CL_STUBS_HEADER = """\
 
 #include "libANGLE/CLtypes.h"
 
-#include "common/PackedCLEnums_autogen.h"
-
 namespace cl
 {{
 {stubs}
@@ -962,6 +963,7 @@ LIBCL_SOURCE_INCLUDES = """\
 #include "libANGLE/validationCL_autogen.h"
 #include "libGLESv2/cl_stubs_autogen.h"
 #include "libGLESv2/entry_points_cl_utils.h"
+#include "libGLESv2/global_state.h"
 """
 
 TEMPLATE_EVENT_COMMENT = """\
@@ -2915,6 +2917,32 @@ def main():
         cl_validation_protos += [comment] + eps.validation_protos
         libcl_windows_def_exports += [win_def_comment] + get_exports(clxml.commands[version])
 
+    clxml.AddExtensionCommands(registry_xml.supported_cl_extensions, ['cl'])
+    for extension_name, ext_cmd_names in sorted(clxml.ext_data.items()):
+
+        # Extensions may have no new commands.
+        if not ext_cmd_names:
+            continue
+
+        # Detect and filter duplicate extensions.
+        eps = CLEntryPoints(clxml, ext_cmd_names)
+
+        comment = "\n// %s" % extension_name
+        win_def_comment = "\n    ; %s" % (extension_name)
+
+        cl_commands += ext_cmd_names
+
+        cl_decls += [comment] + eps.decls
+        cl_defs += [comment] + eps.defs
+        libcl_ep_defs += [comment] + eps.export_defs
+        cl_validation_protos += [comment] + eps.validation_protos
+        libcl_windows_def_exports += [win_def_comment] + get_exports(ext_cmd_names)
+
+        # Avoid writing out entry points defined by a prior extension.
+        for dupe in clxml.ext_dupes[extension_name]:
+            msg = "// %s is already defined.\n" % strip_api_prefix(dupe)
+            cl_defs.append(msg)
+
     cl_decls.append("}  // namespace cl")
     cl_defs.append("}  // namespace cl")
 

scripts/registry_xml.py

@@ -192,6 +192,11 @@ supported_egl_extensions = [
     "EGL_NV_stream_consumer_gltexture_yuv",
 ]
 
+supported_cl_extensions = [
+    "cl_khr_extended_versioning",
+    "cl_khr_icd",
+]
+
 # Strip these suffixes from Context entry point names. NV is excluded (for now).
 strip_suffixes = ["ANGLE", "EXT", "KHR", "OES", "CHROMIUM"]
 
@@ -326,6 +331,8 @@ class RegistryXML:
             return 'eglext'
         elif 'wgl' in supported:
             return 'wglext'
+        elif 'cl' in supported:
+            return 'clext'
         else:
             assert False
             return 'unknown'

src/common/PackedCLEnums_autogen.cpp

@@ -1824,6 +1824,8 @@ PlatformInfo FromCLenum<PlatformInfo>(CLenum from)
             return PlatformInfo::NumericVersion;
         case CL_PLATFORM_EXTENSIONS_WITH_VERSION:
             return PlatformInfo::ExtensionsWithVersion;
+        case CL_PLATFORM_ICD_SUFFIX_KHR:
+            return PlatformInfo::IcdSuffix;
         default:
             return PlatformInfo::InvalidEnum;
     }
@@ -1849,6 +1851,8 @@ CLenum ToCLenum(PlatformInfo from)
             return CL_PLATFORM_NUMERIC_VERSION;
         case PlatformInfo::ExtensionsWithVersion:
             return CL_PLATFORM_EXTENSIONS_WITH_VERSION;
+        case PlatformInfo::IcdSuffix:
+            return CL_PLATFORM_ICD_SUFFIX_KHR;
         default:
             UNREACHABLE();
             return 0;
@@ -1883,6 +1887,9 @@ std::ostream &operator<<(std::ostream &os, PlatformInfo value)
         case PlatformInfo::ExtensionsWithVersion:
             os << "CL_PLATFORM_EXTENSIONS_WITH_VERSION";
             break;
+        case PlatformInfo::IcdSuffix:
+            os << "CL_PLATFORM_ICD_SUFFIX_KHR";
+            break;
         default:
             os << "GL_INVALID_ENUM";
             break;

src/common/PackedCLEnums_autogen.h

@@ -403,9 +403,10 @@ enum class PlatformInfo : uint8_t
     HostTimerResolution   = 5,
     NumericVersion        = 6,
     ExtensionsWithVersion = 7,
+    IcdSuffix             = 8,
 
-    InvalidEnum = 8,
-    EnumCount   = 8,
+    InvalidEnum = 9,
+    EnumCount   = 9,
 };
 
 template <>

src/common/entry_points_enum_autogen.cpp

@@ -180,6 +180,8 @@ const char *GetEntryPointName(EntryPoint ep)
             return "clGetSamplerInfo";
         case EntryPoint::CLGetSupportedImageFormats:
             return "clGetSupportedImageFormats";
+        case EntryPoint::CLIcdGetPlatformIDsKHR:
+            return "clIcdGetPlatformIDsKHR";
         case EntryPoint::CLLinkProgram:
             return "clLinkProgram";
         case EntryPoint::CLReleaseCommandQueue:

src/common/entry_points_enum_autogen.h

@@ -96,6 +96,7 @@ enum class EntryPoint
     CLGetProgramInfo,
     CLGetSamplerInfo,
     CLGetSupportedImageFormats,
+    CLIcdGetPlatformIDsKHR,
     CLLinkProgram,
     CLReleaseCommandQueue,
     CLReleaseContext,

src/common/packed_cl_enums.json

@@ -7,7 +7,8 @@
         "Extensions": "CL_PLATFORM_EXTENSIONS",
         "HostTimerResolution": "CL_PLATFORM_HOST_TIMER_RESOLUTION",
         "NumericVersion": "CL_PLATFORM_NUMERIC_VERSION",
-        "ExtensionsWithVersion": "CL_PLATFORM_EXTENSIONS_WITH_VERSION"
+        "ExtensionsWithVersion": "CL_PLATFORM_EXTENSIONS_WITH_VERSION",
+        "IcdSuffix": "CL_PLATFORM_ICD_SUFFIX_KHR"
     },
     "DeviceInfo": {
         "Type": "CL_DEVICE_TYPE",

src/libANGLE/CLObject.h

@@ -8,7 +8,7 @@
 #ifndef LIBANGLE_CLOBJECT_H_
 #define LIBANGLE_CLOBJECT_H_
 
-#include "libANGLE/CLtypes.h"
+#include "libANGLE/renderer/CLtypes.h"
 
 namespace cl
 {

src/libANGLE/CLPlatform.cpp

@@ -10,6 +10,18 @@
 namespace cl
 {
 
-Platform::Platform(const cl_icd_dispatch &dispatch) : _cl_platform_id(dispatch) {}
+Platform::~Platform() = default;
+
+void Platform::CreatePlatform(const cl_icd_dispatch &dispatch, rx::CLPlatformImpl::Ptr &&impl)
+{
+    GetList().emplace_back(new Platform(dispatch, std::move(impl)));
+}
+
+Platform::Platform(const cl_icd_dispatch &dispatch, rx::CLPlatformImpl::Ptr &&impl)
+    : _cl_platform_id(dispatch), mImpl(std::move(impl))
+{}
+
+constexpr char Platform::kVendor[];
+constexpr char Platform::kIcdSuffix[];
 
 }  // namespace cl

src/libANGLE/CLPlatform.h

@@ -10,6 +10,12 @@
 #define LIBANGLE_CLPLATFORM_H_
 
 #include "libANGLE/CLObject.h"
+#include "libANGLE/renderer/CLPlatformImpl.h"
+
+#include "anglebase/no_destructor.h"
+
+#include <algorithm>
+#include <string>
 
 namespace cl
 {
@@ -17,10 +23,99 @@ namespace cl
 class Platform final : public _cl_platform_id, public Object
 {
   public:
-    Platform(const cl_icd_dispatch &dispatch);
-    ~Platform() = default;
+    using Ptr  = std::unique_ptr<Platform>;
+    using List = std::vector<Ptr>;
+
+    ~Platform();
+
+    const char *getProfile() const;
+    const char *getVersionString() const;
+    cl_version getVersion() const;
+    const char *getName() const;
+    const char *getExtensions() const;
+    const rx::CLPlatformImpl::ExtensionList &getExtensionsWithVersion() const;
+    cl_ulong getHostTimerResolution() const;
+
+    static void CreatePlatform(const cl_icd_dispatch &dispatch, rx::CLPlatformImpl::Ptr &&impl);
+    static const List &GetPlatforms();
+    static bool IsValid(const Platform *platform);
+    static constexpr const char *GetVendor();
+    static constexpr const char *GetIcdSuffix();
+
+  private:
+    Platform(const cl_icd_dispatch &dispatch, rx::CLPlatformImpl::Ptr &&impl);
+
+    static List &GetList();
+
+    const rx::CLPlatformImpl::Ptr mImpl;
+
+    static constexpr char kVendor[]    = "ANGLE";
+    static constexpr char kIcdSuffix[] = "ANGLE";
 };
 
+inline const char *Platform::getProfile() const
+{
+    return mImpl->getInfo().mProfile.c_str();
+}
+
+inline const char *Platform::getVersionString() const
+{
+    return mImpl->getInfo().mVersionStr.c_str();
+}
+
+inline cl_version Platform::getVersion() const
+{
+    return mImpl->getInfo().mVersion;
+}
+
+inline const char *Platform::getName() const
+{
+    return mImpl->getInfo().mName.c_str();
+}
+
+inline const char *Platform::getExtensions() const
+{
+    return mImpl->getInfo().mExtensions.c_str();
+}
+
+inline const rx::CLPlatformImpl::ExtensionList &Platform::getExtensionsWithVersion() const
+{
+    return mImpl->getInfo().mExtensionList;
+}
+
+inline cl_ulong Platform::getHostTimerResolution() const
+{
+    return mImpl->getInfo().mHostTimerRes;
+}
+
+inline Platform::List &Platform::GetList()
+{
+    static angle::base::NoDestructor<List> sList;
+    return *sList;
+}
+
+inline const Platform::List &Platform::GetPlatforms()
+{
+    return GetList();
+}
+
+inline bool Platform::IsValid(const Platform *platform)
+{
+    const List &platforms = GetPlatforms();
+    return std::find_if(platforms.cbegin(), platforms.cend(),
+                        [=](const Ptr &ptr) { return ptr.get() == platform; }) != platforms.cend();
+}
+
+constexpr const char *Platform::GetVendor()
+{
+    return kVendor;
+}
+
+constexpr const char *Platform::GetIcdSuffix()
+{
+    return kIcdSuffix;
+}
+
 }  // namespace cl
 
 #endif  // LIBANGLE_CLPLATFORM_H_

src/libANGLE/CLtypes.h

@@ -10,7 +10,10 @@
 
 #include "angle_cl.h"
 
+#include "common/PackedCLEnums_autogen.h"
+
 #include <memory>
+#include <utility>
 
 namespace cl
 {

src/libANGLE/renderer/CLContextImpl.h

@@ -9,7 +9,7 @@
 #ifndef LIBANGLE_RENDERER_CLCONTEXTIMPL_H_
 #define LIBANGLE_RENDERER_CLCONTEXTIMPL_H_
 
-#include "common/angleutils.h"
+#include "libANGLE/renderer/CLtypes.h"
 
 namespace rx
 {

src/libANGLE/renderer/CLDeviceImpl.h

@@ -9,7 +9,7 @@
 #ifndef LIBANGLE_RENDERER_CLDEVICEIMPL_H_
 #define LIBANGLE_RENDERER_CLDEVICEIMPL_H_
 
-#include "common/angleutils.h"
+#include "libANGLE/renderer/CLtypes.h"
 
 namespace rx
 {

src/libANGLE/renderer/CLPlatformImpl.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.
+//
+// CLPlatformImpl.cpp: Implements the class methods for CLPlatformImpl.
+
+#include "libANGLE/renderer/CLPlatformImpl.h"
+
+namespace rx
+{
+
+CLPlatformImpl::Info::Info() = default;
+
+CLPlatformImpl::Info::~Info() = default;
+
+CLPlatformImpl::Info::Info(Info &&) = default;
+
+CLPlatformImpl::Info &CLPlatformImpl::Info::operator=(Info &&) = default;
+
+CLPlatformImpl::Info::Info(std::string &&profile,
+                           std::string &&versionStr,
+                           cl_version version,
+                           std::string &&name,
+                           std::string &&extensions,
+                           rx::CLPlatformImpl::ExtensionList &&extensionList,
+                           cl_ulong hostTimerRes)
+    : mProfile(std::move(profile)),
+      mVersionStr(std::move(versionStr)),
+      mVersion(version),
+      mName(std::move(name)),
+      mExtensions(std::move(extensions)),
+      mExtensionList(std::move(extensionList)),
+      mHostTimerRes(hostTimerRes)
+{}
+
+CLPlatformImpl::CLPlatformImpl(Info &&info) : mInfo(std::move(info)) {}
+
+CLPlatformImpl::~CLPlatformImpl() = default;
+
+}  // namespace rx

src/libANGLE/renderer/CLPlatformImpl.h

@@ -9,7 +9,10 @@
 #ifndef LIBANGLE_RENDERER_CLPLATFORMIMPL_H_
 #define LIBANGLE_RENDERER_CLPLATFORMIMPL_H_
 
-#include "common/angleutils.h"
+#include "libANGLE/renderer/CLtypes.h"
+
+#include <list>
+#include <vector>
 
 namespace rx
 {
@@ -17,10 +20,52 @@ namespace rx
 class CLPlatformImpl : angle::NonCopyable
 {
   public:
-    CLPlatformImpl() {}
-    virtual ~CLPlatformImpl() {}
+    using Ptr           = std::unique_ptr<CLPlatformImpl>;
+    using ImplList      = std::list<Ptr>;
+    using ExtensionList = std::vector<cl_name_version>;
+
+    struct Info
+    {
+        Info();
+        ~Info();
+
+        Info(const Info &) = delete;
+        Info &operator=(const Info &) = delete;
+
+        Info(Info &&);
+        Info &operator=(Info &&);
+
+        Info(std::string &&profile,
+             std::string &&versionStr,
+             cl_version version,
+             std::string &&name,
+             std::string &&extensions,
+             rx::CLPlatformImpl::ExtensionList &&extensionList,
+             cl_ulong hostTimerRes);
+
+        std::string mProfile;
+        std::string mVersionStr;
+        cl_version mVersion;
+        std::string mName;
+        std::string mExtensions;
+        rx::CLPlatformImpl::ExtensionList mExtensionList;
+        cl_ulong mHostTimerRes;
+    };
+
+    explicit CLPlatformImpl(Info &&info);
+    virtual ~CLPlatformImpl();
+
+    const Info &getInfo();
+
+  protected:
+    const Info mInfo;
 };
 
+inline const CLPlatformImpl::Info &CLPlatformImpl::getInfo()
+{
+    return mInfo;
+}
+
 }  // namespace rx
 
 #endif  // LIBANGLE_RENDERER_CLPLATFORMIMPL_H_

src/libANGLE/renderer/CLtypes.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.
+//
+
+// CLtypes.h: Defines common types for the OpenCL support in ANGLE.
+
+#ifndef LIBANGLE_RENDERER_CLTYPES_H_
+#define LIBANGLE_RENDERER_CLTYPES_H_
+
+#include "libANGLE/CLtypes.h"
+
+#include "common/angleutils.h"
+
+namespace rx
+{
+class CLContextImpl;
+class CLDeviceImpl;
+class CLPlatformImpl;
+}  // namespace rx
+
+#endif  // LIBANGLE_RENDERER_CLTYPES_H_

src/libANGLE/renderer/cl/BUILD.gn

@@ -23,6 +23,10 @@ config("angle_cl_backend_config") {
 
 angle_source_set("angle_cl_backend") {
   sources = _cl_backend_sources
+  configs += [ "$angle_root/src/libOpenCL:opencl_icd_library_name" ]
 
-  public_deps = [ "$angle_root:libANGLE_headers" ]
+  public_deps = [
+    "$angle_root:libANGLE_headers",
+    "$angle_root/third_party/OpenCL-ICD-Loader:opencl_icd_loader",
+  ]
 }

src/libANGLE/renderer/cl/CLPlatformCL.cpp

@@ -9,11 +9,230 @@
 
 #include "libANGLE/renderer/cl/CLPlatformCL.h"
 
+#include "libANGLE/CLPlatform.h"
+#include "libANGLE/Debug.h"
+
+#include "anglebase/no_destructor.h"
+#include "common/angle_version.h"
+
+extern "C" {
+#include "icd.h"
+}  // extern "C"
+
+#include <cstdlib>
+#include <unordered_set>
+#include <vector>
+
 namespace rx
 {
 
-CLPlatformCL::CLPlatformCL() {}
+namespace
+{
+using ExtensionSet = std::unordered_set<std::string>;
+
+const ExtensionSet &GetSupportedExtensions()
+{
+    static angle::base::NoDestructor<ExtensionSet> sExtensions(
+        {"cl_khr_extended_versioning", "cl_khr_icd"});
+    return *sExtensions;
+}
+}  // namespace
+
+CLPlatformCL::~CLPlatformCL() = default;
+
+CLPlatformCL::ImplList CLPlatformCL::GetPlatforms(bool isIcd)
+{
+    // Using khrIcdInitialize() of the third party Khronos OpenCL ICD Loader to enumerate the
+    // available OpenCL implementations on the system. They will be stored in the singly linked
+    // list khrIcdVendors of the C struct KHRicdVendor.
+    if (khrIcdVendors == nullptr)
+    {
+        // Our OpenCL entry points are not reentrant, so we have to prevent khrIcdInitialize()
+        // from querying ANGLE's OpenCL ICD library. We store a dummy entry with the library in
+        // the khrIcdVendors list, because the ICD Loader skips the libries which are already in
+        // the list as it assumes they were already enumerated.
+        static angle::base::NoDestructor<KHRicdVendor> sVendorAngle({});
+        sVendorAngle->library = khrIcdOsLibraryLoad(ANGLE_OPENCL_ICD_LIB_NAME);
+        khrIcdVendors         = sVendorAngle.get();
+        if (khrIcdVendors->library != nullptr)
+        {
+            khrIcdInitialize();
+            // After the enumeration we don't need ANGLE's OpenCL ICD library any more,
+            // but we keep the dummy entry int the list to prevent another enumeration.
+            khrIcdOsLibraryUnload(khrIcdVendors->library);
+            khrIcdVendors->library = nullptr;
+        }
+    }
+
+    // Iterating through the singly linked list khrIcdVendors to create an ANGLE CL pass-through
+    // platform for each found ICD platform. Skipping our dummy entry that has an invalid platform.
+    ImplList implList;
+    for (KHRicdVendor *vendorIt = khrIcdVendors; vendorIt != nullptr; vendorIt = vendorIt->next)
+    {
+        if (vendorIt->platform != nullptr)
+        {
+            rx::CLPlatformImpl::Ptr impl = Create(vendorIt->platform);
+            if (impl)
+            {
+                implList.emplace_back(std::move(impl));
+            }
+        }
+    }
+    return implList;
+}
+
+CLPlatformCL::CLPlatformCL(cl_platform_id platform, Info &&info)
+    : CLPlatformImpl(std::move(info)), mPlatform(platform)
+{}
+
+#define ANGLE_GET_INFO(info, size, param, size_ret) \
+    result = platform->getDispatch().clGetPlatformInfo(platform, info, size, param, size_ret)
+
+#define ANGLE_TRY_GET_INFO(info, size, param, size_ret)  \
+    do                                                   \
+    {                                                    \
+        ANGLE_GET_INFO(info, size, param, size_ret);     \
+        if (result != CL_SUCCESS)                        \
+        {                                                \
+            ERR() << "Failed to query CL platform info"; \
+            return std::unique_ptr<CLPlatformCL>();      \
+        }                                                \
+    } while (0)
+
+std::unique_ptr<CLPlatformCL> CLPlatformCL::Create(cl_platform_id platform)
+{
+    cl_int result    = 0;
+    size_t paramSize = 0u;
+    std::vector<std::string::value_type> param;
+    CLPlatformImpl::Info info;
+
+    // Skip ANGLE CL implementation to prevent passthrough loop
+    ANGLE_TRY_GET_INFO(CL_PLATFORM_VENDOR, 0u, nullptr, &paramSize);
+    param.resize(paramSize, '\0');
+    ANGLE_TRY_GET_INFO(CL_PLATFORM_VENDOR, paramSize, param.data(), nullptr);
+    if (std::string(param.data()).compare(cl::Platform::GetVendor()) == 0)
+    {
+        ERR() << "Tried to create CL pass-through back end for ANGLE library";
+        return std::unique_ptr<CLPlatformCL>();
+    }
+
+    // Skip platform if it is not ICD compatible
+    ANGLE_TRY_GET_INFO(CL_PLATFORM_EXTENSIONS, 0u, nullptr, &paramSize);
+    param.resize(paramSize, '\0');
+    ANGLE_TRY_GET_INFO(CL_PLATFORM_EXTENSIONS, paramSize, param.data(), nullptr);
+    info.mExtensions.assign(param.data());
+    if (info.mExtensions.find("cl_khr_icd") == std::string::npos)
+    {
+        WARN() << "CL platform is not ICD compatible";
+        return std::unique_ptr<CLPlatformCL>();
+    }
+
+    // Filter out extensions which are not (yet) supported to be passed through
+    if (!info.mExtensions.empty())
+    {
+        const ExtensionSet &supported   = GetSupportedExtensions();
+        std::string::size_type extStart = 0u;
+        do
+        {
+            const std::string::size_type spacePos = info.mExtensions.find(' ', extStart);
+            const bool foundSpace                 = spacePos != std::string::npos;
+            const std::string::size_type length =
+                (foundSpace ? spacePos : info.mExtensions.length()) - extStart;
+            if (supported.find(info.mExtensions.substr(extStart, length)) != supported.cend())
+            {
+                extStart = foundSpace && spacePos + 1u < info.mExtensions.length()
+                               ? spacePos + 1u
+                               : std::string::npos;
+            }
+            else
+            {
+                info.mExtensions.erase(extStart, length + (foundSpace ? 1u : 0u));
+                if (extStart >= info.mExtensions.length())
+                {
+                    extStart = std::string::npos;
+                }
+            }
+        } while (extStart != std::string::npos);
+        while (!info.mExtensions.empty() && info.mExtensions.back() == ' ')
+        {
+            info.mExtensions.pop_back();
+        }
+    }
+
+    // Fetch common platform info
+    ANGLE_TRY_GET_INFO(CL_PLATFORM_PROFILE, 0u, nullptr, &paramSize);
+    param.resize(paramSize, '\0');
+    ANGLE_TRY_GET_INFO(CL_PLATFORM_PROFILE, paramSize, param.data(), nullptr);
+    info.mProfile.assign(param.data());
+
+    ANGLE_TRY_GET_INFO(CL_PLATFORM_VERSION, 0u, nullptr, &paramSize);
+    param.resize(paramSize, '\0');
+    ANGLE_TRY_GET_INFO(CL_PLATFORM_VERSION, paramSize, param.data(), nullptr);
+    info.mVersionStr.assign(param.data());
+    info.mVersionStr += " (ANGLE " ANGLE_VERSION_STRING ")";
+
+    const std::string::size_type spacePos = info.mVersionStr.find(' ');
+    const std::string::size_type dotPos   = info.mVersionStr.find('.');
+    if (spacePos == std::string::npos || dotPos == std::string::npos)
+    {
+        ERR() << "Failed to extract version from OpenCL version string: " << info.mVersionStr;
+        return std::unique_ptr<CLPlatformCL>();
+    }
+    const cl_uint major =
+        static_cast<cl_uint>(std::strtol(&info.mVersionStr[spacePos + 1u], nullptr, 10));
+    const cl_uint minor =
+        static_cast<cl_uint>(std::strtol(&info.mVersionStr[dotPos + 1u], nullptr, 10));
+    if (major == 0)
+    {
+        ERR() << "Failed to extract version from OpenCL version string: " << info.mVersionStr;
+        return std::unique_ptr<CLPlatformCL>();
+    }
+
+    ANGLE_GET_INFO(CL_PLATFORM_NUMERIC_VERSION, sizeof(info.mVersion), &info.mVersion, nullptr);
+    if (result != CL_SUCCESS)
+    {
+        info.mVersion = CL_MAKE_VERSION(major, minor, 0);
+    }
+    else if (CL_VERSION_MAJOR(info.mVersion) != major || CL_VERSION_MINOR(info.mVersion) != minor)
+    {
+        WARN() << "CL_PLATFORM_NUMERIC_VERSION = " << CL_VERSION_MAJOR(info.mVersion) << '.'
+               << CL_VERSION_MINOR(info.mVersion)
+               << " does not match version string: " << info.mVersionStr;
+    }
+
+    ANGLE_TRY_GET_INFO(CL_PLATFORM_NAME, 0u, nullptr, &paramSize);
+    param.resize(paramSize, '\0');
+    ANGLE_TRY_GET_INFO(CL_PLATFORM_NAME, paramSize, param.data(), nullptr);
+    info.mName.assign("ANGLE pass-through -> ");
+    info.mName += param.data();
+
+    ANGLE_GET_INFO(CL_PLATFORM_EXTENSIONS_WITH_VERSION, 0u, nullptr, &paramSize);
+    if (result == CL_SUCCESS)
+    {
+        info.mExtensionList.resize(paramSize);
+        ANGLE_TRY_GET_INFO(CL_PLATFORM_EXTENSIONS_WITH_VERSION, paramSize,
+                           info.mExtensionList.data(), nullptr);
+
+        // Filter out extensions which are not (yet) supported to be passed through
+        const ExtensionSet &supported       = GetSupportedExtensions();
+        ExtensionList::const_iterator extIt = info.mExtensionList.cbegin();
+        while (extIt != info.mExtensionList.cend())
+        {
+            if (supported.find(extIt->name) != supported.cend())
+            {
+                ++extIt;
+            }
+            else
+            {
+                extIt = info.mExtensionList.erase(extIt);
+            }
+        }
+    }
+
+    ANGLE_GET_INFO(CL_PLATFORM_HOST_TIMER_RESOLUTION, sizeof(info.mHostTimerRes),
+                   &info.mHostTimerRes, nullptr);
 
-CLPlatformCL::~CLPlatformCL() {}
+    return std::unique_ptr<CLPlatformCL>(new CLPlatformCL(platform, std::move(info)));
+}
 
 }  // namespace rx

src/libANGLE/renderer/cl/CLPlatformCL.h

@@ -12,16 +12,33 @@
 
 #include "libANGLE/renderer/CLPlatformImpl.h"
 
+#include <string>
+
 namespace rx
 {
 
 class CLPlatformCL : public CLPlatformImpl
 {
   public:
-    CLPlatformCL();
     ~CLPlatformCL() override;
+
+    cl_platform_id getNative();
+
+    static ImplList GetPlatforms(bool isIcd);
+
+  private:
+    CLPlatformCL(cl_platform_id platform, Info &&info);
+
+    static std::unique_ptr<CLPlatformCL> Create(cl_platform_id platform);
+
+    const cl_platform_id mPlatform;
 };
 
+inline cl_platform_id CLPlatformCL::getNative()
+{
+    return mPlatform;
+}
+
 }  // namespace rx
 
 #endif  // LIBANGLE_RENDERER_CL_CLPLATFORMCL_H_

src/libANGLE/renderer/vulkan/CLPlatformVk.cpp

@@ -9,11 +9,57 @@
 
 #include "libANGLE/renderer/vulkan/CLPlatformVk.h"
 
+#include "anglebase/no_destructor.h"
+#include "common/angle_version.h"
+
+#include <algorithm>
+
 namespace rx
 {
 
-CLPlatformVk::CLPlatformVk() {}
+namespace
+{
+std::string CreateExtensionString(const CLPlatformImpl::ExtensionList &extList)
+{
+    std::string extensions;
+    for (const cl_name_version &ext : extList)
+    {
+        extensions += ext.name;
+        extensions += ' ';
+    }
+    if (!extensions.empty())
+    {
+        extensions.pop_back();
+    }
+    return extensions;
+}
+}  // anonymous namespace
+
+CLPlatformVk::CLPlatformVk(Info &&info) : CLPlatformImpl(std::move(info)) {}
 
 CLPlatformVk::~CLPlatformVk() = default;
 
+CLPlatformVk::ImplList CLPlatformVk::GetPlatforms()
+{
+    ExtensionList extList = {
+        cl_name_version{CL_MAKE_VERSION(1, 0, 0), "cl_khr_icd"},
+        cl_name_version{CL_MAKE_VERSION(1, 0, 0), "cl_khr_extended_versioning"}};
+    std::string extensions = CreateExtensionString(extList);
+
+    Info info("FULL_PROFILE", std::string(GetVersionString()), GetVersion(), "ANGLE Vulkan",
+              std::move(extensions), std::move(extList), 0u);
+
+    ImplList implList;
+    implList.emplace_back(new CLPlatformVk(std::move(info)));
+    return implList;
+}
+
+const std::string &CLPlatformVk::GetVersionString()
+{
+    static const angle::base::NoDestructor<const std::string> sVersion(
+        "OpenCL " + std::to_string(CL_VERSION_MAJOR(GetVersion())) + "." +
+        std::to_string(CL_VERSION_MINOR(GetVersion())) + " ANGLE " ANGLE_VERSION_STRING);
+    return *sVersion;
+}
+
 }  // namespace rx

src/libANGLE/renderer/vulkan/CLPlatformVk.h

@@ -12,16 +12,29 @@
 
 #include "libANGLE/renderer/CLPlatformImpl.h"
 
+#include <string>
+
 namespace rx
 {
 
 class CLPlatformVk : public CLPlatformImpl
 {
   public:
-    CLPlatformVk();
     ~CLPlatformVk() override;
+
+    static ImplList GetPlatforms();
+    static constexpr cl_version GetVersion();
+    static const std::string &GetVersionString();
+
+  private:
+    explicit CLPlatformVk(Info &&info);
 };
 
+constexpr cl_version CLPlatformVk::GetVersion()
+{
+    return CL_MAKE_VERSION(1, 2, 0);
+}
+
 }  // namespace rx
 
 #endif  // LIBANGLE_RENDERER_VULKAN_CLPLATFORMVK_H_

src/libANGLE/validationCL.cpp

@@ -15,6 +15,11 @@ cl_int ValidateGetPlatformIDs(cl_uint num_entries,
                               Platform *const *platformsPacked,
                               const cl_uint *num_platforms)
 {
+    if ((num_entries == 0u && platformsPacked != nullptr) ||
+        (platformsPacked == nullptr && num_platforms == nullptr))
+    {
+        return CL_INVALID_VALUE;
+    }
     return CL_SUCCESS;
 }
 
@@ -24,6 +29,15 @@ cl_int ValidateGetPlatformInfo(const Platform *platformPacked,
                                const void *param_value,
                                const size_t *param_value_size_ret)
 {
+    if (!Platform::IsValid(platformPacked))
+    {
+        return CL_INVALID_PLATFORM;
+    }
+    if (param_namePacked == PlatformInfo::InvalidEnum ||
+        (param_value_size == 0u && param_value != nullptr))
+    {
+        return CL_INVALID_VALUE;
+    }
     return CL_SUCCESS;
 }
 
@@ -561,7 +575,7 @@ cl_int ValidateUnloadCompiler()
 
 bool ValidateGetExtensionFunctionAddress(const char *func_name)
 {
-    return true;
+    return func_name != nullptr && *func_name != '\0';
 }
 
 bool ValidateCreateCommandQueue(const Context *contextPacked,
@@ -816,7 +830,7 @@ cl_int ValidateEnqueueBarrierWithWaitList(const CommandQueue *command_queuePacke
 bool ValidateGetExtensionFunctionAddressForPlatform(const Platform *platformPacked,
                                                     const char *func_name)
 {
-    return true;
+    return Platform::IsValid(platformPacked) && func_name != nullptr && *func_name != '\0';
 }
 
 // CL 2.0
@@ -1045,4 +1059,18 @@ bool ValidateCreateImageWithProperties(const Context *contextPacked,
 {
     return true;
 }
+
+// cl_khr_icd
+cl_int ValidateIcdGetPlatformIDsKHR(cl_uint num_entries,
+                                    Platform *const *platformsPacked,
+                                    const cl_uint *num_platforms)
+{
+    if ((num_entries == 0u && platformsPacked != nullptr) ||
+        (platformsPacked == nullptr && num_platforms == nullptr))
+    {
+        return CL_INVALID_VALUE;
+    }
+    return CL_SUCCESS;
+}
+
 }  // namespace cl

src/libANGLE/validationCL.h

@@ -19,8 +19,6 @@
 #include "libANGLE/CLProgram.h"
 #include "libANGLE/CLSampler.h"
 
-#include "common/PackedCLEnums_autogen.h"
-
 #define ANGLE_CL_VALIDATE_VOID(EP, ...) \
     do                                  \
     {                                   \

src/libANGLE/validationCL_autogen.h

@@ -602,6 +602,11 @@ bool ValidateCreateImageWithProperties(const Context *contextPacked,
                                        const cl_image_desc *image_desc,
                                        const void *host_ptr,
                                        cl_int *errcode_ret);
+
+// cl_khr_icd
+cl_int ValidateIcdGetPlatformIDsKHR(cl_uint num_entries,
+                                    Platform *const *platformsPacked,
+                                    const cl_uint *num_platforms);
 }  // namespace cl
 
 #endif  // LIBANGLE_VALIDATION_CL_AUTOGEN_H_

src/libGLESv2.gni

@@ -467,6 +467,7 @@ libangle_cl_headers = [
   "src/libANGLE/renderer/CLContextImpl.h",
   "src/libANGLE/renderer/CLDeviceImpl.h",
   "src/libANGLE/renderer/CLPlatformImpl.h",
+  "src/libANGLE/renderer/CLtypes.h",
   "src/libANGLE/validationCL.h",
   "src/libANGLE/validationCL_autogen.h",
 ]
@@ -481,6 +482,7 @@ libangle_cl_sources = [
   "src/libANGLE/CLPlatform.cpp",
   "src/libANGLE/CLProgram.cpp",
   "src/libANGLE/CLSampler.cpp",
+  "src/libANGLE/renderer/CLPlatformImpl.cpp",
   "src/libANGLE/validationCL.cpp",
 ]
 
@@ -588,6 +590,8 @@ libglesv2_cl_sources = [
   "src/libGLESv2/entry_points_cl_autogen.cpp",
   "src/libGLESv2/entry_points_cl_autogen.h",
   "src/libGLESv2/entry_points_cl_utils.h",
+  "src/libGLESv2/proc_table_cl.h",
+  "src/libGLESv2/proc_table_cl_autogen.cpp",
 ]
 
 if (is_win) {

src/libGLESv2/cl_dispatch_table.cpp

@@ -12,7 +12,7 @@
 // clang-format off
 
 // The correct order is required as defined in 'include/CL/cl_icd.h'.
-cl_icd_dispatch gCLIcdDispatchTable = {
+const cl_icd_dispatch gCLIcdDispatchTable = {
 
     // OpenCL 1.0
     cl::clGetPlatformIDs,

src/libGLESv2/cl_dispatch_table.h

@@ -13,7 +13,7 @@
 
 extern "C" {
 
-ANGLE_EXPORT extern cl_icd_dispatch gCLIcdDispatchTable;
+ANGLE_EXPORT extern const cl_icd_dispatch gCLIcdDispatchTable;
 
 }  // extern "C"
 

src/libGLESv2/cl_stubs.cpp

@@ -4,15 +4,87 @@
 // found in the LICENSE file.
 //
 // cl_stubs.cpp: Stubs for CL entry points.
-//
 
 #include "libGLESv2/cl_stubs_autogen.h"
 
+#include "libANGLE/CLPlatform.h"
+#include "libGLESv2/cl_dispatch_table.h"
+#include "libGLESv2/proc_table_cl.h"
+
+#ifdef ANGLE_ENABLE_CL_PASSTHROUGH
+#    include "libANGLE/renderer/cl/CLPlatformCL.h"
+#endif
+#ifdef ANGLE_ENABLE_VULKAN
+#    include "libANGLE/renderer/vulkan/CLPlatformVk.h"
+#endif
+
+#include "libANGLE/Debug.h"
+
+#include <cstring>
+
 namespace cl
 {
+
+namespace
+{
+
+const Platform::List &InitializePlatforms(bool isIcd)
+{
+    static bool initialized = false;
+    if (!initialized)
+    {
+        initialized = true;
+
+#ifdef ANGLE_ENABLE_CL_PASSTHROUGH
+        rx::CLPlatformImpl::ImplList implListCL = rx::CLPlatformCL::GetPlatforms(isIcd);
+        while (!implListCL.empty())
+        {
+            Platform::CreatePlatform(gCLIcdDispatchTable, std::move(implListCL.front()));
+            implListCL.pop_front();
+        }
+#endif
+
+#ifdef ANGLE_ENABLE_VULKAN
+        rx::CLPlatformImpl::ImplList implListVk = rx::CLPlatformVk::GetPlatforms();
+        while (!implListVk.empty())
+        {
+            Platform::CreatePlatform(gCLIcdDispatchTable, std::move(implListVk.front()));
+            implListVk.pop_front();
+        }
+#endif
+    }
+    return Platform::GetPlatforms();
+}
+
+cl_int GetPlatforms(cl_uint num_entries, Platform **platforms, cl_uint *num_platforms, bool isIcd)
+{
+    const Platform::List &platformList = InitializePlatforms(isIcd);
+    if (num_platforms != nullptr)
+    {
+        *num_platforms = static_cast<cl_uint>(platformList.size());
+    }
+    if (platforms != nullptr)
+    {
+        cl_uint entry   = 0u;
+        auto platformIt = platformList.cbegin();
+        while (entry < num_entries && platformIt != platformList.cend())
+        {
+            platforms[entry++] = (*platformIt++).get();
+        }
+    }
+    return CL_SUCCESS;
+}
+
+}  // anonymous namespace
+
+cl_int IcdGetPlatformIDsKHR(cl_uint num_entries, Platform **platforms, cl_uint *num_platforms)
+{
+    return GetPlatforms(num_entries, platforms, num_platforms, true);
+}
+
 cl_int GetPlatformIDs(cl_uint num_entries, Platform **platforms, cl_uint *num_platforms)
 {
-    return 0;
+    return GetPlatforms(num_entries, platforms, num_platforms, false);
 }
 
 cl_int GetPlatformInfo(Platform *platform,
@@ -21,7 +93,73 @@ cl_int GetPlatformInfo(Platform *platform,
                        void *param_value,
                        size_t *param_value_size_ret)
 {
-    return 0;
+    cl_version version    = 0u;
+    cl_ulong hostTimerRes = 0u;
+    const void *value     = nullptr;
+    size_t value_size     = 0u;
+    switch (param_name)
+    {
+        case PlatformInfo::Profile:
+            value      = platform->getProfile();
+            value_size = std::strlen(platform->getProfile()) + 1u;
+            break;
+        case PlatformInfo::Version:
+            value      = platform->getVersionString();
+            value_size = std::strlen(platform->getVersionString()) + 1u;
+            break;
+        case PlatformInfo::NumericVersion:
+            version    = platform->getVersion();
+            value      = &version;
+            value_size = sizeof(version);
+            break;
+        case PlatformInfo::Name:
+            value      = platform->getName();
+            value_size = std::strlen(platform->getName()) + 1u;
+            break;
+        case PlatformInfo::Vendor:
+            value      = Platform::GetVendor();
+            value_size = std::strlen(Platform::GetVendor()) + 1u;
+            break;
+        case PlatformInfo::Extensions:
+            value      = platform->getExtensions();
+            value_size = std::strlen(platform->getExtensions()) + 1u;
+            break;
+        case PlatformInfo::ExtensionsWithVersion:
+            if (platform->getExtensionsWithVersion().empty())
+            {
+                return CL_INVALID_VALUE;
+            }
+            value      = platform->getExtensionsWithVersion().data();
+            value_size = platform->getExtensionsWithVersion().size() * sizeof(cl_name_version);
+            break;
+        case PlatformInfo::HostTimerResolution:
+            hostTimerRes = platform->getHostTimerResolution();
+            value        = &hostTimerRes;
+            value_size   = sizeof(hostTimerRes);
+            break;
+        case PlatformInfo::IcdSuffix:
+            value      = Platform::GetIcdSuffix();
+            value_size = std::strlen(Platform::GetIcdSuffix()) + 1u;
+            break;
+        default:
+            return CL_INVALID_VALUE;
+    }
+    if (param_value != nullptr)
+    {
+        if (param_value_size < value_size)
+        {
+            return CL_INVALID_VALUE;
+        }
+        if (value != nullptr)
+        {
+            std::memcpy(param_value, value, value_size);
+        }
+    }
+    if (param_value_size_ret != nullptr)
+    {
+        *param_value_size_ret = value_size;
+    }
+    return CL_SUCCESS;
 }
 
 cl_int GetDeviceIDs(Platform *platform,
@@ -924,7 +1062,7 @@ cl_int EnqueueSVMMigrateMem(CommandQueue *command_queue,
 
 void *GetExtensionFunctionAddressForPlatform(Platform *platform, const char *func_name)
 {
-    return 0;
+    return GetExtensionFunctionAddress(func_name);
 }
 
 cl_int SetCommandQueueProperty(CommandQueue *command_queue,
@@ -985,7 +1123,13 @@ cl_int UnloadCompiler()
 
 void *GetExtensionFunctionAddress(const char *func_name)
 {
-    return 0;
+    if (func_name == nullptr)
+    {
+        return nullptr;
+    }
+    const ProcTable &procTable = GetProcTable();
+    const auto it              = procTable.find(func_name);
+    return it != procTable.cend() ? it->second : nullptr;
 }
 
 CommandQueue *CreateCommandQueue(Context *context,
@@ -1013,4 +1157,5 @@ cl_int EnqueueTask(CommandQueue *command_queue,
 {
     return 0;
 }
+
 }  // namespace cl

src/libGLESv2/cl_stubs_autogen.h

@@ -12,10 +12,11 @@
 
 #include "libANGLE/CLtypes.h"
 
-#include "common/PackedCLEnums_autogen.h"
-
 namespace cl
 {
+cl_int IcdGetPlatformIDsKHR(cl_uint num_entries,
+                            Platform **platformsPacked,
+                            cl_uint *num_platforms);
 cl_int GetPlatformIDs(cl_uint num_entries, Platform **platformsPacked, cl_uint *num_platforms);
 cl_int GetPlatformInfo(Platform *platformPacked,
                        PlatformInfo param_namePacked,

src/libGLESv2/entry_points_cl_autogen.h

@@ -608,6 +608,11 @@ cl_mem CL_API_CALL clCreateImageWithProperties(cl_context context,
                                                const cl_image_desc *image_desc,
                                                void *host_ptr,
                                                cl_int *errcode_ret);
+
+// cl_khr_icd
+cl_int CL_API_CALL clIcdGetPlatformIDsKHR(cl_uint num_entries,
+                                          cl_platform_id *platforms,
+                                          cl_uint *num_platforms);
 }  // namespace cl
 
 #endif  // LIBGLESV2_ENTRY_POINTS_CL_AUTOGEN_H_

src/libGLESv2/proc_table_cl.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.
+//
+// proc_table:
+//   Mapping from a string entry point name to function address.
+//
+
+#ifndef LIBGLESV2_PROC_TABLE_CL_H_
+#define LIBGLESV2_PROC_TABLE_CL_H_
+
+#include <string>
+#include <unordered_map>
+
+namespace cl
+{
+
+using ProcTable = std::unordered_map<std::string, void *>;
+
+const ProcTable &GetProcTable();
+
+}  // namespace cl
+
+#endif  // LIBGLESV2_PROC_TABLE_CL_H_

src/libGLESv2/proc_table_cl_autogen.cpp

+// GENERATED FILE - DO NOT EDIT.
+// Generated by gen_proc_table.py using data from cl.xml.
+//
+// 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.
+//
+// proc_table:
+//   Mapping from a string entry point name to function address.
+//
+
+#include "libGLESv2/proc_table_cl.h"
+
+#include "libGLESv2/entry_points_cl_autogen.h"
+
+#include "anglebase/no_destructor.h"
+
+// Using fully qualified entry point identifiers to make sure that missing entry points would not
+// pick up the global declarations of OpenCL
+
+#define P(FUNC) reinterpret_cast<void *>(FUNC)
+
+namespace cl
+{
+
+const ProcTable &GetProcTable()
+{
+    static angle::base::NoDestructor<ProcTable> sProcTable(
+        {{"clGetPlatformIDs", P(::cl::clGetPlatformIDs)},
+         {"clGetPlatformInfo", P(::cl::clGetPlatformInfo)},
+         {"clGetDeviceIDs", P(::cl::clGetDeviceIDs)},
+         {"clGetDeviceInfo", P(::cl::clGetDeviceInfo)},
+         {"clCreateContext", P(::cl::clCreateContext)},
+         {"clCreateContextFromType", P(::cl::clCreateContextFromType)},
+         {"clRetainContext", P(::cl::clRetainContext)},
+         {"clReleaseContext", P(::cl::clReleaseContext)},
+         {"clGetContextInfo", P(::cl::clGetContextInfo)},
+         {"clRetainCommandQueue", P(::cl::clRetainCommandQueue)},
+         {"clReleaseCommandQueue", P(::cl::clReleaseCommandQueue)},
+         {"clGetCommandQueueInfo", P(::cl::clGetCommandQueueInfo)},
+         {"clCreateBuffer", P(::cl::clCreateBuffer)},
+         {"clRetainMemObject", P(::cl::clRetainMemObject)},
+         {"clReleaseMemObject", P(::cl::clReleaseMemObject)},
+         {"clGetSupportedImageFormats", P(::cl::clGetSupportedImageFormats)},
+         {"clGetMemObjectInfo", P(::cl::clGetMemObjectInfo)},
+         {"clGetImageInfo", P(::cl::clGetImageInfo)},
+         {"clRetainSampler", P(::cl::clRetainSampler)},
+         {"clReleaseSampler", P(::cl::clReleaseSampler)},
+         {"clGetSamplerInfo", P(::cl::clGetSamplerInfo)},
+         {"clCreateProgramWithSource", P(::cl::clCreateProgramWithSource)},
+         {"clCreateProgramWithBinary", P(::cl::clCreateProgramWithBinary)},
+         {"clRetainProgram", P(::cl::clRetainProgram)},
+         {"clReleaseProgram", P(::cl::clReleaseProgram)},
+         {"clBuildProgram", P(::cl::clBuildProgram)},
+         {"clGetProgramInfo", P(::cl::clGetProgramInfo)},
+         {"clGetProgramBuildInfo", P(::cl::clGetProgramBuildInfo)},
+         {"clCreateKernel", P(::cl::clCreateKernel)},
+         {"clCreateKernelsInProgram", P(::cl::clCreateKernelsInProgram)},
+         {"clRetainKernel", P(::cl::clRetainKernel)},
+         {"clReleaseKernel", P(::cl::clReleaseKernel)},
+         {"clSetKernelArg", P(::cl::clSetKernelArg)},
+         {"clGetKernelInfo", P(::cl::clGetKernelInfo)},
+         {"clGetKernelWorkGroupInfo", P(::cl::clGetKernelWorkGroupInfo)},
+         {"clWaitForEvents", P(::cl::clWaitForEvents)},
+         {"clGetEventInfo", P(::cl::clGetEventInfo)},
+         {"clRetainEvent", P(::cl::clRetainEvent)},
+         {"clReleaseEvent", P(::cl::clReleaseEvent)},
+         {"clGetEventProfilingInfo", P(::cl::clGetEventProfilingInfo)},
+         {"clFlush", P(::cl::clFlush)},
+         {"clFinish", P(::cl::clFinish)},
+         {"clEnqueueReadBuffer", P(::cl::clEnqueueReadBuffer)},
+         {"clEnqueueWriteBuffer", P(::cl::clEnqueueWriteBuffer)},
+         {"clEnqueueCopyBuffer", P(::cl::clEnqueueCopyBuffer)},
+         {"clEnqueueReadImage", P(::cl::clEnqueueReadImage)},
+         {"clEnqueueWriteImage", P(::cl::clEnqueueWriteImage)},
+         {"clEnqueueCopyImage", P(::cl::clEnqueueCopyImage)},
+         {"clEnqueueCopyImageToBuffer", P(::cl::clEnqueueCopyImageToBuffer)},
+         {"clEnqueueCopyBufferToImage", P(::cl::clEnqueueCopyBufferToImage)},
+         {"clEnqueueMapBuffer", P(::cl::clEnqueueMapBuffer)},
+         {"clEnqueueMapImage", P(::cl::clEnqueueMapImage)},
+         {"clEnqueueUnmapMemObject", P(::cl::clEnqueueUnmapMemObject)},
+         {"clEnqueueNDRangeKernel", P(::cl::clEnqueueNDRangeKernel)},
+         {"clEnqueueNativeKernel", P(::cl::clEnqueueNativeKernel)},
+         {"clSetCommandQueueProperty", P(::cl::clSetCommandQueueProperty)},
+         {"clCreateImage2D", P(::cl::clCreateImage2D)},
+         {"clCreateImage3D", P(::cl::clCreateImage3D)},
+         {"clEnqueueMarker", P(::cl::clEnqueueMarker)},
+         {"clEnqueueWaitForEvents", P(::cl::clEnqueueWaitForEvents)},
+         {"clEnqueueBarrier", P(::cl::clEnqueueBarrier)},
+         {"clUnloadCompiler", P(::cl::clUnloadCompiler)},
+         {"clGetExtensionFunctionAddress", P(::cl::clGetExtensionFunctionAddress)},
+         {"clCreateCommandQueue", P(::cl::clCreateCommandQueue)},
+         {"clCreateSampler", P(::cl::clCreateSampler)},
+         {"clEnqueueTask", P(::cl::clEnqueueTask)},
+         {"clCreateSubBuffer", P(::cl::clCreateSubBuffer)},
+         {"clSetMemObjectDestructorCallback", P(::cl::clSetMemObjectDestructorCallback)},
+         {"clCreateUserEvent", P(::cl::clCreateUserEvent)},
+         {"clSetUserEventStatus", P(::cl::clSetUserEventStatus)},
+         {"clSetEventCallback", P(::cl::clSetEventCallback)},
+         {"clEnqueueReadBufferRect", P(::cl::clEnqueueReadBufferRect)},
+         {"clEnqueueWriteBufferRect", P(::cl::clEnqueueWriteBufferRect)},
+         {"clEnqueueCopyBufferRect", P(::cl::clEnqueueCopyBufferRect)},
+         {"clCreateSubDevices", P(::cl::clCreateSubDevices)},
+         {"clRetainDevice", P(::cl::clRetainDevice)},
+         {"clReleaseDevice", P(::cl::clReleaseDevice)},
+         {"clCreateImage", P(::cl::clCreateImage)},
+         {"clCreateProgramWithBuiltInKernels", P(::cl::clCreateProgramWithBuiltInKernels)},
+         {"clCompileProgram", P(::cl::clCompileProgram)},
+         {"clLinkProgram", P(::cl::clLinkProgram)},
+         {"clUnloadPlatformCompiler", P(::cl::clUnloadPlatformCompiler)},
+         {"clGetKernelArgInfo", P(::cl::clGetKernelArgInfo)},
+         {"clEnqueueFillBuffer", P(::cl::clEnqueueFillBuffer)},
+         {"clEnqueueFillImage", P(::cl::clEnqueueFillImage)},
+         {"clEnqueueMigrateMemObjects", P(::cl::clEnqueueMigrateMemObjects)},
+         {"clEnqueueMarkerWithWaitList", P(::cl::clEnqueueMarkerWithWaitList)},
+         {"clEnqueueBarrierWithWaitList", P(::cl::clEnqueueBarrierWithWaitList)},
+         {"clGetExtensionFunctionAddressForPlatform",
+          P(::cl::clGetExtensionFunctionAddressForPlatform)},
+         {"clCreateCommandQueueWithProperties", P(::cl::clCreateCommandQueueWithProperties)},
+         {"clCreatePipe", P(::cl::clCreatePipe)},
+         {"clGetPipeInfo", P(::cl::clGetPipeInfo)},
+         {"clSVMAlloc", P(::cl::clSVMAlloc)},
+         {"clSVMFree", P(::cl::clSVMFree)},
+         {"clCreateSamplerWithProperties", P(::cl::clCreateSamplerWithProperties)},
+         {"clSetKernelArgSVMPointer", P(::cl::clSetKernelArgSVMPointer)},
+         {"clSetKernelExecInfo", P(::cl::clSetKernelExecInfo)},
+         {"clEnqueueSVMFree", P(::cl::clEnqueueSVMFree)},
+         {"clEnqueueSVMMemcpy", P(::cl::clEnqueueSVMMemcpy)},
+         {"clEnqueueSVMMemFill", P(::cl::clEnqueueSVMMemFill)},
+         {"clEnqueueSVMMap", P(::cl::clEnqueueSVMMap)},
+         {"clEnqueueSVMUnmap", P(::cl::clEnqueueSVMUnmap)},
+         {"clSetDefaultDeviceCommandQueue", P(::cl::clSetDefaultDeviceCommandQueue)},
+         {"clGetDeviceAndHostTimer", P(::cl::clGetDeviceAndHostTimer)},
+         {"clGetHostTimer", P(::cl::clGetHostTimer)},
+         {"clCreateProgramWithIL", P(::cl::clCreateProgramWithIL)},
+         {"clCloneKernel", P(::cl::clCloneKernel)},
+         {"clGetKernelSubGroupInfo", P(::cl::clGetKernelSubGroupInfo)},
+         {"clEnqueueSVMMigrateMem", P(::cl::clEnqueueSVMMigrateMem)},
+         {"clSetProgramSpecializationConstant", P(::cl::clSetProgramSpecializationConstant)},
+         {"clSetProgramReleaseCallback", P(::cl::clSetProgramReleaseCallback)},
+         {"clSetContextDestructorCallback", P(::cl::clSetContextDestructorCallback)},
+         {"clCreateBufferWithProperties", P(::cl::clCreateBufferWithProperties)},
+         {"clCreateImageWithProperties", P(::cl::clCreateImageWithProperties)},
+         {"clIcdGetPlatformIDsKHR", P(::cl::clIcdGetPlatformIDsKHR)}});
+    return *sProcTable;
+}
+
+}  // namespace cl

src/libOpenCL/BUILD.gn

@@ -28,7 +28,17 @@ angle_shared_library("OpenCL_ANGLE") {
   ]
 }
 
-angle_shared_library("OpenCL_ICD_ANGLE") {
+icd_library_name = "OpenCL_ICD_ANGLE"
+
+config("opencl_icd_library_name") {
+  if (is_win) {
+    defines = [ "ANGLE_OPENCL_ICD_LIB_NAME=\"" + icd_library_name + "\"" ]
+  } else {
+    defines = [ "ANGLE_OPENCL_ICD_LIB_NAME=\"lib" + icd_library_name + ".so\"" ]
+  }
+}
+
+angle_shared_library(icd_library_name) {
   defines = [ "LIBCL_IMPLEMENTATION" ]
 
   sources = [

src/libOpenCL/dispatch.cpp

@@ -18,26 +18,30 @@ 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()
+IcdDispatch CreateDispatch()
 {
-    static cl_icd_dispatch *sDispatch = nullptr;
+    IcdDispatch dispatch;
 
-    if (sDispatch == nullptr)
-    {
     EntryPointsLib().reset(
         angle::OpenSharedLibrary(ANGLE_GLESV2_LIBRARY_NAME, angle::SearchType::ApplicationDir));
     if (EntryPointsLib())
     {
-            sDispatch = reinterpret_cast<cl_icd_dispatch *>(
+        auto clIcdDispatch = reinterpret_cast<const cl_icd_dispatch *>(
             EntryPointsLib()->getSymbol("gCLIcdDispatchTable"));
-            if (sDispatch == nullptr)
+        if (clIcdDispatch != nullptr)
+        {
+            static_cast<cl_icd_dispatch &>(dispatch) = *clIcdDispatch;
+            dispatch.clIcdGetPlatformIDsKHR          = reinterpret_cast<clIcdGetPlatformIDsKHR_fn>(
+                clIcdDispatch->clGetExtensionFunctionAddress("clIcdGetPlatformIDsKHR"));
+        }
+        else
         {
             std::cerr << "Error loading CL dispatch table." << std::endl;
         }
@@ -46,9 +50,16 @@ cl_icd_dispatch &GetDispatch()
     {
         std::cerr << "Error opening GLESv2 library." << std::endl;
     }
-    }
 
-    return *sDispatch;
+    return dispatch;
+}
+
+}  // anonymous namespace
+
+const IcdDispatch &GetDispatch()
+{
+    static const IcdDispatch sDispatch(CreateDispatch());
+    return sDispatch;
 }
 
 }  // namespace cl

src/libOpenCL/dispatch.h

@@ -18,7 +18,12 @@
 namespace cl
 {
 
-cl_icd_dispatch &GetDispatch();
+struct IcdDispatch : public _cl_icd_dispatch
+{
+    clIcdGetPlatformIDsKHR_fn clIcdGetPlatformIDsKHR;
+};
+
+const IcdDispatch &GetDispatch();
 
 }  // namespace cl
 

src/libOpenCL/libOpenCL_ICD.cpp

@@ -3,29 +3,12 @@
 // 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.
+// libOpenCL_ICD.cpp: Implements the CL entry point 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);

src/libOpenCL/libOpenCL_autogen.cpp

@@ -1150,4 +1150,12 @@ cl_mem CL_API_CALL clCreateImageWithProperties(cl_context context,
                                                          image_desc, host_ptr, errcode_ret);
 }
 
+// cl_khr_icd
+cl_int CL_API_CALL clIcdGetPlatformIDsKHR(cl_uint num_entries,
+                                          cl_platform_id *platforms,
+                                          cl_uint *num_platforms)
+{
+    return cl::GetDispatch().clIcdGetPlatformIDsKHR(num_entries, platforms, num_platforms);
+}
+
 }  // extern "C"

third_party/OpenCL-ICD-Loader/BUILD.gn

+# 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.
+
+import("../../gni/angle.gni")
+
+config("opencl_icd_loader_config") {
+  include_dirs = [ "src/loader" ]
+}
+
+source_set("opencl_icd_loader") {
+  sources = [ "src/loader/icd.c" ]
+  defines = [
+    "CL_TARGET_OPENCL_VERSION=300",
+    "OPENCL_ICD_LOADER_DISABLE_OPENCLON12",
+  ]
+
+  if (is_win) {
+    sources += [
+      "src/loader/windows/icd_windows.c",
+      "src/loader/windows/icd_windows_apppackage.cpp",
+      "src/loader/windows/icd_windows_dxgk.c",
+      "src/loader/windows/icd_windows_envvars.c",
+      "src/loader/windows/icd_windows_hkr.c",
+    ]
+
+    configs -= [ "//build/config/win:lean_and_mean" ]
+    libs = [ "cfgmgr32.lib" ]
+
+    if (is_clang) {
+      cflags = [ "-Wno-incompatible-pointer-types" ]
+    } else {
+      cflags = [ "/wd4133" ]
+    }
+  } else {
+    sources += [
+      "src/loader/linux/icd_linux.c",
+      "src/loader/linux/icd_linux_envvars.c",
+    ]
+    include_dirs = [ "." ]
+
+    if (is_android) {
+      cflags = [ "-Wno-#pragma-messages" ]
+    } else {
+      defines += [ "HAVE_SECURE_GETENV" ]
+    }
+  }
+
+  configs -= [ "//build/config/compiler:chromium_code" ]
+  configs += [ "//build/config/compiler:no_chromium_code" ]
+
+  public_configs = [ ":opencl_icd_loader_config" ]
+  public_deps = [ "$angle_root:cl_includes" ]
+}

third_party/OpenCL-ICD-Loader/README.chromium

+Name: Khronos OpenCL ICD Loader
+Short Name: OpenCL-ICD-Loader
+Version: v2021.04.29
+URL: https://chromium.googlesource.com/external/github.com/KhronosGroup/OpenCL-ICD-Loader
+SOURCE CODE: git clone -b OpenCL-ICD-Loader https://chromium.googlesource.com/external/github.com/KhronosGroup/OpenCL-ICD-Loader
+Date: 29/04/2021
+Revision: 9b5e3849b49a1448996c8b96ba086cd774d987db
+Security Critical: no
+License: Apache 2.0
+License File: LICENSE
+
+Description:
+OpenCL defines an Installable Client Driver (ICD) mechanism to allow developers to build
+applications against an Installable Client Driver loader (ICD loader) rather than linking
+their applications against a specific OpenCL implementation.

third_party/OpenCL-ICD-Loader/icd_cmake_config.h

+/* Empty on purpose, configured by BUILD.gn */