Merge branch 'sycl' into pietro/vecz

clang/lib/CodeGen/BackendUtil.cpp

@@ -70,8 +70,6 @@
 #include "llvm/TargetParser/Triple.h"
 #include "llvm/Transforms/HipStdPar/HipStdPar.h"
 #include "llvm/Transforms/IPO/DeadArgumentElimination.h"
-#include "llvm/Transforms/HipStdPar/HipStdPar.h"
-#include "llvm/Transforms/IPO/DeadArgumentElimination.h"
 #include "llvm/Transforms/IPO/EmbedBitcodePass.h"
 #include "llvm/Transforms/IPO/LowerTypeTests.h"
 #include "llvm/Transforms/IPO/ThinLTOBitcodeWriter.h"
@@ -117,21 +115,22 @@ static cl::opt<bool> ClSanitizeOnOptimizerEarlyEP(
     "sanitizer-early-opt-ep", cl::Optional,
     cl::desc("Insert sanitizers on OptimizerEarlyEP."), cl::init(false));
 
+extern cl::opt<InstrProfCorrelator::ProfCorrelatorKind> ProfileCorrelate;
+
+// Re-link builtin bitcodes after optimization
+cl::opt<bool> ClRelinkBuiltinBitcodePostop(
+    "relink-builtin-bitcode-postop", cl::Optional,
+    cl::desc("Re-link builtin bitcodes after optimization."), cl::init(false));
+
 static cl::opt<bool> SYCLNativeCPUBackend(
     "sycl-native-cpu-backend", cl::init(false),
     cl::desc("Run the backend passes for SYCL Native CPU"));
+} // namespace llvm
 
 static cl::opt<bool> SYCLNativeCPUNoVecz(
     "sycl-native-cpu-no-vecz", cl::init(false),
     cl::desc("Disable vectorizer for SYCL Native CPU"));
 
-// Re-link builtin bitcodes after optimization
-cl::opt<bool> ClRelinkBuiltinBitcodePostop(
-    "relink-builtin-bitcode-postop", cl::Optional,
-    cl::desc("Re-link builtin bitcodes after optimization."), cl::init(false));
-
-}
-
 namespace {
 
 // Default filename used for profile generation.

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5608,8 +5608,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     if (IsSYCLOffloadDevice && IsSYCLNativeCPU) {
       CmdArgs.push_back("-mllvm");
       CmdArgs.push_back("-sycl-native-cpu-backend");
-      CmdArgs.push_back("-mllvm");
-      CmdArgs.push_back("-inline-threshold=10000");
     }
 
     // Also ignore explicit -force_cpusubtype_ALL option.

llvm/include/llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h

@@ -10,8 +10,8 @@
 //
 //===----------------------------------------------------------------------===//
 #pragma once
-#include "llvm/IR/PassManager.h"
 #include "llvm/ADT/Twine.h"
+#include "llvm/IR/PassManager.h"
 
 namespace llvm {
 namespace sycl {
@@ -40,7 +40,7 @@ constexpr char SYCLNATIVECPUSUFFIX[] = ".SYCLNCPU";
 constexpr char SYCLNATIVECPUKERNEL[] = ".NativeCPUKernel";
 constexpr char SYCLNATIVECPUPREFIX[] = "__dpcpp_nativecpu";
 inline llvm::Twine addSYCLNativeCPUSuffix(StringRef S) {
-  if (S.startswith(SYCLNATIVECPUPREFIX) || S.endswith(SYCLNATIVECPUKERNEL))
+  if (S.starts_with(SYCLNATIVECPUPREFIX) || S.ends_with(SYCLNATIVECPUKERNEL))
     return S;
   return llvm::Twine(S, SYCLNATIVECPUSUFFIX);
 }

llvm/lib/SYCLLowerIR/ConvertToMuxBuiltinsSYCLNativeCPU.cpp

@@ -83,7 +83,7 @@ static constexpr char MuxBarrier[] = "__mux_work_group_barrier";
 
 Function *getReplaceFunc(Module &M, StringRef Name) {
   LLVMContext &Ctx = M.getContext();
-  static auto *MuxFTy =
+  auto *MuxFTy =
       FunctionType::get(Type::getInt64Ty(Ctx), {Type::getInt32Ty(Ctx)}, false);
   auto F = M.getOrInsertFunction(Name, MuxFTy);
   return cast<Function>(F.getCallee());
@@ -93,11 +93,11 @@ Function *getMuxBarrierFunc(Module &M) {
   // void __mux_work_group_barrier(i32 %id, i32 %scope, i32 %semantics)
   LLVMContext &Ctx = M.getContext();
   auto *Int32Ty = Type::getInt32Ty(Ctx);
-  static auto *MuxFTy = FunctionType::get(Type::getVoidTy(Ctx),
-                                          {Int32Ty, Int32Ty, Int32Ty}, false);
+  auto *MuxFTy = FunctionType::get(Type::getVoidTy(Ctx),
+                                   {Int32Ty, Int32Ty, Int32Ty}, false);
   auto FCallee = M.getOrInsertFunction(MuxBarrier, MuxFTy);
   auto *F = dyn_cast<Function>(FCallee.getCallee());
-  if(!F) {
+  if (!F) {
     report_fatal_error("Error while inserting mux builtins");
   }
   return F;

llvm/lib/SYCLLowerIR/PipelineSYCLNativeCPU.cpp

@@ -24,18 +24,22 @@
 #include "compiler/utils/work_item_loops_pass.h"
 #include "vecz/pass.h"
 #include "vecz/vecz_target_info.h"
+#include "compiler/utils/prepare_barriers_pass.h"
+#include "compiler/utils/sub_group_analysis.h"
+#include "compiler/utils/work_item_loops_pass.h"
 #include "llvm/Transforms/IPO/AlwaysInliner.h"
 #endif
 
 using namespace llvm;
 using namespace sycl::utils;
+
 cl::opt<bool> ForceNoTail("native-cpu-force-no-tail", cl::init(false),
     cl::desc("Never emit the peeling loop for vectorized kernels,"
     "even when the local size is not known to be a multiple of the vector width"));
 
 cl::opt<bool> IsDebug("native-cpu-debug", cl::init(false),
     cl::desc("Emit extra alloca instructions to preserve the value of live"
-    "vriables between barriers"));
+    "variables between barriers"));
 cl::opt<unsigned> NativeCPUVeczWidth("ncpu-vecz-width", cl::init(8), cl::desc("Vector width for SYCL Native CPU vectorizer, defaults to 8"));
 void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(llvm::ModulePassManager &MPM,
                                    ModuleAnalysisManager &MAM, unsigned OptLevel, bool DisableVecz) {
@@ -65,6 +69,7 @@ void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(llvm::ModulePassManager &M
   Opts.ForceNoTail = ForceNoTail;
   MAM.registerPass([&] { return compiler::utils::BuiltinInfoAnalysis(); });
   MAM.registerPass([&] { return compiler::utils::SubgroupAnalysis(); });
+  MPM.addPass(compiler::utils::PrepareBarriersPass());
   MPM.addPass(compiler::utils::WorkItemLoopsPass(Opts));
   MPM.addPass(AlwaysInlinerPass());
 #endif

llvm/lib/SYCLLowerIR/PrepareSYCLNativeCPU.cpp

@@ -12,7 +12,6 @@
 //===----------------------------------------------------------------------===//
 
 #include "llvm/SYCLLowerIR/PrepareSYCLNativeCPU.h"
-#include "llvm/BinaryFormat/MsgPack.h"
 #include "llvm/IR/BasicBlock.h"
 #include "llvm/IR/Constant.h"
 #include "llvm/IR/DebugInfoMetadata.h"
@@ -30,22 +29,13 @@
 #include "llvm/IR/Instruction.h"
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/LLVMContext.h"
-#include "llvm/IR/Operator.h"
 #include "llvm/IR/Value.h"
-#include "llvm/InitializePasses.h"
 #include "llvm/Pass.h"
 #include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h"
 #include "llvm/Support/Casting.h"
-#include "llvm/Support/CodeGen.h"
-#include "llvm/Support/CommandLine.h"
 #include "llvm/Support/ErrorHandling.h"
-#include "llvm/Support/FileSystem.h"
-#include "llvm/Support/raw_ostream.h"
 #include "llvm/Transforms/Utils/Cloning.h"
 #include "llvm/Transforms/Utils/ValueMapper.h"
-#include <functional>
-#include <numeric>
-#include <set>
 #include <utility>
 #include <vector>
 
@@ -61,9 +51,7 @@ using namespace sycl::utils;
 
 namespace {
 
-void fixCallingConv(Function *F) {
-  F->setCallingConv(llvm::CallingConv::C);
-}
+void fixCallingConv(Function *F) { F->setCallingConv(llvm::CallingConv::C); }
 
 void emitSubkernelForKernel(Function *F, Type *NativeCPUArgDescType,
                             Type *StatePtrType, llvm::Constant *StateArgTLS) {
@@ -225,8 +213,7 @@ static Function *addSetLocalIdFunc(Module &M, StringRef Name, Type *StateType) {
   Type *DimTy = I32Ty;
   Type *ValTy = I64Ty;
   Type *PtrTy = PointerType::get(Ctx, NativeCPUGlobalAS);
-  static FunctionType *FTy =
-      FunctionType::get(RetTy, {DimTy, ValTy, PtrTy}, false);
+  FunctionType *FTy = FunctionType::get(RetTy, {DimTy, ValTy, PtrTy}, false);
   auto FCallee = M.getOrInsertFunction(Name, FTy);
   auto *F = cast<Function>(FCallee.getCallee());
   IRBuilder<> Builder(Ctx);
@@ -308,6 +295,7 @@ static Function *addReplaceFunc(Module &M, StringRef Name, Type *StateType) {
     Builder.CreateRetVoid();
     Res = F;
   }
+  Res->setLinkage(GlobalValue::LinkageTypes::InternalLinkage);
   return Res;
 }
 
@@ -459,10 +447,11 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M,
       // CallInstructions in it have debug info, otherwise we end up with
       // invalid IR after inlining.
       if (I->getFunction()->hasMetadata("dbg")) {
-        I->setDebugLoc(DILocation::get(M.getContext(), 0, 0,
-                                       I->getFunction()->getSubprogram()));
-        if (I->getMetadata("dbg"))
-          NewI->setDebugLoc(I->getDebugLoc());
+        if (!I->getMetadata("dbg")) {
+          I->setDebugLoc(DILocation::get(M.getContext(), 0, 0,
+                                         I->getFunction()->getSubprogram()));
+        }
+        NewI->setDebugLoc(I->getDebugLoc());
       }
       ToRemove.push_back(std::make_pair(I, NewI));
     }
@@ -495,7 +484,7 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M,
   // function will not be executed (since it has been inlined) and so we can
   // just define __mux_work_group_barrier as a no-op to avoid linker errors.
   // Todo: currently we can't remove the function here even if it has no uses,
-  // because we may still emit a declaration for in the offload-wrapper.
+  // because we may still emit a declaration for it in the offload-wrapper.
   auto BarrierF =
       M.getFunction(compiler::utils::MuxBuiltins::work_group_barrier);
   if (BarrierF && BarrierF->isDeclaration()) {

sycl/doc/design/SYCLNativeCPU.md

@@ -43,7 +43,7 @@ SYCL Native CPU uses the [oneAPI Construction Kit](https://github.com/codeplayso
 ```
 python3 buildbot/configure.py \
   --native_cpu \
-  --cmake-opt=-DNATIVE_CPU_USE_OCK=Off
+  --cmake-opt=-DNATIVECPU_USE_OCK=Off
 ```
 
 The SYCL Native CPU device needs to be selected at runtime by setting the environment variable `ONEAPI_DEVICE_SELECTOR=native_cpu:cpu`. 

sycl/plugins/native_cpu/CMakeLists.txt

@@ -44,7 +44,7 @@ if(NATIVECPU_USE_OCK)
     ${oneapi-ck_SOURCE_DIR}/modules/cargo/include
     ${oneapi-ck_SOURCE_DIR}/modules/compiler/vecz/include
     ${oneapi-ck_SOURCE_DIR}/modules/compiler/utils/include)
-  target_link_libraries(LLVMSYCLLowerIR PRIVATE compiler-utils vecz)
+  target_link_libraries(LLVMSYCLLowerIR PRIVATE compiler-pipeline vecz)
   target_compile_definitions(pi_native_cpu PRIVATE NATIVECPU_USE_OCK)
 
 endif()

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -56,14 +56,14 @@ endif()
 if(SYCL_PI_UR_USE_FETCH_CONTENT)
   include(FetchContent)
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/PietroGhg/unified-runtime.git")
-  # commit 659d3f469faa99a886fa680a3d6d20449b109578
-  # Merge: 192e9404 f94550b4
+  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime")
+  # commit 73d85ef9f48ec4d1f213066a31ed7e4402b5499b
+  # Merge: e46dc359 7985d3ee
   # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-  # Date:   Tue Nov 14 16:45:24 2023 +0000
-  #     Merge pull request #1059 from martygrant/martin/moveNativeCPUAdapterToUR
-  #     [NATIVECPU] Move Native CPU adapter to UR.
-  set(UNIFIED_RUNTIME_TAG pietro/barriers)
+  # Date:   Mon Jan 29 14:26:08 2024 +0000
+  #     Merge pull request #1289 from nrspruit/fix_multiDevice
+  #     [L0] Fix native kernel usage, multi device kernel pointer and WorkSize
+  set(UNIFIED_RUNTIME_TAG 73d85ef9f48ec4d1f213066a31ed7e4402b5499b)
 
   if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
     set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}")
@@ -207,7 +207,7 @@ endif()
 if ("native_cpu" IN_LIST SYCL_ENABLE_PLUGINS)
     add_dependencies(sycl-runtime-libraries ur_adapter_native_cpu)
 
-  option(NATIVECPU_USE_OCK "Use the oneAPI Construction Kit for Native CPU" On)
+  option(NATIVECPU_USE_OCK "Use the oneAPI Construction Kit for Native CPU" ON)
 
   if(NATIVECPU_USE_OCK)
     message(STATUS "Compiling Native CPU adapter with OCK support.")