[SYCL][NATIVECPU] Support specialization constants on Native CPU (#14446

) Adds support to specialization constants by scheduling the `SpecConstantsPass` in the Native CPU pass pipeline. A small change is needed in that pass since, due to ABI, the return type of the spec const builtins can't be used to determine the type of the spec constant, so for Native CPU we use the default value instead. UR PR: oneapi-src/unified-runtime#1821
intel · Sep 16, 2024 · f011892 · f011892
1 parent 7fc3ac2
commit f011892
Show file tree

Hide file tree

Showing 6 changed files with 15 additions and 8 deletions.
diff --git a/llvm/lib/SYCLLowerIR/SpecConstants.cpp b/llvm/lib/SYCLLowerIR/SpecConstants.cpp
@@ -951,8 +951,9 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
             updatePaddingInLastMDNode(Ctx, SCMetadata, Padding);
           }
 
+          auto *DefValTy = DefaultValue->getType();
           SCMetadata[SymID] = generateSpecConstantMetadata(
-              M, SymID, SCTy, NextID, /* is native spec constant */ false);
+              M, SymID, DefValTy, NextID, /* is native spec constant */ false);
 
           ++NextID.ID;
           NextOffset += Size;

diff --git a/llvm/lib/SYCLNativeCPUUtils/CMakeLists.txt b/llvm/lib/SYCLNativeCPUUtils/CMakeLists.txt
@@ -13,6 +13,7 @@ add_llvm_component_library(LLVMSYCLNativeCPUUtils
   Core
   Support
   Passes
+  SYCLLowerIR
   Target
   TargetParser
   TransformUtils

diff --git a/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp b/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp
@@ -14,6 +14,7 @@
 #include "llvm/SYCLLowerIR/ConvertToMuxBuiltinsSYCLNativeCPU.h"
 #include "llvm/SYCLLowerIR/PrepareSYCLNativeCPU.h"
 #include "llvm/SYCLLowerIR/RenameKernelSYCLNativeCPU.h"
+#include "llvm/SYCLLowerIR/SpecConstants.h"
 #include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h"
 #include "llvm/Support/CommandLine.h"
 
@@ -60,6 +61,7 @@ static cl::opt<bool>
 void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(
     llvm::ModulePassManager &MPM, ModuleAnalysisManager &MAM,
     OptimizationLevel OptLevel) {
+  MPM.addPass(SpecConstantsPass(SpecConstantsPass::HandlingMode::emulation));
   MPM.addPass(ConvertToMuxBuiltinsSYCLNativeCPUPass());
 #ifdef NATIVECPU_USE_OCK
   MPM.addPass(compiler::utils::TransferKernelMetadataPass());

diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake
@@ -117,13 +117,13 @@ if(SYCL_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit f31160dea6d142014f441bc4ca5e58e48827490e
-  # Merge: 2bbe9526 64068799
-  # Author: Piotr Balcer <piotr.balcer@intel.com>
-  # Date:   Thu Sep 12 14:19:48 2024 +0200
-  #     Merge pull request #2083 from kswiecicki/xpti-init-fix
-  #     Fix XPTI initialization bug
-  set(UNIFIED_RUNTIME_TAG f31160dea6d142014f441bc4ca5e58e48827490e)
+  # commit fa9ebe7bd3d9bd11dd5ea8a59eff12f5746411d3
+  # Merge: 92638b2a 9eb1c74f
+  # Author: Omar Ahmed <omar.ahmed@codeplay.com>
+  # Date:   Fri Sep 13 14:44:27 2024 +0100
+  #     Merge pull request #1821 from PietroGhg/pietro/native_cpu_specconstants
+  #     [NATIVECPU] Initial support for spec constants on Native CPU
+  set(UNIFIED_RUNTIME_TAG fa9ebe7bd3d9bd11dd5ea8a59eff12f5746411d3)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
   # Due to the use of dependentloadflag and no installer for UMF and hwloc we need

diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp
@@ -329,6 +329,8 @@ bool is_compatible(const std::vector<kernel_id> &KernelIDs, const device &Dev) {
       return BE == sycl::backend::ext_oneapi_cuda;
     } else if (strcmp(Target, __SYCL_DEVICE_BINARY_TARGET_AMDGCN) == 0) {
       return BE == sycl::backend::ext_oneapi_hip;
+    } else if (strcmp(Target, __SYCL_DEVICE_BINARY_TARGET_NATIVE_CPU) == 0) {
+      return BE == sycl::backend::ext_oneapi_native_cpu;
     }
 
     return false;

diff --git a/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp b/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp
@@ -11,6 +11,7 @@
 // RUN: %{run} %t.out
 //
 // UNSUPPORTED: hip
+// UNSUPPORTED: native_cpu
 
 #include <cstdlib>
 #include <iostream>