Merge from 'sycl' to 'sycl-web'

Author: iclsrc

clang/include/clang/Driver/Driver.h

@@ -801,6 +801,11 @@ class Driver {
   /// targets.
   mutable llvm::StringMap<StringRef> SYCLUniqueIDList;
 
+  /// Vector of Macros that need to be added to the Host compilation in a
+  /// SYCL based offloading scenario.  These macros are gathered during
+  /// construction of the device compilations.
+  mutable std::vector<std::string> SYCLTargetMacroArgs;
+
   /// Return the typical executable name for the specified driver \p Mode.
   static const char *getExecutableForDriverMode(DriverMode Mode);
 
@@ -871,6 +876,17 @@ class Driver {
   void createAppendedFooterInput(Action *&Input, Compilation &C,
                                  const llvm::opt::ArgList &Args) const;
 
+  /// addSYCLTargetMacroArg - Add the given macro to the vector of args to be
+  /// added to the host compilation step.
+  void addSYCLTargetMacroArg(const llvm::opt::ArgList &Args,
+                             StringRef Macro) const {
+    SYCLTargetMacroArgs.push_back(Args.MakeArgString(Macro));
+  }
+  /// getSYCLTargetMacroArgs - return the previously gathered macro target args.
+  llvm::ArrayRef<std::string> getSYCLTargetMacroArgs() const {
+    return SYCLTargetMacroArgs;
+  }
+
   /// setSYCLUniqueID - set the Unique ID that is used for all FE invocations
   /// when performing compilations for SYCL.
   void addSYCLUniqueID(StringRef UniqueID, StringRef FileName) const {

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5115,19 +5115,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
 
     // Forward -fsycl-default-sub-group-size if in SYCL mode.
     Args.AddLastArg(CmdArgs, options::OPT_fsycl_default_sub_group_size);
-
-    // Add any predefined macros associated with intel_gpu* type targets
-    // passed in with -fsycl-targets
-    if (RawTriple.isSPIR() &&
-        RawTriple.getSubArch() == llvm::Triple::SPIRSubArch_gen) {
-      StringRef Device = JA.getOffloadingArch();
-      if (!Device.empty())
-        CmdArgs.push_back(Args.MakeArgString(
-            Twine("-D") + SYCL::gen::getGenDeviceMacro(Device)));
-    }
-    if (RawTriple.isSPIR() &&
-        RawTriple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64)
-      CmdArgs.push_back("-D__SYCL_TARGET_INTEL_X86_64__");
   }
 
   if (IsSYCL) {
@@ -5213,6 +5200,35 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
         }
       }
     }
+    // Add any predefined macros associated with intel_gpu* type targets
+    // passed in with -fsycl-targets
+    // TODO: Macros are populated during device compilations and saved for
+    // addition to the host compilation. There is no dependence connection
+    // between device and host where we should be able to use the offloading
+    // arch to add the macro to the host compile.
+    auto addTargetMacros = [&](const llvm::Triple &Triple) {
+      if (!Triple.isSPIR())
+        return;
+      SmallString<64> Macro;
+      if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen) {
+        StringRef Device = JA.getOffloadingArch();
+        if (!Device.empty()) {
+          Macro = "-D";
+          Macro += SYCL::gen::getGenDeviceMacro(Device);
+        }
+      } else if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64)
+        Macro = "-D__SYCL_TARGET_INTEL_X86_64__";
+      if (Macro.size()) {
+        CmdArgs.push_back(Args.MakeArgString(Macro));
+        D.addSYCLTargetMacroArg(Args, Macro);
+      }
+    };
+    if (IsSYCLOffloadDevice)
+      addTargetMacros(RawTriple);
+    else {
+      for (auto &Macro : D.getSYCLTargetMacroArgs())
+        CmdArgs.push_back(Args.MakeArgString(Macro));
+    }
   }
 
   if (IsOpenMPDevice) {

clang/test/Driver/sycl-intel-gpu.cpp

@@ -84,6 +84,8 @@
 // MACRO: clang{{.*}} "-triple" "spir64_gen-unknown-unknown"
 // MACRO: "-D__SYCL_TARGET_INTEL_GPU_[[MAC_STR]]__"
 // DEVICE: ocloc{{.*}} "-device" "[[DEV_STR]]"
+// MACRO: clang{{.*}} "-fsycl-is-host"
+// MACRO: "-D__SYCL_TARGET_INTEL_GPU_[[MAC_STR]]__"
 
 /// -fsycl-targets=spir64_x86_64 should set a specific macro
 // RUN: %clangxx -c -fsycl -fsycl-targets=spir64_x86_64 -### %s 2>&1 | \
@@ -92,6 +94,8 @@
 // RUN:   FileCheck %s --check-prefix=MACRO_X86_64
 // MACRO_X86_64: clang{{.*}} "-triple" "spir64_x86_64-unknown-unknown"
 // MACRO_X86_64: "-D__SYCL_TARGET_INTEL_X86_64__"
+// MACRO_X86_64: clang{{.*}} "-fsycl-is-host"
+// MACRO_X86_64: "-D__SYCL_TARGET_INTEL_X86_64__"
 
 /// test for invalid arch
 // RUN: %clangxx -c -fsycl -fsycl-targets=intel_gpu_bad -### %s 2>&1 | \

devops/test_configs.json

@@ -9,7 +9,7 @@
       ],
       "image": "${{ inputs.intel_drivers_image }}",
       "container_options": "-u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN",
-      "check_sycl_all": "level_zero:gpu,host",
+      "check_sycl_all": "level_zero:gpu",
       "cmake_args": ""
     },
     {
@@ -21,7 +21,7 @@
       ],
       "image": "${{ inputs.intel_drivers_image }}",
       "container_options": "-u 1001 --device=/dev/dri --privileged --cap-add SYS_ADMIN",
-      "check_sycl_all": "opencl:gpu,host",
+      "check_sycl_all": "opencl:gpu",
       "cmake_args": ""
     },
     {
@@ -33,7 +33,7 @@
       ],
       "image": "${{ inputs.intel_drivers_image }}",
       "container_options": "-u 1001",
-      "check_sycl_all": "opencl:cpu,host",
+      "check_sycl_all": "opencl:cpu",
       "cmake_args": ""
     },
     {
@@ -57,7 +57,7 @@
       ],
       "image": "${{ inputs.amdgpu_image }}",
       "container_options": "--device=/dev/dri --device=/dev/kfd",
-      "check_sycl_all": "hip:gpu,host",
+      "check_sycl_all": "hip:gpu",
       "cmake_args": "-DHIP_PLATFORM=\"AMD\" -DAMD_ARCH=\"gfx1031\""
     },
     {
@@ -69,7 +69,7 @@
       ],
       "image": "${{ inputs.cuda_image }}",
       "container_options": "--gpus all",
-      "check_sycl_all": "cuda:gpu,host",
+      "check_sycl_all": "cuda:gpu",
       "cmake_args": ""
     }
   ],
@@ -83,7 +83,7 @@
       ],
       "image": "${{ inputs.cuda_image }}",
       "container_options": "--gpus all",
-      "sycl_device_filter": "ext_oneapi_cuda:gpu,host",
+      "sycl_device_filter": "ext_oneapi_cuda:gpu",
       "cmake_args": "-DDPCPP_TARGET_TRIPLES=nvptx64-nvidia-cuda"
     }
   ]

libclc/ptx-nvidiacl/libspirv/group/collectives.cl

@@ -426,18 +426,16 @@ __CLC_GROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, long, 0l)
 __CLC_GROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, long, 0l)
 
 // half requires additional mangled entry points
-_CLC_DEF _CLC_CONVERGENT half _Z17__spirv_GroupFAddjjDF16_(uint scope, uint op,
-                                                           half x) {
-  return __spirv_GroupFAdd(scope, op, x);
-}
-_CLC_DEF _CLC_CONVERGENT half _Z17__spirv_GroupFMinjjDF16_(uint scope, uint op,
-                                                           half x) {
-  return __spirv_GroupFMin(scope, op, x);
-}
-_CLC_DEF _CLC_CONVERGENT half _Z17__spirv_GroupFMaxjjDF16_(uint scope, uint op,
-                                                           half x) {
-  return __spirv_GroupFMax(scope, op, x);
-}
+#define __CLC_GROUP_COLLECTIVE__DF16(MANGLED_NAME, SPIRV_DISPATCH)             \
+  _CLC_DEF _CLC_CONVERGENT half MANGLED_NAME(uint scope, uint op, half x) {    \
+    return SPIRV_DISPATCH(scope, op, x);                                       \
+  }
+__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFAddjjDF16_, __spirv_GroupFAdd)
+__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFMinjjDF16_, __spirv_GroupFMin)
+__CLC_GROUP_COLLECTIVE__DF16(_Z17__spirv_GroupFMaxjjDF16_, __spirv_GroupFMax)
+__CLC_GROUP_COLLECTIVE__DF16(_Z20__spirv_GroupFMulKHRjjDF16_,
+                             __spirv_GroupFMulKHR)
+#undef __CLC_GROUP_COLLECTIVE__DF16
 
 #undef __CLC_GROUP_COLLECTIVE_4
 #undef __CLC_GROUP_COLLECTIVE_5

sycl/ReleaseNotes.md

@@ -48,22 +48,22 @@ Release notes for commit range [`4043dda3..0f579bae`](https://github.com/intel/l
 ### Documentation
 
 - Added stateful to stateless memory access conversion
-  [design document](sycl/doc/design/ESIMDStatelesAccessors.md). [3e03f300]
-- Added [`sycl_ext_oneapi_complex`](sycl/doc/extensions/proposed/sycl_ext_oneapi_complex.asciidoc)
+  [design document](doc/design/ESIMDStatelesAccessors.md). [3e03f300]
+- Added [`sycl_ext_oneapi_complex`](doc/extensions/proposed/sycl_ext_oneapi_complex.asciidoc)
   extension proposal. [01589da5]
-- Updated [`sycl_ext_intel_fpga_device_selector`](sycl/doc/extensions/supported/sycl_ext_intel_fpga_device_selector.asciidoc)
+- Updated [`sycl_ext_intel_fpga_device_selector`](doc/extensions/supported/sycl_ext_intel_fpga_device_selector.asciidoc)
   extension to add `fpga_simulator_accessor`. [9bef890d]
-- Added [`sycl_ext_intel_fpga_kernel_interface_properties`](sycl/doc/extension/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc) extension proposal. [4b6bd14b]
-- Updated [`sycl_ext_oneapi_complex_algorithms`](sycl/doc/extensions/proposed/sycl_ext_oneapi_complex_algorithms.asciidoc)
+- Added [`sycl_ext_intel_fpga_kernel_interface_properties`](doc/extension/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc) extension proposal. [4b6bd14b]
+- Updated [`sycl_ext_oneapi_complex_algorithms`](doc/extensions/proposed/sycl_ext_oneapi_complex_algorithms.asciidoc)
   extension to include `sycl::complex` as supported type for algorithms. [07c5b48f]
-- Clarified sub-group size calculation in [`sycl_ext_oneapi_invoke_simd`](sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc) extension spec. [9b33ad0f]
-- Updated [`sycl_ext_oneapi_accessor_properties`](sycl/doc/extensions/supported/sycl_ext_oneapi_accessor_properties.asciidoc)
+- Clarified sub-group size calculation in [`sycl_ext_oneapi_invoke_simd`](doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc) extension spec. [9b33ad0f]
+- Updated [`sycl_ext_oneapi_accessor_properties`](doc/extensions/supported/sycl_ext_oneapi_accessor_properties.asciidoc)
   to mark `has_property` API as `noexcept`. [7805aa3f]
-- Updated [`sycl_ext_intel_device_info`](sycl/doc/extensions/supported/sycl_ext_intel_device_info.md)
+- Updated [`sycl_ext_intel_device_info`](doc/extensions/supported/sycl_ext_intel_device_info.md)
   to support querying free device memory. [0eeef2b3]
-- Updated [`sycl_ext_oneapi_matrix`](sycl/doc/extensions/experimental/sycl_ext_oneapi_matrix.asciidoc)
+- Updated [`sycl_ext_oneapi_matrix`](doc/extensions/experimental/sycl_ext_oneapi_matrix.asciidoc)
   with description of new matrix features. [770f540d]
-- Moved [`sycl_ext_oneapi_invoke_simd`](sycl/doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc)
+- Moved [`sycl_ext_oneapi_invoke_simd`](doc/extensions/experimental/sycl_ext_oneapi_invoke_simd.asciidoc)
   extensions specification from `proposed` to `experimental` because
   implementation is available. [6bee3440]
 

sycl/doc/EnvironmentVariables.md

@@ -90,7 +90,7 @@ fixed order of properties in the pattern.
 This environment variable limits the SYCL RT to use only a subset of the system's devices. Setting this environment variable affects all of the device query functions (`platform::get_devices()` and `platform::get_platforms()`) and all of the device selectors. 
 
 The value of this environment variable is a comma separated list of filters, where each filter is a triple of the form "`backend`:`device_type`:`device_num`" (without the quotes). Each element of the triple is optional, but each filter must have at least one value. Possible values of `backend` are:
-- `host`
+- `host` (Deprecated)
 - `level_zero`
 - `opencl`
 - `cuda`
@@ -99,17 +99,17 @@ The value of this environment variable is a comma separated list of filters, whe
 - `*`
 
 Possible values of `device_type` are:
-- `host`
+- `host` (Deprecated)
 - `cpu`
 - `gpu`
 - `acc`
 - `*`
 
 `device_num` is an integer that indexes the enumeration of devices from the sycl-ls utility tool, where the first device in that enumeration has index zero in each backend. For example, `SYCL_DEVICE_FILTER=2` will return all devices with index '2' from all different backends. If multiple devices satisfy this device number (e.g., GPU and CPU devices can be assigned device number '2'), then default_selector will choose the device with the highest heuristic point.  When `SYCL_DEVICE_ALLOWLIST` is set, it is applied before enumerating devices and affects `device_num` values.
 
-Assuming a filter has all three elements of the triple, it selects only those devices that come from the given backend, have the specified device type, AND have the given device index. If more than one filter is specified, the RT is restricted to the union of devices selected by all filters. The RT does not include the `host` backend and the `host` device automatically unless one of the filters explicitly specifies the `host` device type. Therefore, `SYCL_DEVICE_FILTER=host` should be set to enforce SYCL to use the `host` device only. 
+Assuming a filter has all three elements of the triple, it selects only those devices that come from the given backend, have the specified device type, AND have the given device index. If more than one filter is specified, the RT is restricted to the union of devices selected by all filters.
 
-Note that all device selectors will throw an exception if the filtered list of devices does not include a device that satisfies the selector. For instance, `SYCL_DEVICE_FILTER=cpu,level_zero` will cause `host_selector()` to throw an exception. `SYCL_DEVICE_FILTER` also limits loading only specified plugins into the SYCL RT. In particular, `SYCL_DEVICE_FILTER=level_zero` will cause the `cpu_selector` to throw an exception since SYCL RT will only load the `level_zero` backend which does not support any CPU devices at this time. When multiple devices satisfy the filter (e..g, `SYCL_DEVICE_FILTER=gpu`), only one of them will be selected.
+Note that all device selectors will throw an exception if the filtered list of devices does not include a device that satisfies the selector. For instance, `SYCL_DEVICE_FILTER=cpu` will cause `gpu_selector()` to throw an exception. `SYCL_DEVICE_FILTER` also limits loading only specified plugins into the SYCL RT. In particular, `SYCL_DEVICE_FILTER=level_zero` will cause the `cpu_selector` to throw an exception since SYCL RT will only load the `level_zero` backend which does not support any CPU devices at this time. When multiple devices satisfy the filter (e..g, `SYCL_DEVICE_FILTER=gpu`), only one of them will be selected.
 
 ## `SYCL_REDUCTION_PREFERRED_WORKGROUP_SIZE`
 

sycl/doc/extensions/proposed/sycl_ext_intel_device_architecture.asciidoc renamed to sycl/doc/extensions/experimental/sycl_ext_intel_device_architecture.asciidoc

@@ -90,18 +90,25 @@ supports.
 
 |2
 |fpga_simulator_selector added.
+
+|3
+|SYCL 2020 selector variants `fpga_selector_v`, `fpga_simulator_selector_v`, and
+`fpga_emulator_selector_v` added.  Old selectors `fpga_selector`,
+`fpga_simulator_selector`, and `fpga_emulator_selector` deprecated.
 |===
 
 === Select FPGA hardware device
 ....
 // select FPGA hardware device
-sycl::queue deviceQueue{sycl::ext::intel::fpga_selector{}};
+sycl::queue deviceQueue1{sycl::ext::intel::fpga_selector{}}; // Deprecated
+sycl::queue deviceQueue2{sycl::ext::intel::fpga_selector_v};
 ....
 
 === Select FPGA simulator device
 ....
 // select FPGA simulator device
-sycl::queue deviceQueue{sycl::ext::intel::fpga_simulator_selector{}};
+sycl::queue deviceQueue1{sycl::ext::intel::fpga_simulator_selector{}}; // Deprecated
+sycl::queue deviceQueue2{sycl::ext::intel::fpga_simulator_selector_v};
 ....
 
 [NOTE]
@@ -112,14 +119,16 @@ Added in version 2 of this extension.
 === Select FPGA emulator device
 ....
 // select FPGA emulator device
-sycl::queue deviceQueue{sycl::ext::intel::fpga_emulator_selector{}};
+sycl::queue deviceQueue1{sycl::ext::intel::fpga_emulator_selector{}}; // Deprecated
+sycl::queue deviceQueue2{sycl::ext::intel::fpga_emulator_selector_v};
 ....
 
 == Implementation notes
 
 The current implementation has a restriction on the use of
-`fpga_simulator_selector`.  If an object of `fpga_simulator_selector` is
-defined in the application, FPGA hardware devices selected using
-`fpga_selector` will select a simulator device. This behaviour is expected to
+`fpga_simulator_selector` and `fpga_simulator_selector_v`.  If an object of
+`fpga_simulator_selector` is defined or `fpga_simulator_selector_v` is used in
+the application, FPGA hardware devices selected using fpga_selector` and
+`fpga_selector_v` will select a simulator device. This behaviour is expected to
 be eliminated in the future.
 

sycl/doc/extensions/supported/sycl_ext_intel_fpga_device_selector.asciidoc

@@ -28,7 +28,7 @@ enum class backend : char {
   cuda __SYCL2020_DEPRECATED("use 'ext_oneapi_cuda' instead") = ext_oneapi_cuda,
   all = 4,
   ext_intel_esimd_emulator = 5,
-  esimd_cpu __SYCL2020_DEPRECATED("use 'ext_oneapi_esimd_emulator' instead") =
+  esimd_cpu __SYCL2020_DEPRECATED("use 'ext_intel_esimd_emulator' instead") =
       ext_intel_esimd_emulator,
   ext_oneapi_hip = 6,
   hip __SYCL2020_DEPRECATED("use 'ext_oneapi_hip' instead") = ext_oneapi_hip,