Merge from 'sycl' to 'sycl-web' (6 commits)

Author: iclsrc

libclc/libspirv/lib/native_cpu-unknown-linux/cl_khr_int64_extended_atomics/minmax_helpers.ll

@@ -28,15 +28,15 @@ entry:
   ret i64 %0
 }
 
-define i64 @__clc__sync_fetch_and_min_generic_8(i64 addrspace(1)* nocapture %ptr, i64 %value) nounwind alwaysinline {
+define i64 @__clc__sync_fetch_and_min_generic_8(i64* nocapture %ptr, i64 %value) nounwind alwaysinline {
 entry:
-  %0 = atomicrmw volatile min i64 addrspace(1)* %ptr, i64 %value seq_cst
+  %0 = atomicrmw volatile min i64* %ptr, i64 %value seq_cst
   ret i64 %0
 }
 
-define i64 @__clc__sync_fetch_and_umin_generic_8(i64 addrspace(1)* nocapture %ptr, i64 %value) nounwind alwaysinline {
+define i64 @__clc__sync_fetch_and_umin_generic_8(i64* nocapture %ptr, i64 %value) nounwind alwaysinline {
 entry:
-  %0 = atomicrmw volatile umin i64 addrspace(1)* %ptr, i64 %value seq_cst
+  %0 = atomicrmw volatile umin i64* %ptr, i64 %value seq_cst
   ret i64 %0
 }
 
@@ -64,14 +64,14 @@ entry:
   ret i64 %0
 }
 
-define i64 @__clc__sync_fetch_and_max_generic_8(i64 addrspace(1)* nocapture %ptr, i64 %value) nounwind alwaysinline {
+define i64 @__clc__sync_fetch_and_max_generic_8(i64* nocapture %ptr, i64 %value) nounwind alwaysinline {
 entry:
-  %0 = atomicrmw volatile max i64 addrspace(1)* %ptr, i64 %value seq_cst
+  %0 = atomicrmw volatile max i64* %ptr, i64 %value seq_cst
   ret i64 %0
 }
 
-define i64 @__clc__sync_fetch_and_umax_generic_8(i64 addrspace(1)* nocapture %ptr, i64 %value) nounwind alwaysinline {
+define i64 @__clc__sync_fetch_and_umax_generic_8(i64* nocapture %ptr, i64 %value) nounwind alwaysinline {
 entry:
-  %0 = atomicrmw volatile umax i64 addrspace(1)* %ptr, i64 %value seq_cst
+  %0 = atomicrmw volatile umax i64* %ptr, i64 %value seq_cst
   ret i64 %0
 }

llvm/test/SYCLLowerIR/work_group_static.ll

@@ -38,5 +38,5 @@ attributes #1 = { convergent norecurse }
 !2 = !{i32 4, i32 100000}
 !3 = !{!"clang version 13.0.0"}
 !4 = !{}
-; ![[ADDR_SPACE_MD]] = !{i32 1, i32 3}
+; CHECK: ![[ADDR_SPACE_MD]] = !{i32 1, i32 3}
 !5 = !{i32 1}

llvm/test/SYCLLowerIR/work_group_static_nv.ll

@@ -8,7 +8,7 @@ target triple = "nvptx64-nvidia-cuda"
 ; CHECK: @__sycl_dynamicLocalMemoryPlaceholder_GV = external local_unnamed_addr addrspace(3) global [0 x i8], align 128
 
 ; Function Attrs: convergent norecurse
-; CHECK: @_ZTS7KernelA(ptr addrspace(1) %0)
+; CHECK: @_ZTS7KernelA(ptr addrspace(1) %0){{.*}} !kernel_arg_addr_space ![[ADDR_SPACE_MD:[0-9]+]]
 define void @_ZTS7KernelA(ptr addrspace(1) %0) local_unnamed_addr #0 !kernel_arg_addr_space !5 {
 entry:
 ; CHECK: getelementptr inbounds i8, ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder_GV
@@ -34,5 +34,5 @@ attributes #1 = { convergent norecurse }
 !2 = !{i32 4, i32 100000}
 !3 = !{!"clang version 13.0.0"}
 !4 = !{}
-; ![[ADDR_SPACE_MD]] = !{i32 1, i32 3}
+; CHECK: ![[ADDR_SPACE_MD]] = !{i32 1}
 !5 = !{i32 1}

sycl-jit/jit-compiler/lib/KernelFusion.cpp

@@ -17,7 +17,13 @@
 #include "rtc/DeviceCompilation.h"
 #include "translation/KernelTranslation.h"
 #include "translation/SPIRVLLVMTranslation.h"
+
+#include <llvm/ADT/StringExtras.h>
 #include <llvm/Support/Error.h>
+#include <llvm/Support/TimeProfiler.h>
+
+#include <clang/Driver/Options.h>
+
 #include <sstream>
 
 using namespace jit_compiler;
@@ -237,14 +243,34 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
 extern "C" KF_EXPORT_SYMBOL RTCResult
 compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
             View<const char *> UserArgs) {
+  std::string BuildLog;
+
   auto UserArgListOrErr = parseUserArgs(UserArgs);
   if (!UserArgListOrErr) {
     return errorTo<RTCResult>(UserArgListOrErr.takeError(),
                               "Parsing of user arguments failed");
   }
   llvm::opt::InputArgList UserArgList = std::move(*UserArgListOrErr);
 
-  std::string BuildLog;
+  llvm::StringRef TraceFileName;
+  if (auto *Arg =
+          UserArgList.getLastArg(clang::driver::options::OPT_ftime_trace_EQ)) {
+    TraceFileName = Arg->getValue();
+    int Granularity =
+        500; // microseconds. Same default as in `clang::FrontendOptions`.
+    if (auto *Arg = UserArgList.getLastArg(
+            clang::driver::options::OPT_ftime_trace_granularity_EQ)) {
+      if (!llvm::to_integer(Arg->getValue(), Granularity)) {
+        BuildLog += "warning: ignoring malformed argument: '" +
+                    Arg->getAsString(UserArgList) + "'\n";
+      }
+    }
+    bool Verbose =
+        UserArgList.hasArg(clang::driver::options::OPT_ftime_trace_verbose);
+
+    llvm::timeTraceProfilerInitialize(Granularity, /*ProcName=*/"sycl-rtc",
+                                      Verbose);
+  }
 
   auto ModuleOrErr =
       compileDeviceCode(SourceFile, IncludeFiles, UserArgList, BuildLog);
@@ -279,6 +305,15 @@ compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
     DevImgInfo.BinaryInfo = std::move(*BinaryInfoOrError);
   }
 
+  if (llvm::timeTraceProfilerEnabled()) {
+    auto Error = llvm::timeTraceProfilerWrite(
+        TraceFileName, /*FallbackFileName=*/"trace.json");
+    llvm::timeTraceProfilerCleanup();
+    if (Error) {
+      return errorTo<RTCResult>(std::move(Error), "Trace file writing failed");
+    }
+  }
+
   return RTCResult{std::move(BundleInfo), BuildLog.c_str()};
 }
 

sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp

@@ -33,6 +33,7 @@
 #include <llvm/SYCLLowerIR/ModuleSplitter.h>
 #include <llvm/SYCLLowerIR/SYCLJointMatrixTransform.h>
 #include <llvm/Support/PropertySetIO.h>
+#include <llvm/Support/TimeProfiler.h>
 
 #include <algorithm>
 #include <array>
@@ -225,6 +226,8 @@ class LLVMDiagnosticWrapper : public llvm::DiagnosticHandler {
 Expected<std::unique_ptr<llvm::Module>> jit_compiler::compileDeviceCode(
     InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
     const InputArgList &UserArgList, std::string &BuildLog) {
+  TimeTraceScope TTS{"compileDeviceCode"};
+
   const std::string &DPCPPRoot = getDPCPPRoot();
   if (DPCPPRoot == InvalidDPCPPRoot) {
     return createStringError("Could not locate DPCPP root directory");
@@ -244,6 +247,9 @@ Expected<std::unique_ptr<llvm::Module>> jit_compiler::compileDeviceCode(
   // linking).
   DAL.eraseArg(OPT_fsycl_device_lib_EQ);
   DAL.eraseArg(OPT_fno_sycl_device_lib_EQ);
+  DAL.eraseArg(OPT_ftime_trace_EQ);
+  DAL.eraseArg(OPT_ftime_trace_granularity_EQ);
+  DAL.eraseArg(OPT_ftime_trace_verbose);
 
   SmallVector<std::string> CommandLine;
   for (auto *Arg : DAL) {
@@ -382,6 +388,8 @@ static bool getDeviceLibraries(const ArgList &Args,
 Error jit_compiler::linkDeviceLibraries(llvm::Module &Module,
                                         const InputArgList &UserArgList,
                                         std::string &BuildLog) {
+  TimeTraceScope TTS{"linkDeviceLibraries"};
+
   const std::string &DPCPPRoot = getDPCPPRoot();
   if (DPCPPRoot == InvalidDPCPPRoot) {
     return createStringError("Could not locate DPCPP root directory");
@@ -458,6 +466,8 @@ static IRSplitMode getDeviceCodeSplitMode(const InputArgList &UserArgList) {
 Expected<PostLinkResult>
 jit_compiler::performPostLink(std::unique_ptr<llvm::Module> Module,
                               const InputArgList &UserArgList) {
+  TimeTraceScope TTS{"performPostLink"};
+
   // This is a simplified version of `processInputModule` in
   // `llvm/tools/sycl-post-link.cpp`. Assertions/TODOs point to functionality
   // left out of the algorithm for now.

sycl-jit/jit-compiler/lib/translation/KernelTranslation.cpp

@@ -17,6 +17,7 @@
 #include "llvm/MC/TargetRegistry.h"
 #include "llvm/Support/MemoryBuffer.h"
 #include "llvm/Support/TargetSelect.h"
+#include "llvm/Support/TimeProfiler.h"
 #include "llvm/Target/TargetMachine.h"
 #include "llvm/Target/TargetOptions.h"
 
@@ -225,6 +226,8 @@ llvm::Error KernelTranslator::translateKernel(SYCLKernelInfo &Kernel,
 llvm::Expected<RTCDevImgBinaryInfo>
 KernelTranslator::translateDevImgToSPIRV(llvm::Module &Mod,
                                          JITContext &JITCtx) {
+  llvm::TimeTraceScope TTS{"translateDevImgToSPIRV"};
+
   llvm::Expected<KernelBinary *> BinaryOrError = translateToSPIRV(Mod, JITCtx);
   if (auto Error = BinaryOrError.takeError()) {
     return Error;

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -1,7 +1,7 @@
-# commit f058cb230c65fe8094f74043d0c9afd5ba0e8325
-# Merge: 871061f1 3bf76246
-# Author: Ross Brunton <ross@codeplay.com>
-# Date:   Tue Jan 21 15:49:41 2025 +0000
-#     Merge pull request #2593 from RossBrunton/ross/formatall
-#     Ensure all files are clang formatted
-set(UNIFIED_RUNTIME_TAG f058cb230c65fe8094f74043d0c9afd5ba0e8325)
+# commit d3e97040b67ee6b45655eccdfca19185bd9115b6
+# Merge: 2ef69591 37ad03a7
+# Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
+# Date:   Thu Jan 23 10:23:08 2025 +0000
+#     Merge pull request #2605 from yingcong-wu/yc/0123-ur-quick-fix
+#     [DeviceASAN] Bugfix for GetDeviceType
+set(UNIFIED_RUNTIME_TAG d3e97040b67ee6b45655eccdfca19185bd9115b6)

sycl/doc/extensions/proposed/sycl_ext_oneapi_device_image_backend_content.asciidoc

@@ -101,9 +101,7 @@ class device_image {
   backend ext_oneapi_get_backend() const noexcept;
   std::vector<std::byte> ext_oneapi_get_backend_content() const;
 
-#if defined(__cpp_lib_span)
-  std::span<std::byte> ext_oneapi_get_backend_content_view() const;
-#endif
+  std::span<std::byte> ext_oneapi_get_backend_content_view() const; // Requires C++20
 
   /*...*/
 };
@@ -154,8 +152,7 @@ std::span<std::byte> ext_oneapi_get_content_backend_view() const;
 ----
 !====
 
-Available only when the compiler defines the `__cpp_lib_span` feature-test macro
-(which is defined in {cpp}20 and higher).
+Minimum C++ Version: {cpp}20
 
 _Constraints:_ Available only when `State` is `bundle_state::executable`.
 

sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit_time_trace.cpp

@@ -0,0 +1,68 @@
+//==----- kernel_compiler_sycl_jit_time_trace.cpp --- time-tracing test ----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: (opencl || level_zero)
+// UNSUPPORTED: accelerator
+// UNSUPPORTED-INTENDED: SYCL-RTC is not available for accelerator devices
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+
+int test_tracing() {
+  namespace syclex = sycl::ext::oneapi::experimental;
+  using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
+
+  sycl::queue q;
+  sycl::context ctx = q.get_context();
+
+  bool ok =
+      q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl_jit);
+  if (!ok) {
+    std::cout << "Apparently this device does not support `sycl_jit` source "
+                 "kernel bundle extension: "
+              << q.get_device().get_info<sycl::info::device::name>()
+              << std::endl;
+    return -1;
+  }
+
+  source_kb kbSrc = syclex::create_kernel_bundle_from_source(
+      ctx, syclex::source_language::sycl_jit, "");
+
+  auto props = syclex::properties{
+      syclex::build_options{std::vector<std::string>{
+          "-ftime-trace=-", "-ftime-trace-granularity=1000" /* us */,
+          "-ftime-trace-verbose"}},
+  };
+
+  syclex::build(kbSrc, props);
+  // CHECK: {"traceEvents":
+
+  std::string log;
+  auto props2 = syclex::properties{
+      syclex::build_options{std::vector<std::string>{
+          "-ftime-trace=-", "-ftime-trace-granularity=invalid_int"}},
+      syclex::save_log{&log}};
+  syclex::build(kbSrc, props2);
+  std::cout << log << std::endl;
+  // CHECK: {"traceEvents":
+  // CHECK: warning: ignoring malformed argument: '-ftime-trace-granularity=invalid_int'
+
+  return 0;
+}
+
+int main() {
+#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER
+  return test_tracing();
+#else
+  static_assert(false, "Kernel Compiler feature test macro undefined");
+#endif
+  return 0;
+}