Merge remote-tracking branch 'origin/sycl' into aaron/optionalInfoQueries

Author: kbenzie

.github/workflows/sycl-containers-igc-dev.yaml

@@ -6,11 +6,13 @@ on:
       - sycl
     paths:
       - 'devops/actions/build_container/**'
+      - 'devops/scripts/**'
       - 'devops/dependencies-igc-dev.json'
       - '.github/workflows/sycl-containers-igc-dev.yaml'
   pull_request:
     paths:
       - 'devops/actions/build_container/**'
+      - 'devops/scripts/**'
       - 'devops/dependencies-igc-dev.json'
       - '.github/workflows/sycl-containers-igc-dev.yaml'
 

.github/workflows/sycl-containers.yaml

@@ -78,6 +78,7 @@ jobs:
           file: ${{ matrix.file }}
           username: ${{ github.repository_owner }}
           password: ${{ secrets.GITHUB_TOKEN }}
+          sycl_ci_passwd: ${{ secrets.DOCKER_SUDO_PASSWORD }}
           tags: |
             ghcr.io/${{ github.repository }}/${{ matrix.file }}:${{ matrix.tag }}-${{ github.sha }}
             ghcr.io/${{ github.repository }}/${{ matrix.file }}:${{ matrix.tag }}

.github/workflows/sycl-linux-precommit.yml

@@ -134,7 +134,8 @@ jobs:
       install_dev_igc_driver: >-
         ${{ !contains(matrix.target_devices, 'ext_oneapi_cuda') &&
         !contains(matrix.target_devices, 'ext_oneapi_hip') &&
-        matrix.use_igc_dev && contains(needs.detect_changes.outputs.filters, 'devigccfg') ||
+        matrix.use_igc_dev &&
+        (contains(needs.detect_changes.outputs.filters, 'devigccfg') || contains(needs.detect_changes.outputs.filters, 'drivers')) ||
         'false' }}
       # Run only if the PR does not have the 'ci-no-devigc' label.
       skip_run: ${{matrix.use_igc_dev && contains(github.event.pull_request.labels.*.name, 'ci-no-devigc') || 'false'}}

devops/actions/build_container/action.yml

@@ -21,6 +21,9 @@ inputs:
   file:
     description: "Dockerfile"
     required: true
+  sycl_ci_passwd:
+    description: "Password to assign to sycl_ci user within a container"
+    required: true
 
 runs:
   using: "composite"
@@ -41,3 +44,4 @@ runs:
       file: ${{ github.workspace }}/devops/containers/${{ inputs.file }}.Dockerfile
       secrets: |
         github_token=${{ github.token }}
+        sycl_ci_passwd=${{ inputs.sycl_ci_passwd }}

devops/dependencies-igc-dev.json

@@ -1,10 +1,10 @@
 {
   "linux": {
     "igc_dev": {
-      "github_tag": "igc-dev-3db59df",
-      "version": "3db59df",
-      "updated_at": "2024-12-03T20:43:00Z",
-      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2248119261/zip",
+      "github_tag": "igc-dev-e0d826a",
+      "version": "e0d826a",
+      "updated_at": "2024-12-17T21:18:30Z",
+      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2327583926/zip",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     }
   }

devops/scripts/install_drivers.sh

@@ -162,6 +162,9 @@ InstallIGFX () {
     echo "Install libopencl-clang"
     # Workaround only, will download deb and install with dpkg once fixed.
     cp -d libopencl-clang.so.14*  /usr/local/lib/
+    rm /usr/local/lib/libigc.so /usr/local/lib/libigc.so.1* && \
+       ln -s /usr/local/lib/libigc.so.2 /usr/local/lib/libigc.so && \
+       ln -s /usr/local/lib/libigc.so.2 /usr/local/lib/libigc.so.1
     echo "Clean up"
     rm *.deb libopencl-clang.so.14*
     echo "$IGC_DEV_TAG" > /usr/local/lib/igc/IGCTAG.txt

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -189,11 +189,26 @@ class IntelTargetInfo<string Name, list<Aspect> Aspects, list<int> subGroupSizes
 // Note: only the "canonical" target names are listed here - see
 // SYCL::gen::resolveGenDevice().
 //
-// TODO: instructions on how to get "???" values below are required. Currently
-// device architectures below which have "???" in it are not fully supported.
-//
-// def : IntelTargetInfo<"intel_gpu_ptl_u", ???, ???>;
-// def : IntelTargetInfo<"intel_gpu_ptl_h", ???, ???>;
+// When adding the new Intel architectures, the ideal way to get these 
+// values is to run sycl-ls --verbose on the target machine. However,
+// if you don't have access to the target machine, then you can look in
+// the intel/compute-runtime source code to get the values. For 
+// the fp64 and atomic64 aspects, look for the definition 
+// of specific RuntimeCapabilityTable's. These will be defined in 
+// a file called hw_info_<device>.cpp. For example, for PVC, it 
+// would be hw_info_pvc.cpp. For fp16, this is supported by all 
+// Intel GPUs currently. (see definition of getDeviceExtensions in 
+// compiler_product_helper_base.inl). For the supported sub-group sizes,
+// it is not as straightforward. They are defined per family in 
+// GfxCoreHelperHw<Family>::getDeviceSubGroupSizes and by default
+// the default sub-group sizes are 8, 16, and 32. (see gfx_core_helper_base.inl)
+// However, this can be overriden by specific families. For example,
+// gfx_core_helper_xe_hpc_core.cpp defines the supported sub-group 
+// sizes to be 16 and 32, which corresponds to PVC having 16 and 32 
+// as sub-group sizes the supported sub-group sizes. 
+
+def : IntelTargetInfo<"intel_gpu_ptl_u", Fp16Fp64Atomic64, Sg16_32>;
+def : IntelTargetInfo<"intel_gpu_ptl_h", Fp16Fp64Atomic64, Sg16_32>;
 def : IntelTargetInfo<"intel_gpu_lnl_m", Fp16Fp64Atomic64, Sg16_32>;
 def : IntelTargetInfo<"intel_gpu_bmg_g21", Fp16Fp64Atomic64, Sg16_32>;
 def : IntelTargetInfo<"intel_gpu_arl_h", Fp16Fp64Atomic64, Sg8_16_32>;

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -2075,6 +2075,45 @@ recording mode, as opposed to throwing.
 This section defines the interaction of `sycl_ext_oneapi_graph` with other
 extensions.
 
+==== sycl_ext_codeplay_enqueue_native_command
+
+`ext_codeplay_enqueue_native_command`, defined in
+link:../experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc[sycl_ext_codeplay_enqueue_native_command]
+cannot be used in graph nodes. A synchronous exception will be thrown with error
+code `invalid` if a user tries to add them to a graph.
+
+Removing this restriction is something we may look at for future revisions of
+`sycl_ext_oneapi_graph`.
+
+==== sycl_ext_intel_queue_index
+
+The compute index queue property defined by
+link:../supported/sycl_ext_intel_queue_index.asciidoc[sycl_ext_intel_queue_index]
+is ignored during queue recording.
+
+Using this information is something we may look at for future revisions of
+`sycl_ext_oneapi_graph`.
+
+==== sycl_ext_oneapi_bindless_images
+
+The new handler methods, and queue shortcuts, defined by
+link:../experimental/sycl_ext_oneapi_bindless_images.asciidoc[sycl_ext_oneapi_bindless_images]
+cannot be used in graph nodes. A synchronous exception will be thrown with error
+code `invalid` if a user tries to add them to a graph.
+
+Removing this restriction is something we may look at for future revisions of
+`sycl_ext_oneapi_graph`.
+
+==== sycl_ext_oneapi_device_global
+
+The new handler methods, and queue shortcuts, defined by
+link:../experimental/sycl_ext_oneapi_device_global.asciidoc[sycl_ext_oneapi_device_global].
+cannot be used in graph nodes. A synchronous exception will be thrown with error
+code `invalid` if a user tries to add them to a graph.
+
+Removing this restriction is something we may look at for future revisions of
+`sycl_ext_oneapi_graph`.
+
 ==== sycl_ext_oneapi_discard_queue_events
 
 When recording a `sycl::queue` which has been created with the
@@ -2108,37 +2147,25 @@ nodes that are recorded from multiple queues and/or added by the explicit API:
 * The only commands which have an implicit dependency on the barrier command
   are those recorded from the same queue the barrier command was submitted to.
 
-==== sycl_ext_oneapi_memcpy2d
-
-The new handler methods, and queue shortcuts, defined by
-link:../supported/sycl_ext_oneapi_memcpy2d.asciidoc[sycl_ext_oneapi_memcpy2d]
-cannot be used in graph nodes. A synchronous exception will be thrown with
-error code `invalid` if a user tries to add them to a graph.
-
-Removing this restriction is something we may look at for future revisions of
-`sycl_ext_oneapi_graph`.
-
-==== sycl_ext_oneapi_queue_priority
-
-The queue priority property defined by
-link:../supported/sycl_ext_oneapi_queue_priority.asciidoc[sycl_ext_oneapi_queue_priority]
-is ignored during queue recording.
+==== sycl_ext_oneapi_enqueue_functions
 
-==== sycl_ext_oneapi_queue_empty
+The command submission functions defined in
+link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions]
+can be used adding nodes to a graph when creating a graph from queue recording.
+New methods are also defined that enable submitting an executable graph,
+e.g. directly to a queue without returning an event.
 
-The `queue::ext_oneapi_empty()` query defined by the
-link:../supported/sycl_ext_oneapi_queue_empty.asciidoc[sycl_ext_oneapi_queue_empty]
-extension behaves as normal during queue recording and is not captured to the graph.
-Recorded commands are not counted as submitted for the purposes of this query.
+==== sycl_ext_oneapi_free_function_kernels
 
-==== sycl_ext_intel_queue_index
+`sycl_ext_oneapi_free_function_kernels`, defined in
+link:../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc[sycl_ext_oneapi_free_function_kernels]
+can be used with SYCL Graphs.
 
-The compute index queue property defined by
-link:../supported/sycl_ext_intel_queue_index.asciidoc[sycl_ext_intel_queue_index]
-is ignored during queue recording.
+==== sycl_ext_oneapi_kernel_compiler_spirv
 
-Using this information is something we may look at for future revisions of
-`sycl_ext_oneapi_graph`.
+The kernels loaded using
+link:../experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc[sycl_ext_oneapi_kernel_compiler_spirv]
+behave as normal when used in graph nodes.
 
 ==== sycl_ext_oneapi_kernel_properties
 
@@ -2147,62 +2174,41 @@ link:../experimental/sycl_ext_oneapi_kernel_properties.asciidoc[sycl_ext_oneapi_
 can be used in graph nodes in the same way as they are used in normal queue
 submission. 
 
-==== sycl_ext_oneapi_prod
+==== sycl_ext_oneapi_local_memory
 
-The new `sycl::queue::ext_oneapi_prod()` method added by
-link:../proposed/sycl_ext_oneapi_prod.asciidoc[sycl_ext_oneapi_prod]
-behaves as normal during queue recording and is not captured to the graph.
-Recorded commands are not counted as submitted for the purposes of its operation.
+Allocating local memory inside a graph kernel node with `sycl::ext::oneapi::group_local_memory()` or
+`sycl::ext::oneapi::group_local_memory_for_overwrite()` is supported. These methods are defined by
+link:../supported/sycl_ext_oneapi_local_memory.asciidoc[sycl_ext_oneapi_local_memory.]
 
-==== sycl_ext_oneapi_device_global
-
-The new handler methods, and queue shortcuts, defined by
-link:../experimental/sycl_ext_oneapi_device_global.asciidoc[sycl_ext_oneapi_device_global].
-cannot be used in graph nodes. A synchronous exception will be thrown with error
-code `invalid` if a user tries to add them to a graph.
-
-Removing this restriction is something we may look at for future revisions of
-`sycl_ext_oneapi_graph`.
-
-==== sycl_ext_oneapi_bindless_images
+==== sycl_ext_oneapi_memcpy2d
 
 The new handler methods, and queue shortcuts, defined by
-link:../experimental/sycl_ext_oneapi_bindless_images.asciidoc[sycl_ext_oneapi_bindless_images]
-cannot be used in graph nodes. A synchronous exception will be thrown with error
-code `invalid` if a user tries to add them to a graph.
+link:../supported/sycl_ext_oneapi_memcpy2d.asciidoc[sycl_ext_oneapi_memcpy2d]
+cannot be used in graph nodes. A synchronous exception will be thrown with
+error code `invalid` if a user tries to add them to a graph.
 
 Removing this restriction is something we may look at for future revisions of
 `sycl_ext_oneapi_graph`.
 
-==== sycl_ext_oneapi_kernel_compiler_spirv
-
-The kernels loaded using
-link:../experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc[sycl_ext_oneapi_kernel_compiler_spirv]
-behave as normal when used in graph nodes.
-
-==== sycl_ext_codeplay_enqueue_native_command
-
-`ext_codeplay_enqueue_native_command`, defined in
-link:../experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc[sycl_ext_codeplay_enqueue_native_command]
-cannot be used in graph nodes. A synchronous exception will be thrown with error
-code `invalid` if a user tries to add them to a graph.
+==== sycl_ext_oneapi_prod
 
-Removing this restriction is something we may look at for future revisions of
-`sycl_ext_oneapi_graph`.
+The new `sycl::queue::ext_oneapi_prod()` method added by
+link:../proposed/sycl_ext_oneapi_prod.asciidoc[sycl_ext_oneapi_prod]
+behaves as normal during queue recording and is not captured to the graph.
+Recorded commands are not counted as submitted for the purposes of its operation.
 
-==== sycl_ext_oneapi_enqueue_functions
+==== sycl_ext_oneapi_queue_empty
 
-The command submission functions defined in
-link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions]
-can be used adding nodes to a graph when creating a graph from queue recording.
-New methods are also defined that enable submitting an executable graph,
-e.g. directly to a queue without returning an event.
+The `queue::ext_oneapi_empty()` query defined by the
+link:../supported/sycl_ext_oneapi_queue_empty.asciidoc[sycl_ext_oneapi_queue_empty]
+extension behaves as normal during queue recording and is not captured to the graph.
+Recorded commands are not counted as submitted for the purposes of this query.
 
-==== sycl_ext_oneapi_free_function_kernels
+==== sycl_ext_oneapi_queue_priority
 
-`sycl_ext_oneapi_free_function_kernels`, defined in
-link:../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc[sycl_ext_oneapi_free_function_kernels]
-can be used with SYCL Graphs.
+The queue priority property defined by
+link:../supported/sycl_ext_oneapi_queue_priority.asciidoc[sycl_ext_oneapi_queue_priority]
+is ignored during queue recording.
 
 ==== sycl_ext_oneapi_work_group_memory
 

sycl/include/sycl/detail/type_traits.hpp

@@ -384,6 +384,22 @@ struct map_type<T, From, To, Rest...> {
 template <typename T, typename... Ts>
 constexpr bool check_type_in_v = ((std::is_same_v<T, Ts> || ...));
 
+#if __has_builtin(__type_pack_element)
+template <int N, typename... Ts>
+using nth_type_t = __type_pack_element<N, Ts...>;
+#else
+template <int N, typename T, typename... Ts> struct nth_type {
+  using type = typename nth_type<N - 1, Ts...>::type;
+};
+
+template <typename T, typename... Ts> struct nth_type<0, T, Ts...> {
+  using type = T;
+};
+
+template <int N, typename... Ts>
+using nth_type_t = typename nth_type<N, Ts...>::type;
+#endif
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl