Merge branch 'sycl' into root_group_query

Author: 0x12CC

clang/lib/Sema/SemaSYCL.cpp

@@ -6969,6 +6969,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
   Policy.adjustForCPlusPlusFwdDecl();
   Policy.SuppressTypedefs = true;
   Policy.SuppressUnwrittenScope = true;
+  Policy.PrintCanonicalTypes = true;
 
   llvm::SmallSet<const VarDecl *, 8> Visited;
   bool EmittedFirstSpecConstant = false;

clang/test/CodeGenSYCL/int_footer_with_explicit_specialization.cpp

@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -emit-llvm %s -o -
+// RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER
+
+// This test checks that integration footer is emitted correctly when a
+// device_global has an explicit template specialization in template arguments.
+
+#include "sycl.hpp"
+
+namespace sycl {
+template <typename T> struct X {};
+template <> struct X<int> {};
+namespace detail {
+struct Y {};
+} // namespace detail
+template <> struct X<detail::Y> {};
+} // namespace sycl
+
+using namespace sycl;
+template <typename T, typename = X<detail::Y>> struct Arg1 { T val; };
+
+using namespace sycl::ext::oneapi;
+template <typename properties_t>
+device_global<properties_t> dev_global;
+
+SYCL_EXTERNAL auto foo() {
+  (void)dev_global<Arg1<int>>;
+}
+
+// CHECK-FOOTER: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::dev_global<Arg1<int, sycl::X<sycl::detail::Y>>>, "_Z10dev_globalI4Arg1IiN4sycl1XINS1_6detail1YEEEEE");
+// CHECK-FOOTER-NEXT: }
+// CHECK-FOOTER-NEXT: } // namespace (unnamed)

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -1,7 +1,7 @@
-# commit 2f479e01f62ed96a9466d19d9df1aabeda77b335
-# Merge: b0f9293d b693389e
+# commit 38ee6ce2a0400573c0c7c5da782bc32ff578fcc4
+# Merge: c7086f7f 91b6db04
 # Author: Callum Fare <[email protected]>
-# Date:   Fri Nov 22 15:27:05 2024 +0000
-#     Merge pull request #2370 from againull/multi_device_sanitizer
-#     [L0] Check that program is in exe state in urProgramGetGlobalVariablePointer
-set(UNIFIED_RUNTIME_TAG f99adf104ef5ce94d66b6ed4f9dccc450f1f638c)
+# Date:   Mon Nov 25 11:04:37 2024 +0000
+#     Merge pull request #2303 from nrspruit/zeInitDrivers
+#     [L0] Implement Support for zeInitDrivers
+set(UNIFIED_RUNTIME_TAG 39b72622cf2997459788000ac665e259f728a9ff)

sycl/doc/design/CommandGraph.md

@@ -337,6 +337,62 @@ Backends which are implemented currently are: [Level Zero](#level-zero),
 
 ### Level Zero
 
+The command-buffer implementation for the level-zero adapter has 2 different
+implementation paths which are chosen depending on the device and level-zero
+version:
+
+- Immediate Append path - Relies on
+  [zeCommandListImmediateAppendCommandListsExp](https://oneapi-src.github.io/level-zero-spec/level-zero/latest/core/api.html#zecommandlistimmediateappendcommandlistsexp)
+  to submit the command-buffer. This function is an experimental extension to the level-zero API.
+- Wait event path - Relies on
+  [zeCommandQueueExecuteCommandLists](https://oneapi-src.github.io/level-zero-spec/level-zero/latest/core/api.html#zecommandqueueexecutecommandlists)
+  to submit the command-buffer work. However, this level-zero function has
+  limitations and, as such, this path is used only when the immediate append
+  path is unavailable.
+
+#### Immediate Append Path Implementation Details
+
+This path is only available when the device supports immediate command-lists
+and the [zeCommandListImmediateAppendCommandListsExp](https://oneapi-src.github.io/level-zero-spec/level-zero/latest/core/api.html#zecommandlistimmediateappendcommandlistsexp) 
+API. This API can wait on a list of event dependencies using the `phWaitEvents`
+parameter and can signal a return event when finished using the `hSignalEvent`
+parameter. This allows for a cleaner and more efficient implementation than
+what can be achieved when using the wait-event path
+(see [this section](#wait-event-path-implementation-details) for
+more details about the wait-event path).
+
+This path relies on 3 different command-lists in order to execute the
+command-buffer:
+
+- `ComputeCommandList` - Used to submit command-buffer work that requires
+the compute engine.
+- `CopyCommandList` - Used to submit command-buffer work that requires the
+[copy engine](#copy-engine). This command-list is not created when none of the 
+nodes require the copy engine.
+- `EventResetCommandList` - Used to reset the level-zero events that are
+needed for every submission of the command-buffer. This is executed after
+the compute and copy command-lists have finished executing. For the first
+execution, this command-list is skipped since there is no need to reset events
+at this point. When counter-based events are enabled (i.e. the command-buffer
+is in-order), this command-list is not created since counter-based events do
+not need to be reset.
+
+The following diagram illustrates which commands are executed on
+each command-list when the command-buffer is enqueued:
+![L0 command-buffer diagram](images/diagram_immediate_append.png)
+
+Additionally,
+[zeCommandListImmediateAppendCommandListsExp](https://oneapi-src.github.io/level-zero-spec/level-zero/latest/core/api.html#zecommandlistimmediateappendcommandlistsexp)
+requires an extra command-list which is used to submit the other
+command-lists. This command-list has a specific engine type
+associated to it (i.e. compute or copy engine). Hence, for our implementation,
+we need 2 of these helper command-lists:
+  - The `CommandListHelper` command-list is used to submit the
+`ComputeCommandList`, `CommandListResetEvents` and profiling queries.
+  - The `ZeCopyEngineImmediateListHelper`  command-list is used to submit the
+`CopyCommandList`
+
+#### Wait event Path Implementation Details
 The UR `urCommandBufferEnqueueExp` interface for submitting a command-buffer
 takes a list of events to wait on, and returns an event representing the
 completion of that specific submission of the command-buffer.
@@ -364,7 +420,7 @@ is made only once (during the command-buffer finalization stage). This allows
 the adapter to save time when submitting the command-buffer, by executing only
 this command-list (i.e. without enqueuing any commands of the graph workload).
 
-#### Prefix
+##### Prefix
 
 The prefix's commands aim to:
 1. Handle the list of events to wait on, which is passed by the runtime
@@ -409,7 +465,7 @@ and another reset command for resetting the signal we use to signal the
 completion of the graph workload. This signal is called *SignalEvent* and is
 defined in the `ur_exp_command_buffer_handle_t` class.
 
-#### Suffix
+##### Suffix
 
 The suffix's commands aim to:
 1) Handle the completion of the graph workload and signal a UR return event.
@@ -435,7 +491,7 @@ with extra commands associated with *CB*, and the other after *CB*. These new
 command-lists are retrieved from the UR queue, which will likely reuse existing
 command-lists and only create a new one in the worst case.
 
-#### Drawbacks
+##### Drawbacks
 
 There are three drawbacks of this approach to implementing UR command-buffers for
 Level Zero:

sycl/doc/design/images/diagram_immediate_append.png

@@ -1,4 +1,4 @@
-// REQUIRES: linux
+// REQUIRES: linux, cpu || (gpu && level_zero)
 // RUN: %{build} %device_asan_flags -O0 -g -o %t
 // RUN: %force_device_asan_rt %{run} not %t 2>&1 | FileCheck %s
 #include <sycl/usm.hpp>

sycl/test-e2e/AddressSanitizer/bad-free/bad-free-host.cpp

@@ -1,4 +1,4 @@
-// REQUIRES: linux
+// REQUIRES: linux, cpu || (gpu && level_zero)
 // RUN: %{build} %device_asan_flags -O0 -g -o %t1.out
 // RUN: %force_device_asan_rt %{run} not %t1.out 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_HOST -O0 -g -o %t2.out

sycl/test-e2e/AddressSanitizer/bad-free/bad-free-minus1.cpp

@@ -1,4 +1,4 @@
-// REQUIRES: linux
+// REQUIRES: linux, cpu || (gpu && level_zero)
 // RUN: %{build} %device_asan_flags -O0 -g -o %t1.out
 // RUN: %force_device_asan_rt %{run} not %t1.out 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_HOST -O0 -g -o %t2.out

sycl/test-e2e/AddressSanitizer/bad-free/bad-free-plus1.cpp

@@ -1,4 +1,4 @@
-// REQUIRES: linux
+// REQUIRES: linux, cpu || (gpu && level_zero)
 // RUN: %{build} %device_asan_flags -O2 -g -o %t
 // RUN: %{run} not %t &> %t.txt ; FileCheck --input-file %t.txt %s
 #include <sycl/detail/core.hpp>

sycl/test-e2e/AddressSanitizer/common/demangle-kernel-name.cpp

@@ -1,6 +1,5 @@
-// REQUIRES: linux
+// REQUIRES: linux, cpu || (gpu && level_zero)
 // RUN: %{build} %device_asan_flags -O2 -g -o %t
-// RUN: env UR_LAYER_ASAN_OPTIONS=debug:1 %{run} %t 2>&1 | FileCheck --check-prefixes CHECK-DEBUG %s
 // RUN: env UR_LAYER_ASAN_OPTIONS=debug:0 %{run} %t 2>&1 | FileCheck %s
 
 #include <sycl/usm.hpp>
@@ -23,7 +22,6 @@ int main() {
         });
   });
   Q.wait();
-  // CHECK-DEBUG: [kernel]
   // CHECK-NOT: [kernel]
 
   sycl::free(array, Q);