Merge branch 'sycl' into e2e-regression

Author: ayylol

devops/cts_exclude_filter

@@ -5,4 +5,3 @@ marray
 math_builtin_api
 # https://github.com/intel/llvm/issues/13574
 hierarchical
-accessor

devops/dependencies-igc-dev.json

@@ -1,10 +1,10 @@
 {
   "linux": {
     "igc_dev": {
-      "github_tag": "igc-dev-3bd1d5e",
-      "version": "3bd1d5e",
-      "updated_at": "2024-06-08T23:45:49Z",
-      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/1582291042/zip",
+      "github_tag": "igc-dev-480f8b6",
+      "version": "480f8b6",
+      "updated_at": "2024-06-12T22:42:55Z",
+      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/1595870554/zip",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     }
   }

libclc/amdgcn-amdhsa/libspirv/synchronization/barrier.cl

@@ -10,42 +10,56 @@
 #include <spirv/spirv.h>
 #include <spirv/spirv_types.h>
 
-#define BUILTIN_FENCE(semantics, scope_memory)                                 \
-  if (semantics & Acquire)                                                     \
-    return __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, scope_memory);             \
-  else if (semantics & Release)                                                \
-    return __builtin_amdgcn_fence(__ATOMIC_RELEASE, scope_memory);             \
-  else if (semantics & AcquireRelease)                                         \
-    return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, scope_memory);             \
-  else if (semantics & SequentiallyConsistent)                                 \
-    return __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, scope_memory);             \
-  else                                                                         \
-    return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, scope_memory);
 
-_CLC_DEF _CLC_OVERLOAD void __mem_fence(unsigned int scope_memory,
-                                        unsigned int semantics) {
+#define BUILTIN_FENCE(order, scope_memory)                                     \
+  /* None implies Monotonic (for llvm/AMDGPU), or relaxed in C++.              \
+   * This does not make sense as ordering argument for a fence instruction     \
+   * and is not part of the supported orderings for a fence in AMDGPU. */      \
+  if (order != None) {                                                         \
+    switch (order) {                                                           \
+    case Acquire:                                                              \
+      return __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, scope_memory);           \
+    case Release:                                                              \
+      return __builtin_amdgcn_fence(__ATOMIC_RELEASE, scope_memory);           \
+    case AcquireRelease:                                                       \
+      return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, scope_memory);           \
+    case SequentiallyConsistent:                                               \
+      return __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, scope_memory);           \
+    default:                                                                   \
+      __builtin_trap();                                                        \
+      __builtin_unreachable();                                                 \
+    }                                                                          \
+  }
+
+_CLC_INLINE void builtin_fence_order(unsigned int scope_memory,
+                                     unsigned int order) {
   switch ((enum Scope)scope_memory) {
   case CrossDevice:
-    BUILTIN_FENCE(semantics, "")
+    BUILTIN_FENCE(order, "")
   case Device:
-    BUILTIN_FENCE(semantics, "agent")
+    BUILTIN_FENCE(order, "agent")
   case Workgroup:
-    BUILTIN_FENCE(semantics, "workgroup")
+    BUILTIN_FENCE(order, "workgroup")
   case Subgroup:
-    BUILTIN_FENCE(semantics, "wavefront")
+    BUILTIN_FENCE(order, "wavefront")
   case Invocation:
-    BUILTIN_FENCE(semantics, "singlethread")
+    BUILTIN_FENCE(order, "singlethread")
   }
 }
 #undef BUILTIN_FENCE
 
+_CLC_DEF _CLC_OVERLOAD void __mem_fence(unsigned int scope_memory,
+                                        unsigned int semantics) {
+  builtin_fence_order(scope_memory, semantics & 0x1F);
+}
+
 _CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(unsigned int scope_memory,
                                                   unsigned int semantics) {
   __mem_fence(scope_memory, semantics);
 }
 
 _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT void
-__spirv_ControlBarrier(unsigned int scope_execution, unsigned scope_memory,
+__spirv_ControlBarrier(unsigned int scope_execution, unsigned int scope_memory,
                        unsigned int semantics) {
   if (semantics) {
     __mem_fence(scope_memory, semantics);

libdevice/cmath_wrapper.cpp

@@ -39,6 +39,18 @@ extern "C" SYCL_EXTERNAL float __devicelib_fminf(float, float);
 DEVICE_EXTERN_C_INLINE
 float fminf(float x, float y) { return __devicelib_fminf(x, y); }
 
+DEVICE_EXTERN_C_INLINE
+float truncf(float x) { return __devicelib_truncf(x); }
+
+DEVICE_EXTERN_C_INLINE
+float sinpif(float x) { return __devicelib_sinpif(x); }
+
+DEVICE_EXTERN_C_INLINE
+float rsqrtf(float x) { return __devicelib_rsqrtf(x); }
+
+DEVICE_EXTERN_C_INLINE
+float exp10f(float x) { return __devicelib_exp10f(x); }
+
 DEVICE_EXTERN_C_INLINE
 div_t div(int x, int y) { return __devicelib_div(x, y); }
 

libdevice/cmath_wrapper_fp64.cpp

@@ -36,6 +36,18 @@ extern "C" SYCL_EXTERNAL double __devicelib_fmin(double, double);
 DEVICE_EXTERN_C_INLINE
 double fmin(double x, double y) { return __devicelib_fmin(x, y); }
 
+DEVICE_EXTERN_C_INLINE
+double trunc(double x) { return __devicelib_trunc(x); }
+
+DEVICE_EXTERN_C_INLINE
+double sinpi(double x) { return __devicelib_sinpi(x); }
+
+DEVICE_EXTERN_C_INLINE
+double rsqrt(double x) { return __devicelib_rsqrt(x); }
+
+DEVICE_EXTERN_C_INLINE
+double exp10(double x) { return __devicelib_exp10(x); }
+
 DEVICE_EXTERN_C_INLINE
 double log(double x) { return __devicelib_log(x); }
 

libdevice/device_math.h

@@ -76,6 +76,30 @@ float __devicelib_fminf(float x, float y);
 DEVICE_EXTERN_C
 double __devicelib_fmin(double x, double y);
 
+DEVICE_EXTERN_C
+float __devicelib_truncf(float x);
+
+DEVICE_EXTERN_C
+double __devicelib_trunc(double x);
+
+DEVICE_EXTERN_C
+double __devicelib_sinpi(double x);
+
+DEVICE_EXTERN_C
+float __devicelib_sinpif(float x);
+
+DEVICE_EXTERN_C
+double __devicelib_rsqrt(double x);
+
+DEVICE_EXTERN_C
+float __devicelib_rsqrtf(float x);
+
+DEVICE_EXTERN_C
+double __devicelib_exp10(double x);
+
+DEVICE_EXTERN_C
+float __devicelib_exp10f(float x);
+
 DEVICE_EXTERN_C
 div_t __devicelib_div(int x, int y);
 

libdevice/fallback-cmath-fp64.cpp

@@ -35,6 +35,18 @@ double __devicelib_fmax(double x, double y) { return __spirv_ocl_fmax(x, y); }
 DEVICE_EXTERN_C_INLINE
 double __devicelib_fmin(double x, double y) { return __spirv_ocl_fmin(x, y); }
 
+DEVICE_EXTERN_C_INLINE
+double __devicelib_trunc(double x) { return __spirv_ocl_trunc(x); }
+
+DEVICE_EXTERN_C_INLINE
+double __devicelib_sinpi(double x) { return __spirv_ocl_sinpi(x); }
+
+DEVICE_EXTERN_C_INLINE
+double __devicelib_rsqrt(double x) { return __spirv_ocl_rsqrt(x); }
+
+DEVICE_EXTERN_C_INLINE
+double __devicelib_exp10(double x) { return __spirv_ocl_exp10(x); }
+
 DEVICE_EXTERN_C_INLINE
 double __devicelib_log(double x) { return __spirv_ocl_log(x); }
 

libdevice/fallback-cmath.cpp

@@ -45,6 +45,18 @@ float __devicelib_fmaxf(float x, float y) { return __spirv_ocl_fmax(x, y); }
 DEVICE_EXTERN_C_INLINE
 float __devicelib_fminf(float x, float y) { return __spirv_ocl_fmin(x, y); }
 
+DEVICE_EXTERN_C_INLINE
+float __devicelib_truncf(float x) { return __spirv_ocl_trunc(x); }
+
+DEVICE_EXTERN_C_INLINE
+float __devicelib_sinpif(float x) { return __spirv_ocl_sinpi(x); }
+
+DEVICE_EXTERN_C_INLINE
+float __devicelib_rsqrtf(float x) { return __spirv_ocl_rsqrt(x); }
+
+DEVICE_EXTERN_C_INLINE
+float __devicelib_exp10f(float x) { return __spirv_ocl_exp10(x); }
+
 DEVICE_EXTERN_C_INLINE
 div_t __devicelib_div(int x, int y) { return {x / y, x % y}; }
 

mlir/utils/vscode/package-lock.json

@@ -164,3 +164,39 @@ These tests verify SYCL specification conformance. All implementation details
 are out of scope for the tests.
 See DPC++ compiler invocation definitions at
 [FindIntel_SYCL](https://github.com/KhronosGroup/SYCL-CTS/blob/SYCL-1.2.1/master/cmake/FindIntel_SYCL.cmake))
+
+## Unified Runtime Updates
+
+To integrate changes from the [Unified Runtime][ur] project into DPC++ there
+two main options which depend on the scope of those changes and the current
+state of DPC++.
+
+1. Synchronized update:
+  * When: If the Unified Runtime change touches the API/ABI, more than one
+    adapter, or common code such as the loader.
+  * How: Update the `UNIFIED_RUNTIME_TAG` to point at the desired commit or tag
+    name in the Unified Runtime repository and ensure that any tag for specific
+    adapters are set to use `${UNIFIED_RUNTIME_TAG}`.
+
+2. Decoupled update:
+  * When: If only a single Unified Runtime adatper has changed.
+  * How: Update the tag used in the `fetch_adapter_source()` call for a
+    specific Unified Runtime adapter, e.g. Level Zero, OpenCL, CUDA, HIP, or
+    Native CPU.
+
+In general, a synchronized update should be the default. However, when there
+are a lot of changes in flight in parallel always synchronizing the tag can be
+troublesome. This is when a decoupled update can help sustain the merge
+velocity of Unified Runtime changes.
+
+The [intel/unified-runtime-reviewers][ur-reviewers-team] team is responsible
+for ensuring that the Unified Runtime tag is updated correctly and will only
+provide code owner approval to pull requests once the following criteria are
+met:
+
+* Tags are pointing to a valid commit or tag on Unified Runtime main branch.
+* Changes to additional code owned files are in a good state.
+* GitHub Actions checks are passing.
+
+[ur]: https://github.com/oneapi-src/unified-runtime
+[ur-reviewers-team]: https://github.com/orgs/intel/teams/unified-runtime-reviewers

sycl/doc/developer/ContributeToDPCPP.md

@@ -568,9 +568,14 @@ class [[__sycl_detail__::__uses_aspects__(aspect::atomic64)]] atomic_ref_impl<
 // Partial specialization for pointer types
 // Arithmetic is emulated because target's representation of T* is unknown
 // TODO: Find a way to use intptr_t or uintptr_t atomics instead
-template <typename T, bool IsAspectAtomic64AttrUsed, memory_order DefaultOrder, memory_scope DefaultScope,
-          access::address_space AddressSpace>
-class atomic_ref_impl<T *, IsAspectAtomic64AttrUsed, DefaultOrder, DefaultScope, AddressSpace>
+template <typename T, bool IsAspectAtomic64AttrUsed, memory_order DefaultOrder,
+          memory_scope DefaultScope, access::address_space AddressSpace>
+#ifndef __SYCL_DEVICE_ONLY__
+class atomic_ref_impl<
+#else
+class [[__sycl_detail__::__uses_aspects__(aspect::atomic64)]] atomic_ref_impl<
+#endif
+    T *, IsAspectAtomic64AttrUsed, DefaultOrder, DefaultScope, AddressSpace>
     : public atomic_ref_base<uintptr_t, DefaultOrder, DefaultScope,
                              AddressSpace> {
 

sycl/include/sycl/atomic_ref.hpp

@@ -342,6 +342,8 @@ template <typename T> auto convertToOpenCLType(T &&x) {
                                                    std::declval<ElemTy>()))>,
                             no_ref::size()>;
 #ifdef __SYCL_DEVICE_ONLY__
+
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
     // TODO: for some mysterious reasons on NonUniformGroups E2E tests fail if
     // we use the "else" version only. I suspect that's an issues with
     // non-uniform groups implementation.
@@ -350,6 +352,10 @@ template <typename T> auto convertToOpenCLType(T &&x) {
     else
       return static_cast<typename MatchingVec::vector_t>(
           x.template as<MatchingVec>());
+#else  // __INTEL_PREVIEW_BREAKING_CHANGES
+    return sycl::bit_cast<typename MatchingVec::vector_t>(x);
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
+
 #else
     return x.template as<MatchingVec>();
 #endif