Merge branch 'sycl' into counter-based-4

Author: Alberto Cabrera Pérez

clang/lib/Frontend/InitPreprocessor.cpp

@@ -1513,7 +1513,8 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
     const llvm::Triple::SubArchType DeviceSubArch = DeviceTriple.getSubArch();
     if (DeviceTriple.isNVPTX() || DeviceTriple.isAMDGPU() ||
         (DeviceTriple.isSPIR() &&
-         DeviceSubArch != llvm::Triple::SPIRSubArch_fpga))
+         DeviceSubArch != llvm::Triple::SPIRSubArch_fpga) ||
+        LangOpts.SYCLIsNativeCPU)
       Builder.defineMacro("SYCL_USE_NATIVE_FP_ATOMICS");
     // Enable generation of USM address spaces for FPGA.
     if (DeviceSubArch == llvm::Triple::SPIRSubArch_fpga) {

clang/test/Preprocessor/sycl-macro-target-specific.cpp

@@ -42,6 +42,8 @@
 // RUN: | FileCheck --check-prefix=CHECK-SYCL-FP-ATOMICS %s
 // RUN: %clang_cc1 %s -fsycl-is-device -triple amdgcn-amdhsa-amdhsa -E -dM \
 // RUN: | FileCheck --check-prefix=CHECK-SYCL-FP-ATOMICS %s
+// RUN: %clang_cc1 %s -fsycl-is-device -triple x86_64-unknown-linux-gnu -fsycl-is-native-cpu  \
+// RUN: -E -dM | FileCheck --check-prefix=CHECK-SYCL-FP-ATOMICS %s
 // CHECK-SYCL-FP-ATOMICS: #define SYCL_USE_NATIVE_FP_ATOMICS
 // CHECK-SYCL-FP-ATOMICS-NEG-NOT: #define SYCL_USE_NATIVE_FP_ATOMICS
 

llvm/include/llvm/SYCLLowerIR/FAtomicsNativeCPU.h

@@ -0,0 +1,28 @@
+//===------- FAtomicsNativeCPU.h - Materializes FP Atomics ----------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// A transformation pass that materializes floating points atomics by emitting
+// corresponding atomicrmw instruction.
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "llvm/IR/Module.h"
+#include "llvm/IR/PassManager.h"
+
+namespace llvm {
+
+class ModulePass;
+
+class FAtomicsNativeCPU : public PassInfoMixin<FAtomicsNativeCPU> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM);
+};
+
+} // namespace llvm

llvm/lib/SYCLNativeCPUUtils/CMakeLists.txt

@@ -4,6 +4,7 @@ add_llvm_component_library(LLVMSYCLNativeCPUUtils
   RenameKernelSYCLNativeCPU.cpp
   ConvertToMuxBuiltinsSYCLNativeCPU.cpp
   FixABIMuxBuiltinsSYCLNativeCPU.cpp
+  FAtomicsNativeCPU.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLLowerIR

llvm/lib/SYCLNativeCPUUtils/FAtomicsNativeCPU.cpp

@@ -0,0 +1,50 @@
+//===------- FAtomicsNativeCPU.cpp - Materializes FP Atomics --------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// A transformation pass that materializes floating points atomics by emitting
+// corresponding atomicrmw instruction.
+//
+//===----------------------------------------------------------------------===//
+
+#include "llvm/SYCLLowerIR/FAtomicsNativeCPU.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/LLVMContext.h"
+#include "llvm/Support/Alignment.h"
+#include "llvm/Support/AtomicOrdering.h"
+
+using namespace llvm;
+
+PreservedAnalyses FAtomicsNativeCPU::run(Module &M,
+                                         ModuleAnalysisManager &MAM) {
+  bool ModuleChanged = false;
+  auto &Ctx = M.getContext();
+  // TODO: add checks for windows mangling
+  for (auto &F : M) {
+    AtomicRMWInst::BinOp OpCode;
+    if (F.getName().starts_with("_Z21__spirv_AtomicFAddEXT")) {
+      OpCode = AtomicRMWInst::BinOp::FAdd;
+    } else if (F.getName().starts_with("_Z21__spirv_AtomicFMinEXT")) {
+      OpCode = AtomicRMWInst::BinOp::FMin;
+    } else if (F.getName().starts_with("_Z21__spirv_AtomicFMaxEXT")) {
+      OpCode = AtomicRMWInst::BinOp::FMax;
+    } else {
+      continue;
+    }
+
+    BasicBlock *BB = BasicBlock::Create(Ctx, "entry", &F);
+    IRBuilder<> Builder(BB);
+    // Currently we drop arguments 1 and 2 (scope and memory ordering),
+    // defaulting to Monotonic ordering and System scope.
+    auto A =
+        Builder.CreateAtomicRMW(OpCode, F.getArg(0), F.getArg(3), MaybeAlign(),
+                                AtomicOrdering::Monotonic, SyncScope::System);
+    Builder.CreateRet(A);
+  }
+  return ModuleChanged ? PreservedAnalyses::none() : PreservedAnalyses::all();
+}

llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp

@@ -12,6 +12,7 @@
 //
 //===----------------------------------------------------------------------===//
 #include "llvm/SYCLLowerIR/ConvertToMuxBuiltinsSYCLNativeCPU.h"
+#include "llvm/SYCLLowerIR/FAtomicsNativeCPU.h"
 #include "llvm/SYCLLowerIR/FixABIMuxBuiltinsSYCLNativeCPU.h"
 #include "llvm/SYCLLowerIR/PrepareSYCLNativeCPU.h"
 #include "llvm/SYCLLowerIR/RenameKernelSYCLNativeCPU.h"
@@ -70,6 +71,7 @@ void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(
     OptimizationLevel OptLevel) {
   MPM.addPass(SpecConstantsPass(SpecConstantsPass::HandlingMode::emulation));
   MPM.addPass(ConvertToMuxBuiltinsSYCLNativeCPUPass());
+  MPM.addPass(FAtomicsNativeCPU());
 #ifdef NATIVECPU_USE_OCK
   MPM.addPass(compiler::utils::PrepareBarriersPass());
   MPM.addPass(compiler::utils::TransferKernelMetadataPass());

llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp

@@ -1450,8 +1450,14 @@ PreservedAnalyses AddressSanitizerPass::run(Module &M,
   const StackSafetyGlobalInfo *const SSGI =
       ClUseStackSafety ? &MAM.getResult<StackSafetyGlobalAnalysis>(M) : nullptr;
 
-  if (Triple(M.getTargetTriple()).isSPIROrSPIRV())
+  if (Triple(M.getTargetTriple()).isSPIROrSPIRV()) {
     ExtendSpirKernelArgs(M, FAM);
+    // FIXME: W/A skip instrumentation if this module has ESIMD
+    for (auto &F : M) {
+      if (F.hasMetadata("sycl_explicit_simd"))
+        return PreservedAnalyses::all();
+    }
+  }
 
   for (Function &F : M) {
     AddressSanitizer FunctionSanitizer(
@@ -3497,10 +3503,6 @@ bool AddressSanitizer::instrumentFunction(Function &F,
     // function isn't supported yet in intel-graphics-compiler.
     if (F.hasFnAttribute("referenced-indirectly"))
       return false;
-    // FIXME: ESIMD kernel doesn't support noinline functions, so we can't
-    // support sanitizer for it
-    if (F.hasMetadata("sycl_explicit_simd"))
-      return false;
   }
 
   bool FunctionModified = false;

llvm/test/Instrumentation/AddressSanitizer/SPIRV/sycl_esimd.ll

@@ -3,12 +3,16 @@
 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"
 target triple = "spir64-unknown-unknown"
 
+;;
+;; W/A: We skip asan completely if one module has esimd
+;;
+
 define spir_kernel void @sycl_kernel(ptr addrspace(1) %p) #0 {
 ; CHECK-LABEL: define spir_kernel void @sycl_kernel(ptr addrspace(1) %p, ptr addrspace(1) %__asan_launch) #0
 entry:
   %0 = load i32, ptr addrspace(1) %p, align 4
-  ; CHECK: store ptr addrspace(1) %__asan_launch, ptr addrspace(3) @__AsanLaunchInfo, align 8
-  ; CHECK: call void @__asan_load4
+  ; CHECK-NOT: store ptr addrspace(1) %__asan_launch, ptr addrspace(3) @__AsanLaunchInfo, align 8
+  ; CHECK-NOT: call void @__asan_load4
   ret void
 }
 

sycl/source/detail/kernel_bundle_impl.hpp

@@ -12,6 +12,7 @@
 #include <detail/kernel_compiler/kernel_compiler_opencl.hpp>
 #include <detail/kernel_compiler/kernel_compiler_sycl.hpp>
 #include <detail/kernel_impl.hpp>
+#include <detail/persistent_device_code_cache.hpp>
 #include <detail/program_manager/program_manager.hpp>
 #include <sycl/backend_types.hpp>
 #include <sycl/context.hpp>
@@ -396,6 +397,53 @@ class kernel_bundle_impl {
     return SS.str();
   }
 
+  bool
+  extKernelCompilerFetchFromCache(const std::vector<device> Devices,
+                                  const std::vector<std::string> &BuildOptions,
+                                  const std::string &SourceStr,
+                                  ur_program_handle_t &UrProgram) {
+    using ContextImplPtr = std::shared_ptr<sycl::detail::context_impl>;
+    ContextImplPtr ContextImpl = getSyclObjImpl(MContext);
+    const AdapterPtr &Adapter = ContextImpl->getAdapter();
+
+    std::string UserArgs = syclex::detail::userArgsAsString(BuildOptions);
+
+    std::vector<ur_device_handle_t> DeviceHandles;
+    std::transform(
+        Devices.begin(), Devices.end(), std::back_inserter(DeviceHandles),
+        [](const device &Dev) { return getSyclObjImpl(Dev)->getHandleRef(); });
+
+    std::vector<const uint8_t *> Binaries;
+    std::vector<size_t> Lengths;
+    std::vector<std::vector<std::vector<char>>> PersistentBinaries;
+    for (size_t i = 0; i < Devices.size(); i++) {
+      std::vector<std::vector<char>> BinProg =
+          PersistentDeviceCodeCache::getCompiledKernelFromDisc(
+              Devices[i], UserArgs, SourceStr);
+
+      // exit if any device binary is missing
+      if (BinProg.empty()) {
+        return false;
+      }
+      PersistentBinaries.push_back(BinProg);
+
+      Binaries.push_back((uint8_t *)(PersistentBinaries[i][0].data()));
+      Lengths.push_back(PersistentBinaries[i][0].size());
+    }
+
+    ur_program_properties_t Properties = {};
+    Properties.stype = UR_STRUCTURE_TYPE_PROGRAM_PROPERTIES;
+    Properties.pNext = nullptr;
+    Properties.count = 0;
+    Properties.pMetadatas = nullptr;
+
+    Adapter->call<UrApiKind::urProgramCreateWithBinary>(
+        ContextImpl->getHandleRef(), DeviceHandles.size(), DeviceHandles.data(),
+        Lengths.data(), Binaries.data(), &Properties, &UrProgram);
+
+    return true;
+  }
+
   std::shared_ptr<kernel_bundle_impl>
   build_from_source(const std::vector<device> Devices,
                     const std::vector<std::string> &BuildOptions,
@@ -415,57 +463,68 @@ class kernel_bundle_impl {
       DeviceVec.push_back(Dev);
     }
 
-    const auto spirv = [&]() -> std::vector<uint8_t> {
-      if (Language == syclex::source_language::opencl) {
-        // if successful, the log is empty. if failed, throws an error with the
-        // compilation log.
-        const auto &SourceStr = std::get<std::string>(this->Source);
-        std::vector<uint32_t> IPVersionVec(Devices.size());
-        std::transform(DeviceVec.begin(), DeviceVec.end(), IPVersionVec.begin(),
-                       [&](ur_device_handle_t d) {
-                         uint32_t ipVersion = 0;
-                         Adapter->call<UrApiKind::urDeviceGetInfo>(
-                             d, UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t),
-                             &ipVersion, nullptr);
-                         return ipVersion;
-                       });
-        return syclex::detail::OpenCLC_to_SPIRV(SourceStr, IPVersionVec,
-                                                BuildOptions, LogPtr);
-      }
-      if (Language == syclex::source_language::spirv) {
-        const auto &SourceBytes =
-            std::get<std::vector<std::byte>>(this->Source);
-        std::vector<uint8_t> Result(SourceBytes.size());
-        std::transform(SourceBytes.cbegin(), SourceBytes.cend(), Result.begin(),
-                       [](std::byte B) { return static_cast<uint8_t>(B); });
-        return Result;
-      }
-      if (Language == syclex::source_language::sycl) {
-        const auto &SourceStr = std::get<std::string>(this->Source);
-        return syclex::detail::SYCL_to_SPIRV(SourceStr, IncludePairs,
-                                             BuildOptions, LogPtr,
-                                             RegisteredKernelNames);
-      }
-      if (Language == syclex::source_language::sycl_jit) {
-        const auto &SourceStr = std::get<std::string>(this->Source);
-        return syclex::detail::SYCL_JIT_to_SPIRV(SourceStr, IncludePairs,
-                                                 BuildOptions, LogPtr,
-                                                 RegisteredKernelNames);
-      }
-      throw sycl::exception(
-          make_error_code(errc::invalid),
-          "OpenCL C and SPIR-V are the only supported languages at this time");
-    }();
-
     ur_program_handle_t UrProgram = nullptr;
-    Adapter->call<UrApiKind::urProgramCreateWithIL>(ContextImpl->getHandleRef(),
-                                                    spirv.data(), spirv.size(),
-                                                    nullptr, &UrProgram);
-    // program created by urProgramCreateWithIL is implicitly retained.
-    if (UrProgram == nullptr)
-      throw sycl::exception(
-          sycl::make_error_code(errc::invalid),
-          "urProgramCreateWithIL resulted in a null program handle.");
+    // SourceStrPtr will be null when source is Spir-V bytes.
+    const std::string *SourceStrPtr = std::get_if<std::string>(&this->Source);
+    bool FetchedFromCache = false;
+    if (PersistentDeviceCodeCache::isEnabled() && SourceStrPtr) {
+      FetchedFromCache = extKernelCompilerFetchFromCache(
+          Devices, BuildOptions, *SourceStrPtr, UrProgram);
+    }
+
+    if (!FetchedFromCache) {
+      const auto spirv = [&]() -> std::vector<uint8_t> {
+        if (Language == syclex::source_language::opencl) {
+          // if successful, the log is empty. if failed, throws an error with
+          // the compilation log.
+          std::vector<uint32_t> IPVersionVec(Devices.size());
+          std::transform(DeviceVec.begin(), DeviceVec.end(),
+                         IPVersionVec.begin(), [&](ur_device_handle_t d) {
+                           uint32_t ipVersion = 0;
+                           Adapter->call<UrApiKind::urDeviceGetInfo>(
+                               d, UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t),
+                               &ipVersion, nullptr);
+                           return ipVersion;
+                         });
+          return syclex::detail::OpenCLC_to_SPIRV(*SourceStrPtr, IPVersionVec,
+                                                  BuildOptions, LogPtr);
+        }
+        if (Language == syclex::source_language::spirv) {
+          const auto &SourceBytes =
+              std::get<std::vector<std::byte>>(this->Source);
+          std::vector<uint8_t> Result(SourceBytes.size());
+          std::transform(SourceBytes.cbegin(), SourceBytes.cend(),
+                         Result.begin(),
+                         [](std::byte B) { return static_cast<uint8_t>(B); });
+          return Result;
+        }
+        if (Language == syclex::source_language::sycl) {
+          return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, IncludePairs,
+                                               BuildOptions, LogPtr,
+                                               RegisteredKernelNames);
+        }
+        if (Language == syclex::source_language::sycl_jit) {
+          const auto &SourceStr = std::get<std::string>(this->Source);
+          return syclex::detail::SYCL_JIT_to_SPIRV(SourceStr, IncludePairs,
+                                                   BuildOptions, LogPtr,
+                                                   RegisteredKernelNames);
+        }
+        throw sycl::exception(
+            make_error_code(errc::invalid),
+            "SYCL C++, OpenCL C and SPIR-V are the only supported "
+            "languages at this time");
+      }();
+
+      Adapter->call<UrApiKind::urProgramCreateWithIL>(
+          ContextImpl->getHandleRef(), spirv.data(), spirv.size(), nullptr,
+          &UrProgram);
+      // program created by urProgramCreateWithIL is implicitly retained.
+      if (UrProgram == nullptr)
+        throw sycl::exception(
+            sycl::make_error_code(errc::invalid),
+            "urProgramCreateWithIL resulted in a null program handle.");
+
+    } // if(!FetchedFromCache)
 
     std::string XsFlags = extractXsFlags(BuildOptions);
     auto Res = Adapter->call_nocheck<UrApiKind::urProgramBuildExp>(
@@ -501,6 +560,17 @@ class kernel_bundle_impl {
         nullptr, MContext, MDevices, bundle_state::executable, KernelIDs,
         UrProgram);
     device_image_plain DevImg{DevImgImpl};
+
+    // If caching enabled and kernel not fetched from cache, cache.
+    if (PersistentDeviceCodeCache::isEnabled() && !FetchedFromCache &&
+        SourceStrPtr) {
+      for (const auto &Device : Devices) {
+        PersistentDeviceCodeCache::putCompiledKernelToDisc(
+            Device, syclex::detail::userArgsAsString(BuildOptions),
+            *SourceStrPtr, UrProgram);
+      }
+    }
+
     return std::make_shared<kernel_bundle_impl>(MContext, MDevices, DevImg,
                                                 KernelNames, Language);
   }

sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp

@@ -36,6 +36,15 @@ SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs,
   throw sycl::exception(sycl::errc::build,
                         "kernel_compiler does not support GCC<8");
 }
+
+std::string userArgsAsString(const std::vector<std::string> &UserArguments) {
+  return std::accumulate(UserArguments.begin(), UserArguments.end(),
+                         std::string(""),
+                         [](const std::string &A, const std::string &B) {
+                           return A.empty() ? B : A + " " + B;
+                         });
+}
+
 } // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1