For pulling in PR intel#7567

Author: premanandrao

clang/lib/Driver/ToolChains/Clang.cpp

@@ -9556,7 +9556,8 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA,
                 ",+SPV_INTEL_bfloat16_conversion"
                 ",+SPV_INTEL_joint_matrix"
                 ",+SPV_INTEL_hw_thread_queries"
-                ",+SPV_KHR_uniform_group_instructions";
+                ",+SPV_KHR_uniform_group_instructions"
+                ",+SPV_INTEL_masked_gather_scatter";
     TranslatorArgs.push_back(TCArgs.MakeArgString(ExtArg));
   }
   for (auto I : Inputs) {

clang/test/Driver/sycl-spirv-ext.c

@@ -48,7 +48,8 @@
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_bfloat16_conversion
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_joint_matrix
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_hw_thread_queries
-// CHECK-DEFAULT-SAME:,+SPV_KHR_uniform_group_instructions"
+// CHECK-DEFAULT-SAME:,+SPV_KHR_uniform_group_instructions
+// CHECK-DEFAULT-SAME:,+SPV_INTEL_masked_gather_scatter"
 // CHECK-FPGA-HW: llvm-spirv{{.*}}"-spirv-ext=-all
 // CHECK-FPGA-HW-SAME:,+SPV_EXT_shader_atomic_float_add
 // CHECK-FPGA-HW-SAME:,+SPV_EXT_shader_atomic_float_min_max

sycl/doc/extensions/supported/sycl_ext_oneapi_queue_priority.asciidoc

@@ -0,0 +1,107 @@
+= sycl_ext_oneapi_queue_priority
+
+:source-highlighter: coderay
+:coderay-linenums-mode: table
+
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+:dpcpp: pass:[DPC++]
+
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
+
+
+== Notice
+
+[%hardbreaks]
+Copyright (C) 2022-2022 Intel Corporation.  All rights reserved.
+
+Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
+of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
+permission by Khronos.
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/intel/llvm/issues
+
+
+== Dependencies
+
+This extension is written against the SYCL 2020 revision 6 specification.  All
+references below to the "core SYCL specification" or to section numbers in the
+SYCL specification refer to that revision.
+
+== Status
+
+This extension is implemented and fully supported by {dpcpp}.
+[NOTE]
+====
+Although {dpcpp} supports this extension on all backends, it is currently used
+only on Level Zero. Other backends ignore the properties defined in this specification.
+====
+
+== Overview
+
+Introduce SYCL queue properties specifying the desired priority of a queue.
+These priorities are a hint and may be ignored if not supported by
+underlying backends. 
+
+== Specification
+
+=== Feature test macro
+
+This extension provides a feature-test macro as described in the core SYCL
+specification.  An implementation supporting this extension must predefine
+the macro `SYCL_EXT_ONEAPI_QUEUE_PRIORITY` to one of the values defined
+in the table below.  Applications can test for the existence of this macro
+to determine if the implementation supports this feature, or applications
+can test the macro's value to determine which of the extension's features
+the implementation supports.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|Initial version of this extension.
+|===
+
+=== API of the extension
+
+This extension adds support for new properties for SYCL queue constructors
+taking properties list:
+
+```c++
+namespace sycl::ext::oneapi::property::queue {
+
+  class priority_normal {
+  public:
+    priority_normal() = default;
+  };
+  class priority_low {
+  public:
+    priority_low() = default;
+  };
+  class priority_high {
+  public:
+    priority_high() = default;
+  };
+
+} // namespace
+```
+The new properties hint the SYCL runtime that the queue gets the specified
+priority for execution if supported by underlying target runtimes. These
+properties are hints and may safely be ignored by an implementation.
+
+It is illegal to specify multiple differrent priority hints for the same queue.
+Doing so causes the `queue` constructor to throw a synchronous `exception` with
+the `errc::invalid` error code.

sycl/include/sycl/detail/pi.h

@@ -56,6 +56,8 @@
 // 11.16 Add PI_EXT_INTEL_DEVICE_INFO_MEMORY_CLOCK_RATE and
 // PI_EXT_INTEL_DEVICE_INFO_MEMORY_BUS_WIDTH as an extension for
 // piDeviceGetInfo.
+// 11.17 Added new PI_EXT_ONEAPI_QUEUE_PRIORITY_LOW and
+// PI_EXT_ONEAPI_QUEUE_PRIORITY_HIGH queue properties.
 
 #define _PI_H_VERSION_MAJOR 11
 #define _PI_H_VERSION_MINOR 16
@@ -580,6 +582,8 @@ constexpr pi_queue_properties PI_QUEUE_PROFILING_ENABLE = (1 << 1);
 constexpr pi_queue_properties PI_QUEUE_ON_DEVICE = (1 << 2);
 constexpr pi_queue_properties PI_QUEUE_ON_DEVICE_DEFAULT = (1 << 3);
 constexpr pi_queue_properties PI_EXT_ONEAPI_QUEUE_DISCARD_EVENTS = (1 << 4);
+constexpr pi_queue_properties PI_EXT_ONEAPI_QUEUE_PRIORITY_LOW = (1 << 5);
+constexpr pi_queue_properties PI_EXT_ONEAPI_QUEUE_PRIORITY_HIGH = (1 << 6);
 
 using pi_result = _pi_result;
 using pi_platform_info = _pi_platform_info;

sycl/include/sycl/detail/properties_traits.def

@@ -4,10 +4,16 @@ __SYCL_PARAM_TRAITS_SPEC(sycl::property::buffer::context_bound)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::image::use_host_ptr)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::image::use_mutex)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::image::context_bound)
-__SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::property::buffer::use_pinned_host_memory)
+__SYCL_PARAM_TRAITS_SPEC(
+    sycl::ext::oneapi::property::buffer::use_pinned_host_memory)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::noinit)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::no_init)
-__SYCL_PARAM_TRAITS_SPEC(sycl::property::context::cuda::use_primary_context) // Deprecated
-__SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::cuda::property::context::use_primary_context)
+__SYCL_PARAM_TRAITS_SPEC(
+    sycl::property::context::cuda::use_primary_context) // Deprecated
+__SYCL_PARAM_TRAITS_SPEC(
+    sycl::ext::oneapi::cuda::property::context::use_primary_context)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::queue::in_order)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::reduction::initialize_to_identity)
+__SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::property::queue::priority_low)
+__SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::property::queue::priority_high)
+__SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::property::queue::priority_normal)

sycl/include/sycl/detail/property_helper.hpp

@@ -40,8 +40,11 @@ enum DataLessPropKind {
   FusionNoBarrier = 13,
   FusionEnable = 14,
   FusionForce = 15,
+  QueuePriorityNormal = 16,
+  QueuePriorityLow = 17,
+  QueuePriorityHigh = 18,
   // Indicates the last known dataless property.
-  LastKnownDataLessPropKind = 15,
+  LastKnownDataLessPropKind = 18,
   // Exceeding 32 may cause ABI breaking change on some of OSes.
   DataLessPropKindSize = 32
 };

sycl/include/sycl/ext/oneapi/bfloat16.hpp

@@ -25,9 +25,21 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) {
 namespace ext {
 namespace oneapi {
 
+class bfloat16;
+
+namespace detail {
+using Bfloat16StorageT = uint16_t;
+Bfloat16StorageT bfloat16ToBits(const bfloat16 &Value);
+bfloat16 bitsToBfloat16(const Bfloat16StorageT Value);
+} // namespace detail
+
 class bfloat16 {
-  using storage_t = uint16_t;
-  storage_t value;
+  detail::Bfloat16StorageT value;
+
+  friend inline detail::Bfloat16StorageT
+  detail::bfloat16ToBits(const bfloat16 &Value);
+  friend inline bfloat16
+  detail::bitsToBfloat16(const detail::Bfloat16StorageT Value);
 
 public:
   bfloat16() = default;
@@ -36,7 +48,7 @@ class bfloat16 {
 
 private:
   // Explicit conversion functions
-  static storage_t from_float(const float &a) {
+  static detail::Bfloat16StorageT from_float(const float &a) {
 #if defined(__SYCL_DEVICE_ONLY__)
 #if defined(__NVPTX__)
 #if (__CUDA_ARCH__ >= 800)
@@ -72,7 +84,7 @@ class bfloat16 {
 #endif
   }
 
-  static float to_float(const storage_t &a) {
+  static float to_float(const detail::Bfloat16StorageT &a) {
 #if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__)
     return __devicelib_ConvertBF16ToFINTEL(a);
 #else
@@ -85,12 +97,6 @@ class bfloat16 {
 #endif
   }
 
-  static bfloat16 from_bits(const storage_t &a) {
-    bfloat16 res;
-    res.value = a;
-    return res;
-  }
-
 public:
   // Implicit conversion from float to bfloat16
   bfloat16(const float &a) { value = from_float(a); }
@@ -122,7 +128,7 @@ class bfloat16 {
 #if defined(__SYCL_DEVICE_ONLY__)
 #if defined(__NVPTX__)
 #if (__CUDA_ARCH__ >= 800)
-    return from_bits(__nvvm_neg_bf16(lhs.value));
+    return detail::bitsToBfloat16(__nvvm_neg_bf16(lhs.value));
 #else
     return -to_float(lhs.value);
 #endif
@@ -203,6 +209,23 @@ class bfloat16 {
   // for floating-point types.
 };
 
+namespace detail {
+
+// Helper function for getting the internal representation of a bfloat16.
+inline Bfloat16StorageT bfloat16ToBits(const bfloat16 &Value) {
+  return Value.value;
+}
+
+// Helper function for creating a float16 from a value with the same type as the
+// internal representation.
+inline bfloat16 bitsToBfloat16(const Bfloat16StorageT Value) {
+  bfloat16 res;
+  res.value = Value;
+  return res;
+}
+
+} // namespace detail
+
 } // namespace oneapi
 } // namespace ext
 

sycl/include/sycl/ext/oneapi/experimental/backend/cuda.hpp

@@ -69,6 +69,13 @@ interop_handle::get_native_context<backend::ext_oneapi_cuda>() const {
 template <>
 inline device make_device<backend::ext_oneapi_cuda>(
     const backend_input_t<backend::ext_oneapi_cuda, device> &BackendObject) {
+  auto devs = device::get_devices(info::device_type::gpu);
+  for (auto &dev : devs) {
+    if (dev.get_backend() == backend::ext_oneapi_cuda &&
+        BackendObject == get_native<backend::ext_oneapi_cuda>(dev)) {
+      return dev;
+    }
+  }
   pi_native_handle NativeHandle = static_cast<pi_native_handle>(BackendObject);
   return ext::oneapi::cuda::make_device(NativeHandle);
 }