Merge remote-tracking branch 'origin/sycl' into private/asachkov/track-amount-of-xfailed-tests-without-tracker

Author: AlexeySachkov

.github/workflows/sycl-linux-run-tests.yml

@@ -151,7 +151,7 @@ permissions:
 
 jobs:
   run:
-    if: inputs.skip_run == 'false'
+    if: github.event_name == 'workflow_dispatch' || inputs.skip_run == 'false'
     name: ${{ inputs.name }}
     runs-on: ${{ fromJSON(inputs.runner) }}
     container:

clang/lib/Driver/ToolChains/Cuda.cpp

@@ -946,6 +946,15 @@ void CudaToolChain::addClangTargetOptions(
 
     if (DriverArgs.hasArg(options::OPT_fsycl_fp32_prec_sqrt))
       CC1Args.push_back("-fcuda-prec-sqrt");
+
+    bool FastRelaxedMath = DriverArgs.hasFlag(
+        options::OPT_ffast_math, options::OPT_fno_fast_math, false);
+    bool UnsafeMathOpt =
+        DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations,
+                           options::OPT_fno_unsafe_math_optimizations, false);
+    if (FastRelaxedMath || UnsafeMathOpt)
+      CC1Args.append({"-mllvm", "--nvptx-prec-divf32=0", "-mllvm",
+                      "--nvptx-prec-sqrtf32=0"});
   } else {
     CC1Args.append(
         {"-fcuda-is-device", "-mllvm", "-enable-memcpyopt-without-libcalls"});

clang/test/Driver/sycl-device-traits-macros-amdgcn.cpp

@@ -50,8 +50,10 @@
 // RUN:   FileCheck %s --check-prefix=CHECK-SYCL-AMDGCN-AMD-AMDHSA-DEVICE-TRIPLE
 // RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amd_gpu_gfx940 \
 // RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
+// RUN:   FileCheck %s --check-prefix=CHECK-SYCL-AMDGCN-AMD-AMDHSA-DEVICE-TRIPLE
 // RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amd_gpu_gfx941 \
 // RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
+// RUN:   FileCheck %s --check-prefix=CHECK-SYCL-AMDGCN-AMD-AMDHSA-DEVICE-TRIPLE
 // RUN: %clangxx -fsycl -nogpulib -fsycl-targets=amd_gpu_gfx942 \
 // RUN:   -fsycl-libspirv-path=%S/Inputs/SYCL/libspirv.bc -### %s 2>&1 | \
 // RUN:   FileCheck %s --check-prefix=CHECK-SYCL-AMDGCN-AMD-AMDHSA-DEVICE-TRIPLE

clang/test/Driver/sycl-nvptx-fast-math.cpp

@@ -0,0 +1,18 @@
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang -### -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-DEFAULT %s
+
+// RUN: %clang -### -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -ffast-math %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-FAST %s
+
+// RUN: %clang -### -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -funsafe-math-optimizations %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-FAST %s
+
+// CHECK-FAST: "-mllvm" "--nvptx-prec-divf32=0" "-mllvm" "--nvptx-prec-sqrtf32=0"
+
+// CHECK-DEFAULT-NOT: "nvptx-prec-divf32=0"
+// CHECK-DEFAULT-NOT: "nvptx-prec-sqrtf32=0"

libdevice/nativecpu_utils.cpp

@@ -65,13 +65,13 @@ __spirv_MemoryBarrier(uint32_t Memory, uint32_t Semantics) {
   template <>                                                                  \
   __SYCL_CONVERGENT__ DEVICE_EXTERNAL Type                                     \
   __spirv_SubgroupBlockReadINTEL<Type>(const OCL_GLOBAL PType *Ptr) noexcept { \
-    return *Ptr;                                                               \
+    return Ptr[__spirv_SubgroupLocalInvocationId()];                           \
   }                                                                            \
   template <>                                                                  \
   __SYCL_CONVERGENT__ DEVICE_EXTERNAL void                                     \
   __spirv_SubgroupBlockWriteINTEL<Type>(PType OCL_GLOBAL * ptr,                \
                                         Type v) noexcept {                     \
-    *(Type *)ptr = v;                                                          \
+    ((Type*)ptr)[__spirv_SubgroupLocalInvocationId()]  = v;                    \
   }
 
 #define DefSubgroupBlockINTEL_vt(Type, VT_name)                                \
@@ -92,16 +92,19 @@ template <class T> struct vtypes {
 DefSubgroupBlockINTEL(uint32_t) DefSubgroupBlockINTEL(uint64_t)
 DefSubgroupBlockINTEL(uint8_t) DefSubgroupBlockINTEL(uint16_t)
 
-#define DefineGOp1(spir_sfx, mux_name)\
-DEVICE_EXTERN_C bool mux_name(bool);\
+#define DefineGOp1(spir_sfx, name)\
+DEVICE_EXTERN_C bool __mux_sub_group_##name##_i1(bool);\
+DEVICE_EXTERN_C bool __mux_work_group_##name##_i1(uint32_t id, bool val);\
 DEVICE_EXTERNAL bool __spirv_Group ## spir_sfx(unsigned g, bool val) {\
   if (__spv::Scope::Flag::Subgroup == g)\
-    return mux_name(val);\
+    return __mux_sub_group_##name##_i1(val);\
+  else if (__spv::Scope::Flag::Workgroup == g)\
+    return __mux_work_group_##name##_i1(0, val);\
   return false;\
 }
 
-DefineGOp1(Any, __mux_sub_group_any_i1)
-DefineGOp1(All, __mux_sub_group_all_i1)
+DefineGOp1(Any, any)
+DefineGOp1(All, all)
 
 
 #define DefineGOp(Type, MuxType, spir_sfx, mux_sfx)                            \
@@ -184,18 +187,6 @@ DefineBitwiseGroupOp(uint64_t, int64_t, i64)
 
 DefineLogicalGroupOp(bool, bool, i1)
 
-#define DefineBroadCastImpl(Type, Sfx, MuxType, IDType)                       \
-  DEVICE_EXTERN_C MuxType __mux_work_group_broadcast_##Sfx(                   \
-      int32_t id, MuxType val, int64_t lidx, int64_t lidy, int64_t lidz);     \
-  DEVICE_EXTERN_C MuxType __mux_sub_group_broadcast_##Sfx(MuxType val,        \
-                                                          int32_t sg_lid);    \
-  DEVICE_EXTERNAL Type __spirv_GroupBroadcast(uint32_t g, Type v,             \
-                                              IDType l) {                     \
-    if (__spv::Scope::Flag::Subgroup == g)                                    \
-      return __mux_sub_group_broadcast_##Sfx(v, l);                           \
-    return Type(); /*todo: add support for other flags as they are tested*/   \
-  }
-
 #define DefineBroadcastMuxType(Type, Sfx, MuxType, IDType)                    \
   DEVICE_EXTERN_C MuxType __mux_work_group_broadcast_##Sfx(                   \
       int32_t id, MuxType val, uint64_t lidx, uint64_t lidy, uint64_t lidz);  \
@@ -216,7 +207,7 @@ DefineLogicalGroupOp(bool, bool, i1)
     if (__spv::Scope::Flag::Subgroup == g)                                    \
       return __mux_sub_group_broadcast_##Sfx(v, l[0]);                        \
     else                                                                      \
-      return __mux_work_group_broadcast_##Sfx(0, v, l[0], l[0], 0);           \
+      return __mux_work_group_broadcast_##Sfx(0, v, l[0], l[1], 0);           \
   }                                                                           \
                                                                               \
   DEVICE_EXTERNAL Type __spirv_GroupBroadcast(uint32_t g, Type v,             \

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -85,6 +85,7 @@ def AspectExt_oneapi_virtual_mem : Aspect<"ext_oneapi_virtual_mem">;
 def AspectExt_oneapi_cuda_cluster_group : Aspect<"ext_oneapi_cuda_cluster_group">;
 def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">;
 def AspectExt_oneapi_atomic16 : Aspect<"ext_oneapi_atomic16">;
+def AspectExt_oneapi_virtual_functions : Aspect<"ext_oneapi_virtual_functions">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -148,7 +149,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph, AspectExt_oneapi_private_alloca, 
     AspectExt_oneapi_queue_profiling_tag, AspectExt_oneapi_virtual_mem, AspectExt_oneapi_cuda_cluster_group, 
     AspectExt_intel_fpga_task_sequence,
-    AspectExt_oneapi_atomic16],
+    AspectExt_oneapi_atomic16,
+    AspectExt_oneapi_virtual_functions],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp

@@ -37,6 +37,7 @@ getSYCLESIMDSplitStatusFromMetadata(const Module &M) {
   assert(MDOp && "Unexpected metadata operand");
   const auto &MDConst = MDOp->getOperand(0);
   auto *MDVal = mdconst::dyn_extract_or_null<ConstantInt>(MDConst);
+  assert(MDVal && "Unexpected metadata operand type");
   uint8_t Val = MDVal->getZExtValue();
   assert(Val < 3 && "Unexpected value for split metadata");
   auto AsEnum = static_cast<module_split::SyclEsimdSplitStatus>(Val);

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -117,13 +117,14 @@ if(SYCL_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit b79ebe4e98789144bcdf3832088eb6e6b5ae6967
-  # Merge: 7b4bc761 fbb6e862
-  # Author: Kenneth Benzie (Benie) <[email protected]>
-  # Date:   Fri Oct 4 16:39:59 2024 +0100
-  #     Merge pull request #2018 from wenju-he/L0-bindless-image-device-query
-  #     [L0] Fix device query of bindless image support
-  set(UNIFIED_RUNTIME_TAG b79ebe4e98789144bcdf3832088eb6e6b5ae6967)
+  # commit df6da35d6e67f2383db28dd49ab08c5c0ef541d2
+  # Merge: 67590533 55bd5636
+  # Author: aarongreig <[email protected]>
+  # Date:   Mon Oct 7 12:28:07 2024 +0100
+  #     Merge pull request #2038 from GeorgeWeb/georgi/unsupported-max-coop-wgsize
+  #     [UR][hip][opencl] Mark urKernelSuggestMaxCooperativeGroupCountExp as unsupported
+  #     instead of returning misleading default value
+  set(UNIFIED_RUNTIME_TAG df6da35d6e67f2383db28dd49ab08c5c0ef541d2)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
   # Due to the use of dependentloadflag and no installer for UMF and hwloc we need

sycl/doc/syclcompat/README.md

@@ -855,6 +855,41 @@ public:
 } // syclcompat
 ```
 
+### ptr_to_int
+
+The following cuda backend specific function is introduced in order to
+translate from local memory pointers to `uint32_t` or `size_t` variables that
+contain a byte address to the local (local refers to`.shared` in nvptx) memory
+state space.
+
+``` c++
+namespace syclcompat {
+template <typename T>
+__syclcompat_inline__
+    std::enable_if_t<std::is_same_v<T, uint32_t> || std::is_same_v<T, size_t>,
+                     T>
+    ptr_to_int(void *ptr)
+} // namespace syclcompat
+```
+
+These variables can be used in inline PTX instructions that take address
+operands. Such inline PTX instructions are commonly used in optimized
+libraries. A simplified example usage of the above functions is as follows:
+
+``` c++
+  half *data = syclcompat::local_mem<half[NUM_ELEMENTS]>();
+  // ...
+  // ...
+  T addr =
+      syclcompat::ptr_to_int<T>(reinterpret_cast<char *>(data) + (id % 8) * 16);
+  uint32_t fragment;
+#if defined(__NVPTX__)
+  asm volatile("ldmatrix.sync.aligned.m8n8.x1.shared.b16 {%0}, [%1];\n"
+               : "=r"(fragment)
+               : "r"(addr));
+#endif
+```
+
 ### Device Information
 
 `sycl::device` properties are encapsulated using the `device_info` helper class.
@@ -1544,10 +1579,10 @@ SYCL spec supported by the current SYCL compiler.
 
 The `SYCLCOMPAT_CHECK_ERROR` macro encapsulates an error-handling mechanism for
 expressions that might throw `sycl::exception` and `std::runtime_error`. If no
-exceptions are thrown, it returns `syclcompat::error_code::SUCCESS`. If a
-`sycl::exception` is caught, it returns `syclcompat::error_code::BACKEND_ERROR`.
+exceptions are thrown, it returns `syclcompat::error_code::success`. If a
+`sycl::exception` is caught, it returns `syclcompat::error_code::backend_error`.
 If a `std::runtime_error` exception is caught,
-`syclcompat::error_code::DEFAULT_ERROR` is returned instead. For both cases, it
+`syclcompat::error_code::default_error` is returned instead. For both cases, it
 prints the error message to the standard error stream.
 
 ``` c++
@@ -1580,7 +1615,7 @@ template <int Arg> class syclcompat_kernel_scalar;
 
 
 namespace syclcompat {
-enum error_code { SUCCESS = 0, BACKEND_ERROR = 1, DEFAULT_ERROR = 999 };
+enum error_code { success = 0, backend_error = 1, default_error = 999 };
 }
 
 #define SYCLCOMPAT_CHECK_ERROR(expr)

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -1365,7 +1365,15 @@ __clc_BarrierTestWait(int64_t *state, int64_t arrival) noexcept;
 __SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL __SYCL_EXPORT void
 __clc_BarrierArriveAndWait(int64_t *state) noexcept;
 
-#ifdef __SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__
+#if defined(__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__) &&                         \
+    !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
+#if defined(__clang__)
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wpedantic"
+#warning                                                                       \
+    "__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ is deprecated and will be removed in a future release."
+#pragma clang diagnostic pop
+#endif
 extern __DPCPP_SYCL_EXTERNAL int
 __spirv_ocl_printf(const __attribute__((opencl_constant)) char *Format, ...);
 extern __DPCPP_SYCL_EXTERNAL int __spirv_ocl_printf(const char *Format, ...);