[SYCL] kernel_compiler support of OpenCL queries

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc

@@ -41,6 +41,11 @@ This extension is written against the SYCL 2020 revision 8 specification.
 All references below to the "core SYCL specification" or to section numbers in
 the SYCL specification refer to that revision.
 
+This extension references sections of the OpenCL specification version 3.0.14.
+References below to the "OpenCL specification" refer to that version.
+It also references sections of the OpenCL C specification version 3.0.14.
+References below to the "OpenCL C specification" refer to that version.
+
 This extension also depends on the following other SYCL extensions:
 
 * link:../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc[
@@ -114,6 +119,140 @@ As a result, the application must use the overload of
 `create_kernel_bundle_from_source` taking `std::string` when creating a kernel
 bundle from this language.
 
+=== Queries
+
+==== Version type
+
+This extension adds the following type and constant definitions, which help
+identify the version of OpenCL and its extensions.
+
+|====
+a|
+[frame=all,grid=none]
+!====
+a!
+[source]
+----
+namespace sycl::ext::oneapi::experimental {
+
+struct cl_version {
+  unsigned major:10;
+  unsigned minor:10;
+  unsigned patch:12;
+};
+
+inline constexpr cl_version opencl_c_1_0 = {1,0,0};
+inline constexpr cl_version opencl_c_1_1 = {1,1,0};
+inline constexpr cl_version opencl_c_1_2 = {1,2,0};
+inline constexpr cl_version opencl_c_2_0 = {2,0,0};
+inline constexpr cl_version opencl_c_3_0 = {3,0,0};
+
+} // namespace ext::oneapi::experimental
+----
+!====
+
+The meaning of the `major`, `minor`, and `patch` values are defined by section
+3.4.3.1 "Versions" of the OpenCL specification.
+
+The constant values (e.g. `opencl_c_1_0`) are shorthands that identify various
+OpenCL C versions.
+
+[_Note:_ The OpenCL C version is not the same as the the OpenCL version because
+some minor releases of OpenCL did not change the OpenCL C language.
+For example, there is no version of OpenCL C named "2.1" even though there is
+an OpenCL version named "2.1".
+_{endnote}_]
+|====
+
+==== New member functions for the device class
+
+This extension also adds the following member functions to the `device` class,
+which allow the application to query which OpenCL features and extensions the
+device supports.
+
+|====
+a|
+[frame=all,grid=none]
+!====
+a!
+[source]
+----
+class device {
+  bool ext_oneapi_supports_cl_c_version(
+    const ext::oneapi::experimental::cl_version &version) const;
+};
+----
+!====
+
+_Returns:_ The value `true` only if the device supports kernel bundles written
+in the OpenCL C version identified by `version`.
+Returns `false` if the device does not support kernel bundles written in
+`source_language::opencl`.
+
+a|
+[frame=all,grid=none]
+!====
+a!
+[source]
+----
+class device {
+  bool ext_oneapi_supports_cl_c_feature(const std::string &name) const;
+};
+----
+!====
+
+_Returns:_ The value `true` only if the device supports kernel bundles using
+the OpenCL C feature whose feature macro is `name`.
+The set of possible feature macros are defined in section 6.2.1 "Features" of
+the OpenCL C specification.
+Returns `false` if the device does not support kernel bundles written in
+`source_language::opencl`.
+
+a|
+[frame=all,grid=none]
+!====
+a!
+[source]
+----
+class device {
+  bool ext_oneapi_supports_cl_extension(const std::string &name,
+                                        ext::oneapi::experimental::cl_version *version = nullptr) const;
+};
+----
+!====
+
+_Effects:_ If the device supports kernel bundles using the OpenCL extension
+identified by `name` and if `version` is not a null pointer, the supported
+version of the extension is written to `version`.
+
+_Returns:_ The value `true` only if the device supports kernel bundles using
+the OpenCL extension identified by `name`.
+Returns `false` if the device does not support kernel bundles written in
+`source_language::opencl`.
+
+a|
+[frame=all,grid=none]
+!====
+a!
+[source]
+----
+class device {
+  std::string ext_oneapi_cl_profile() const;
+};
+----
+!====
+
+_Returns:_ If the device supports kernel bundles written in
+`source_language::opencl`, returns the name of the OpenCL profile that is
+supported.
+The profile name is the same string that is returned by the query
+`CL_DEVICE_PROFILE`, as defined in section 4.2 "Querying Devices" of the OpenCL
+specification.
+If the device does not support kernel bundles written in
+`source_language::opencl`, returns the empty string.
+
+|====
+
 === Build options
 
 The `build_options` property accepts any of the compiler or linker options
@@ -122,6 +261,13 @@ creating an OpenCL library.
 The kernel compiler can be used to create an OpenCL program, but not an OpenCL
 library.
 
+The `-cl-std=` option is required when compiling kernels that use OpenCL C 2.0
+or OpenCL C 3.0 features.
+Otherwise, the implementation defaults to the highest OpenCL C 1.x language
+version that each device supports.
+See section 5.8.6.5 "Options Controlling the OpenCL C version" of the OpenCL
+specification for details.
+
 === Obtaining a kernel
 
 OpenCL C kernel functions do not support {cpp} features like overloads or
@@ -187,7 +333,9 @@ _{endnote}_]
 |===
 
 
-== Example
+== Examples
+
+=== Simple example
 
 The following example shows a simple SYCL program that defines an OpenCL C
 kernel as a string and then compiles and launches it.
@@ -243,52 +391,45 @@ int main() {
 }
 ```
 
+=== Querying supported features and extensions
 
-== Issues
-
-* How should we expose the difference between OpenCL C versions?
-  It seems like there are two aspects to the problem.
-  Applications need some way to query which versions the backend (or device)
-  supports.
-  Applications also need some way to tell the runtime which version the kernel
-  is written in.
-+
---
-One option is to define separate enumerators in `source_language` for each
-version like this:
+This example demonstrates how to query the version of OpenCL C that is
+supported, how to query the supported features, and how to query the
+supported extensions.
 
 ```
-enum class source_language : /*unspecified*/ {
-  opencl_1_0,
-  opencl_1_1,
-  opencl_2_0,
-  opencl_3_0,
-};
+#include <iostream>
+#include <sycl/sycl.hpp>
+namespace syclex = sycl::ext::oneapi::experimental;
+
+int main() {
+  sycl::queue q;
+  sycl::device d = q.get_device();
+
+  if (d.ext_oneapi_can_compile(syclex::source_language::opencl))
+    std::cout << "Device supports online compilation of OpenCL C kernels\n";
+
+  if (d.ext_oneapi_supports_cl_c_version(syclex::opencl_c_3_0))
+    std::cout << "Device supports online compilation with OpenCL C 3.0\n";
+
+  if (d.ext_oneapi_supports_cl_c_feature("__opencl_c_fp64"))
+    std::cout << "Device supports online compilation with 64-bit FP in OpenCL C\n";
+
+  syclex::cl_version version;
+  if (d.ext_oneapi_supports_cl_extension("cl_intel_bfloat16_conversions", &version)) {
+    std::cout << "Device supports online compilation of OpenCL C with bfloat16 "
+      "conversions (version: " << version.major << "." << version.minor << "." <<
+      version.patch << ")\n";
+  }
+
+  if (d.ext_oneapi_cl_profile().find("FULL_PROFILE") == 0)
+    std::cout << "Device supports online compilation with the OpenCL full profile\n";
+
+}
 ```
 
-Applications could then query the supported versions via
-`is_source_kernel_bundle_supported`, and applications would identify the
-version of their kernel string via the `lang` parameter to
-`create_kernel_bundle_from_source`.
-
-Alternatively, this extension could define just a single language enumerator
-(`opencl`), but also provide as separate query to get the supported OpenCL C
-versions.
-When building a kernel bundle, applications would be required to pass "-cl-std"
-via the `build_options` property in order to identify the OpenCL C version of
-their source string.
---
-
-* How can an application determine the OpenCL C optional features that are
-  supported and the extensions that are supported?
-  One option is to require the application to use OpenCL APIs for these
-  queries.
-  This seems better than duplicating these queries into this extension.
-  However, this assumes the application is running with an OpenCL backend.
-  Do we want to support the use of OpenCL C kernels also with the Level Zero
-  backend?
-  Currently, the online_compiler does support this case (but it provides no way
-  to query about optional features or extensions).
+
+== Issues
 
 * Do we need to document some restrictions on the OpenCL C
   https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_C.html#work-item-functions[

sycl/include/sycl/device.hpp

@@ -23,6 +23,7 @@
 #include <sycl/device_selector.hpp>
 #include <sycl/ext/oneapi/experimental/device_architecture.hpp>
 #include <sycl/info/info_desc.hpp>
+#include <sycl/kernel_bundle_enums.hpp>
 #include <sycl/platform.hpp>
 
 #include <cstddef>
@@ -282,6 +283,57 @@ class __SYCL_EXPORT device : public detail::OwnerLessBase<device> {
   /// the function.
   bool ext_oneapi_architecture_is(ext::oneapi::experimental::architecture arch);
 
+  /// kernel_compiler extension
+
+  /// Indicates if the device can compile a kernel for the given language.
+  ///
+  /// \param Language is one of the values from the
+  /// kernel_bundle::source_language enumeration described in the
+  /// sycl_ext_oneapi_kernel_compiler specification
+  ///
+  /// \return true only if the device supports kernel bundles written in the
+  /// source language `lang`.
+  bool
+  ext_oneapi_can_compile(ext::oneapi::experimental::source_language Language);
+
+  /// Indicates if the device supports a given feature when compiling the OpenCL
+  /// C language
+  ///
+  /// \param Feature
+  ///
+  /// \return true if supported
+  bool ext_oneapi_supports_cl_c_feature(const std::string &Feature);
+
+  /// Indicates if the device supports kernel bundles written in a particular
+  /// OpenCL C version
+  ///
+  /// \param Version
+  ///
+  /// \return true only if the device supports kernel bundles written in the
+  /// version identified by `Version`.
+  bool ext_oneapi_supports_cl_c_version(
+      const ext::oneapi::experimental::cl_version &Version) const;
+
+  /// If the device supports kernel bundles using the OpenCL extension
+  /// identified by `name` and if `version` is not a null pointer, the supported
+  /// version of the extension is written to `version`.
+  ///
+  /// \return true only if the device supports kernel bundles using the OpenCL
+  /// extension identified by `name`.
+  bool ext_oneapi_supports_cl_extension(
+      const std::string &name,
+      ext::oneapi::experimental::cl_version *version = nullptr) const;
+
+  /// Retrieve the OpenCl Device Profile
+  ///
+  /// \return If the device supports kernel bundles written in
+  /// `source_language::opencl`, returns the name of the OpenCL profile that is
+  /// supported. The profile name is the same string that is returned by the
+  /// query `CL_DEVICE_PROFILE`, as defined in section 4.2 "Querying Devices" of
+  /// the OpenCL specification. If the device does not support kernel bundles
+  /// written in `source_language::opencl`, returns the empty string.
+  std::string ext_oneapi_cl_profile() const;
+
 // TODO: Remove this diagnostics when __SYCL_WARN_IMAGE_ASPECT is removed.
 #if defined(__clang__)
 #pragma clang diagnostic pop

sycl/include/sycl/kernel_bundle_enums.hpp

@@ -22,6 +22,19 @@ namespace ext::oneapi::experimental {
 
 enum class source_language : int { opencl = 0, spirv = 1 /* sycl, cuda */ };
 
+// opencl versions
+struct cl_version {
+  unsigned major : 10;
+  unsigned minor : 10;
+  unsigned patch : 12;
+};
+
+inline constexpr cl_version opencl_c_1_0 = {1, 0, 0};
+inline constexpr cl_version opencl_c_1_1 = {1, 1, 0};
+inline constexpr cl_version opencl_c_1_2 = {1, 2, 0};
+inline constexpr cl_version opencl_c_2_0 = {2, 0, 0};
+inline constexpr cl_version opencl_c_3_0 = {3, 0, 0};
+
 } // namespace ext::oneapi::experimental
 
 } // namespace _V1

sycl/source/detail/device_impl.cpp

@@ -707,6 +707,15 @@ bool device_impl::isGetDeviceAndHostTimerSupported() {
   return Result != PI_ERROR_INVALID_OPERATION;
 }
 
+bool device_impl::extOneapiCanCompile(
+    ext::oneapi::experimental::source_language Language) {
+  try {
+    return is_source_kernel_bundle_supported(getBackend(), Language);
+  } catch (sycl::exception &) {
+    return false;
+  }
+}
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl

sycl/source/detail/device_impl.hpp

@@ -238,6 +238,8 @@ class device_impl {
     return Arch == getDeviceArch();
   }
 
+  bool extOneapiCanCompile(ext::oneapi::experimental::source_language Language);
+
   /// Gets the current device timestamp
   /// @throw sycl::feature_not_supported if feature is not supported on device
   uint64_t getCurrentDeviceTime();