Merge branch 'urlog' of github.com:mateuszpn/llvm into urlog

Author: mateuszpn

.github/workflows/sycl-ur-perf-benchmarking.yml

@@ -0,0 +1,12 @@
+name: Benchmarks
+
+# This workflow is a WIP: this workflow file acts as a placeholder.
+
+on: [ workflow_dispatch ]
+
+jobs:
+  do-nothing:
+    runs-on: ubuntu-latest
+    steps:
+      - run: echo 'This workflow is a WIP.'
+

.github/workflows/ur-benchmarks-reusable.yml

@@ -1,6 +1,8 @@
 name: Benchmarks Reusable
 
 # This workflow is a WIP: This workflow file acts as a placeholder.
+#
+# This workflow is set to be merged into benchmark.yml
 
 on: [ workflow_call ]
 

.github/workflows/ur-benchmarks.yml

@@ -1,6 +1,8 @@
 name: Benchmarks
 
 # This workflow is a WIP: this workflow file acts as a placeholder.
+#
+# This workflow is set to be merged into benchmark.yml
 
 on: [ workflow_dispatch ]
 

clang/test/Driver/linker-wrapper-sycl.cpp

@@ -47,11 +47,6 @@
 // RUN: clang-linker-wrapper -sycl-device-libraries=%t.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--triple=spir64" "--linker-path=/usr/bin/ld" -shared "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-SHARED %s
 // CHK-SHARED: "{{.*}}clang"{{.*}} -fPIC
 
-// RUN: rm %T/linker_wrapper_dump || true
-// RUN: clang-linker-wrapper -sycl-dump-device-code=%T/linker_wrapper_dump -sycl-device-libraries=%t.devicelib.o "--host-triple=x86_64-unknown-linux-gnu" "--triple=spir64" "--linker-path=/usr/bin/ld" -shared "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t.o --dry-run
-// RUN: ls %T/linker_wrapper_dump | FileCheck -check-prefix=CHK-SYCL-DUMP-DEVICE %s
-// CHK-SYCL-DUMP-DEVICE: {{.*}}.spv
-
 /// Check for list of commands for standalone clang-linker-wrapper run for sycl (AOT for Intel GPU)
 // -------
 // Generate .o file as linker wrapper input.

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

@@ -151,7 +151,7 @@ static std::optional<llvm::module_split::IRSplitMode> SYCLModuleSplitMode;
 
 static bool UseSYCLPostLinkTool;
 
-static SmallString<128> SPIRVDumpDir;
+static SmallString<128> OffloadImageDumpDir;
 
 using OffloadingImage = OffloadBinary::OffloadingImage;
 
@@ -923,29 +923,6 @@ static Expected<StringRef> runLLVMToSPIRVTranslation(StringRef File,
   if (Error Err = executeCommands(*LLVMToSPIRVPath, CmdArgs))
     return std::move(Err);
 
-  if (!SPIRVDumpDir.empty()) {
-    std::error_code EC =
-        llvm::sys::fs::create_directory(SPIRVDumpDir, /*IgnoreExisting*/ true);
-    if (EC)
-      return createStringError(
-          EC,
-          formatv("failed to create dump directory. path: {0}, error_code: {1}",
-                  SPIRVDumpDir, EC.value()));
-
-    StringRef Sep = llvm::sys::path::get_separator();
-    StringRef Path = *TempFileOrErr;
-    StringRef Filename = Path.rsplit(Sep).second;
-    SmallString<128> CopyPath = SPIRVDumpDir;
-    CopyPath.append(Filename);
-    EC = llvm::sys::fs::copy_file(Path, CopyPath);
-    if (EC)
-      return createStringError(
-          EC,
-          formatv(
-              "failed to copy file. original: {0}, copy: {1}, error_code: {2}",
-              Path, CopyPath, EC.value()));
-  }
-
   return *TempFileOrErr;
 }
 
@@ -1103,6 +1080,18 @@ wrapSYCLBinariesFromFile(std::vector<module_split::SplitModule> &SplitModules,
     RegularTarget = "spir64";
 
   for (auto &SI : SplitModules) {
+    if (!OffloadImageDumpDir.empty()) {
+      StringRef CopyFrom = SI.ModuleFilePath;
+      SmallString<128> CopyTo = OffloadImageDumpDir;
+      StringRef Filename = sys::path::filename(CopyFrom);
+      CopyTo.append(Filename);
+      std::error_code EC = sys::fs::copy_file(CopyFrom, CopyTo);
+      if (EC)
+        return createStringError(EC, formatv("failed to copy file. From: "
+                                             "{0} to: {1}, error_code: {2}",
+                                             CopyFrom, CopyTo, EC.value()));
+    }
+
     auto MBOrDesc = MemoryBuffer::getFile(SI.ModuleFilePath);
     if (!MBOrDesc)
       return createFileError(SI.ModuleFilePath, MBOrDesc.getError());
@@ -2618,13 +2607,11 @@ int main(int Argc, char **Argv) {
 
   if (Args.hasArg(OPT_sycl_dump_device_code_EQ)) {
     Arg *A = Args.getLastArg(OPT_sycl_dump_device_code_EQ);
-    SmallString<128> Dir(A->getValue());
-    if (Dir.empty())
-      llvm::sys::path::native(Dir = "./");
+    OffloadImageDumpDir = A->getValue();
+    if (OffloadImageDumpDir.empty())
+      sys::path::native(OffloadImageDumpDir = "./");
     else
-      Dir.append(llvm::sys::path::get_separator());
-
-    SPIRVDumpDir = Dir;
+      OffloadImageDumpDir.append(sys::path::get_separator());
   }
 
   {

clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td

@@ -242,7 +242,7 @@ Flags<[WrapperOnlyOption]>,  HelpText<"Embed LLVM IR for runtime kernel fusion">
 
 def sycl_dump_device_code_EQ : Joined<["--", "-"], "sycl-dump-device-code=">,
   Flags<[WrapperOnlyOption]>,
-  HelpText<"Path to the folder where the tool dumps SPIR-V device code. Other formats aren't dumped.">;
+  HelpText<"Directory to dump offloading images to.">;
 
 // Options to enable/disable device dynamic linking.
 def sycl_allow_device_image_dependencies : Flag<["--", "-"], "sycl-allow-device-image-dependencies">,

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -89,6 +89,19 @@ static const char *LegalSYCLFunctionsInStatelessMode[] = {
 
 namespace {
 
+class BuffDeleter {
+public:
+  BuffDeleter(char *Buffer) : Buff(Buffer) {};
+  ~BuffDeleter() { std::free(Buff); };
+
+  BuffDeleter() = delete;
+  BuffDeleter(const BuffDeleter &) = delete;
+  BuffDeleter(BuffDeleter &&) = delete;
+
+private:
+  char *Buff;
+};
+
 class ESIMDVerifierImpl {
   const Module &M;
   bool MayNeedForceStatelessMemModeAPI;
@@ -149,6 +162,7 @@ class ESIMDVerifierImpl {
             continue;
 
           id::OutputBuffer NameBuf;
+          BuffDeleter NameBufDeleter(NameBuf.getBuffer());
           NameNode->print(NameBuf);
           StringRef Name(NameBuf.getBuffer(), NameBuf.getCurrentPosition());
 

llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp

@@ -118,13 +118,19 @@ struct ThreadSanitizerOnSpirv {
 
   void initialize();
 
-  void instrumentKernelsMetadata();
+  void instrumentModule();
 
   void appendDebugInfoToArgs(Instruction *I, SmallVectorImpl<Value *> &Args);
 
 private:
+  void instrumentGlobalVariables();
+
+  void instrumentKernelsMetadata();
+
   bool isSupportedSPIRKernel(Function &F);
 
+  bool isUnsupportedDeviceGlobal(const GlobalVariable &G);
+
   GlobalVariable *GetOrCreateGlobalString(StringRef Name, StringRef Value,
                                           unsigned AddressSpace);
 
@@ -243,7 +249,7 @@ PreservedAnalyses ModuleThreadSanitizerPass::run(Module &M,
     return PreservedAnalyses::all();
   if (Triple(M.getTargetTriple()).isSPIROrSPIRV()) {
     ThreadSanitizerOnSpirv Spirv(M);
-    Spirv.instrumentKernelsMetadata();
+    Spirv.instrumentModule();
   } else
     insertModuleCtor(M);
   return PreservedAnalyses::none();
@@ -327,6 +333,70 @@ bool ThreadSanitizerOnSpirv::isSupportedSPIRKernel(Function &F) {
   return true;
 }
 
+bool ThreadSanitizerOnSpirv::isUnsupportedDeviceGlobal(
+    const GlobalVariable &G) {
+  if (G.user_empty())
+    return true;
+  // Skip instrumenting on "__TsanKernelMetadata" etc.
+  if (G.getName().starts_with("__Tsan"))
+    return true;
+  if (G.getName().starts_with("__tsan_"))
+    return true;
+  if (G.getName().starts_with("__spirv_BuiltIn"))
+    return true;
+  if (G.getName().starts_with("__usid_str"))
+    return true;
+  // TODO: Will support global variable with local address space later.
+  if (G.getAddressSpace() == kSpirOffloadLocalAS)
+    return true;
+  // Global variables have constant value or constant address space will not
+  // trigger race condition.
+  if (G.isConstant() || G.getAddressSpace() == kSpirOffloadConstantAS)
+    return true;
+  return false;
+}
+
+void ThreadSanitizerOnSpirv::instrumentModule() {
+  instrumentGlobalVariables();
+  instrumentKernelsMetadata();
+}
+
+void ThreadSanitizerOnSpirv::instrumentGlobalVariables() {
+  SmallVector<Constant *, 8> DeviceGlobalMetadata;
+
+  // Device global metadata is described by a structure
+  //  size_t device_global_size
+  //  size_t beginning address of the device global
+  StructType *StructTy = StructType::get(IntptrTy, IntptrTy);
+
+  for (auto &G : M.globals()) {
+    if (isUnsupportedDeviceGlobal(G)) {
+      for (auto *User : G.users())
+        if (auto *Inst = dyn_cast<Instruction>(User))
+          Inst->setNoSanitizeMetadata();
+      continue;
+    }
+
+    DeviceGlobalMetadata.push_back(ConstantStruct::get(
+        StructTy,
+        ConstantInt::get(IntptrTy, DL.getTypeAllocSize(G.getValueType())),
+        ConstantExpr::getPointerCast(&G, IntptrTy)));
+  }
+
+  if (DeviceGlobalMetadata.empty())
+    return;
+
+  // Create meta data global to record device globals' information
+  ArrayType *ArrayTy = ArrayType::get(StructTy, DeviceGlobalMetadata.size());
+  Constant *MetadataInitializer =
+      ConstantArray::get(ArrayTy, DeviceGlobalMetadata);
+  GlobalVariable *MsanDeviceGlobalMetadata = new GlobalVariable(
+      M, MetadataInitializer->getType(), false, GlobalValue::AppendingLinkage,
+      MetadataInitializer, "__TsanDeviceGlobalMetadata", nullptr,
+      GlobalValue::NotThreadLocal, 1);
+  MsanDeviceGlobalMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
+}
+
 void ThreadSanitizerOnSpirv::instrumentKernelsMetadata() {
   SmallVector<Constant *, 8> SpirKernelsMetadata;
 

llvm/test/Instrumentation/ThreadSanitizer/SPIRV/instrument_device_global.ll

@@ -0,0 +1,20 @@
+; RUN: opt < %s -passes='function(tsan),module(tsan-module)' -tsan-instrument-func-entry-exit=0 -tsan-instrument-memintrinsics=0 -S | FileCheck %s
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
+target triple = "spir64-unknown-unknown"
+
+@dev_global = external addrspace(1) global { [4 x i32] }
+@dev_global_no_users = dso_local addrspace(1) global { [4 x i32] } zeroinitializer
+@.str = external addrspace(1) constant [59 x i8]
+@__spirv_BuiltInGlobalInvocationId = external addrspace(1) constant <3 x i64>
+
+; CHECK: @__TsanDeviceGlobalMetadata
+; CHECK-NOT: @dev_global_no_users
+; CHECK-NOT: @.str
+; CHECK-NOT: @__spirv_BuiltInGlobalInvocationId
+; CHECK-SAME: @dev_global
+
+define spir_func void @test() {
+entry:
+  %call = call spir_func ptr addrspace(4) null(ptr addrspace(4) addrspacecast (ptr addrspace(1) @dev_global to ptr addrspace(4)), i64 0)
+  ret void
+}

sycl/doc/design/SYCL2020-SpecializationConstants.md

@@ -80,7 +80,7 @@ as:
 [sycl-2020-spec-constant-glossary]: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#specialization-constant
 [sycl-2020-glossary]: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#glossary
 
-And implementation is based on [SPIR-V speficiation][spirv-spec] support
+And implementation is based on [SPIR-V specification][spirv-spec] support
 for [Specialization][spirv-specialization]. However, the specification also
 states the following:
 
@@ -172,7 +172,7 @@ Based on those limitations, the following mapping design is proposed:
   ```
   namespace detail {
     // assuming user defined the following specialization_id:
-    // constexpr specialiation_id<int> int_const;
+    // constexpr specialization_id<int> int_const;
     // class Wrapper {
     // public:
     //   static constexpr specialization_id<float> float_const;
@@ -341,7 +341,7 @@ used to identify the specialization constants at SPIR-V level.
 As noted above one symbolic ID can have several numeric IDs assigned to it -
 such 1:N mapping comes from the fact that at SPIR-V level, composite
 specialization constants don't have dedicated IDs and they are being identified
-and specialized through their scalar leafs and corresponding numeric IDs.
+and specialized through their scalar leaves and corresponding numeric IDs.
 
 For example, the following code:
 ```
@@ -375,7 +375,7 @@ unique_symbolic_id_for_id_A -> { 1, 2, 3 }
 
 As it is shown in the example above, if a composite specialization constant
 contains another composite within it, that nested composite is also being
-"flattened" and its leafs are considered to be leafs of the parent
+"flattened" and its leaves are considered to be leaves of the parent
 specialization constants. This done by depth-first search through the composite
 elements.
 
@@ -509,8 +509,8 @@ constant in that buffer:
 ```
 [
   0, // for id_int, the first constant is at the beginning of the buffer
-  4, // sizeof(int) == 4, the second constant is located right after the fisrt one
-  16, // sizeof(int) + sizezof(A) == 4, the same approach for the third constant
+  4, // sizeof(int) == 4, the second constant is located right after the first one
+  16, // sizeof(int) + sizeof(A) == 4, the same approach for the third constant
 ]
 ```
 
@@ -661,9 +661,9 @@ While transforming SYCL kernel function into an OpenCL kernel, DPC++ FE should
 - Communicate to DPC++ RT which kernel argument should be used for passing
   a buffer with specialization constant values when they are emulated.
 
-DPC++ FE provides implementation of `__builtin_sycl_unique_id` built-in function and
-it also populates special integration footer with the content required by DPC++
-RT for access to right device image properties describing specialization
+DPC++ FE provides implementation of `__builtin_sycl_unique_id` built-in function
+and it also populates special integration footer with the content required by
+DPC++ RT for access to right device image properties describing specialization
 constants.
 
 #### SYCL Kernel function transformations
@@ -763,7 +763,7 @@ struct A {
 };
 
 constexpr specialization_id<int> id_int;
-struct Wraper {
+struct Wrapper {
 public:
   static constexpr specialization_id<A> id_A;
 };
@@ -839,10 +839,10 @@ constexpr sycl::specialization_id<int> same_name{1};
 
 namespace {
   constexpr sycl::specialization_id<int> same_name{2}:
-  /* application code that referenes ::(unnamed)::same_name */
+  /* application code that references ::(unnamed)::same_name */
   namespace {
     constexpr sycl::specialization_id<int> same_name{3}:
-    /* application code that referenes ::(unnamed)::(unnamed)::same_name */
+    /* application code that references ::(unnamed)::(unnamed)::same_name */
   }
 }
 
@@ -899,7 +899,7 @@ namespace {
 
   namespace __sycl_detail {
     // Sometimes we need a 'shim', which points to another 'shim' in order to
-    // "extract" a variable from an anonymous namespace unambiguosly
+    // "extract" a variable from an anonymous namespace unambiguously
     static constexpr decltype(__sycl_detail::__shim_1()) &__shim_2() {
       // still address of ::(unnamed)::(unnamed)::same_name;
       return __sycl_detail::__shim_1();
@@ -972,7 +972,7 @@ address of the specialization constant provided by user and `offset` field of
 the descriptor as `(char*)(SpecConstantValuesMap[SymbolicID]) + offset`.
 
 That calculation is required, because at SPIR-V level composite
-specialization constants are respresented by several specialization constants
+specialization constants are represented by several specialization constants
 for each element of a composite, whilst on a SYCL level, the whole composite
 is passed by user as a single blob of data. `offset` field from properties is
 used to specify which exact piece of that blob should be extracted to perform
@@ -1053,7 +1053,7 @@ the translator will generate `OpSpecConstant` SPIR-V instructions with proper
               OpDecorate %A.float SpecId 44                        ; ID of the 2nd member
      %A.i32 = OpSpecConstant %int.type 0                           ; 1st member with default value
    %A.float = OpSpecConstant %float.type 0.0                       ; 2nd member with default value
-    %struct = OpSpecConstantComposite %struct.type %A.i32 %A.float ; Composite doens't need IDs or default value
+    %struct = OpSpecConstantComposite %struct.type %A.i32 %A.float ; Composite doesn't need IDs or default value
          %1 = OpTypeFunction %int
 
        %get = OpFunction %int None %1

sycl/include/sycl/detail/spinlock.hpp

@@ -26,6 +26,8 @@ namespace detail {
 /// std::mutex, that doesn't provide such guarantees).
 class SpinLock {
 public:
+  bool try_lock() { return !MLock.test_and_set(std::memory_order_acquire); }
+
   void lock() {
     while (MLock.test_and_set(std::memory_order_acquire))
       std::this_thread::yield();