From ee7c4ba2b49358d196860049ea929b2d4d70d1cc Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Mon, 3 Jun 2024 12:48:48 +0100 Subject: [PATCH 1/6] Initial support for spec constants on Native CPU --- llvm/lib/SYCLLowerIR/SpecConstants.cpp | 13 +++++++++++-- .../SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp | 2 ++ sycl/cmake/modules/FetchUnifiedRuntime.cmake | 4 ++-- sycl/source/kernel_bundle.cpp | 2 ++ 4 files changed, 17 insertions(+), 4 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/SpecConstants.cpp b/llvm/lib/SYCLLowerIR/SpecConstants.cpp index 4f43a22e95fd9..02ea4c7ca9234 100644 --- a/llvm/lib/SYCLLowerIR/SpecConstants.cpp +++ b/llvm/lib/SYCLLowerIR/SpecConstants.cpp @@ -9,7 +9,9 @@ //===----------------------------------------------------------------------===// #include "llvm/SYCLLowerIR/SpecConstants.h" +#include "llvm/IR/DerivedTypes.h" #include "llvm/SYCLLowerIR/Support.h" +#include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h" #include "llvm/ADT/APInt.h" #include "llvm/ADT/StringMap.h" @@ -937,6 +939,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, unsigned Size = M.getDataLayout().getTypeStoreSize(SCTy); uint64_t Align = M.getDataLayout().getABITypeAlign(SCTy).value(); + // Ensure correct alignment if (CurrentOffset % Align != 0) { // Compute necessary padding to correctly align the constant. @@ -951,8 +954,14 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, updatePaddingInLastMDNode(Ctx, SCMetadata, Padding); } - SCMetadata[SymID] = generateSpecConstantMetadata( - M, SymID, SCTy, NextID, /* is native spec constant */ false); + if (sycl::utils::isSYCLNativeCPU(M) && isa(DefaultValue->getType())) { + auto STy = cast(DefaultValue->getType()); + SCMetadata[SymID] = generateSpecConstantMetadata( + M, SymID, STy, NextID, /* is native spec constant */ false); + } else { + SCMetadata[SymID] = generateSpecConstantMetadata( + M, SymID, SCTy, NextID, /* is native spec constant */ false); + } ++NextID.ID; NextOffset += Size; diff --git a/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp b/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp index 1454c10fc4200..eee17720c283a 100644 --- a/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp +++ b/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp @@ -16,6 +16,7 @@ #include "llvm/SYCLLowerIR/RenameKernelSYCLNativeCPU.h" #include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h" #include "llvm/Support/CommandLine.h" +#include "llvm/SYCLLowerIR/SpecConstants.h" #ifdef NATIVECPU_USE_OCK #include "compiler/utils/builtin_info.h" @@ -60,6 +61,7 @@ static cl::opt 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 index 4993e5791ebf0..9ec486047ec67 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -156,8 +156,8 @@ if(SYCL_UR_USE_FETCH_CONTENT) ) fetch_adapter_source(native_cpu - ${UNIFIED_RUNTIME_REPO} - ${UNIFIED_RUNTIME_TAG} + "https://github.com/PietroGhg/unified-runtime.git" + pietro/native_cpu_specconstants ) if(SYCL_UR_OVERRIDE_FETCH_CONTENT_REPO) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index c666f437b30dd..4f6d27b83bd57 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -329,6 +329,8 @@ bool is_compatible(const std::vector &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_PI_DEVICE_BINARY_TARGET_NATIVE_CPU) == 0) { + return BE == sycl::backend::ext_oneapi_native_cpu; } return false; From 6a18136a3d3544f71f7cdef8d03dc9ebc46c0a33 Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Wed, 26 Jun 2024 14:24:24 +0100 Subject: [PATCH 2/6] Mark kernel-bundle-api.cpp test unsupported --- llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp | 1 - sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp | 1 + 2 files changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp b/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp index b3888db8a7b50..c8b6ef67d5aba 100644 --- a/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp +++ b/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp @@ -473,6 +473,5 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M, ModuleChanged = true; } } - return ModuleChanged ? PreservedAnalyses::none() : PreservedAnalyses::all(); } diff --git a/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp b/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp index 0345fa801c51d..560fd22ce73f8 100644 --- 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 #include From dd32fd0e23785f512210ab3602db060ef3bb2a2f Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Thu, 4 Jul 2024 16:03:39 +0100 Subject: [PATCH 3/6] Link to SYCLLowerIR --- llvm/lib/SYCLNativeCPUUtils/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/llvm/lib/SYCLNativeCPUUtils/CMakeLists.txt b/llvm/lib/SYCLNativeCPUUtils/CMakeLists.txt index 31211a38105a6..bea5f1fac7cb1 100644 --- 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 From 21c278cdd96882449905c7a5faec7fbc99ef67f3 Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Fri, 5 Jul 2024 10:05:18 +0100 Subject: [PATCH 4/6] Formatting --- llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp b/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp index eee17720c283a..e44afd21526d8 100644 --- a/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp +++ b/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp @@ -17,6 +17,7 @@ #include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h" #include "llvm/Support/CommandLine.h" #include "llvm/SYCLLowerIR/SpecConstants.h" +#include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h" #ifdef NATIVECPU_USE_OCK #include "compiler/utils/builtin_info.h" From fe5edcad93aa2db63927745f01c25af4a3345aa8 Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Tue, 9 Jul 2024 14:51:33 +0100 Subject: [PATCH 5/6] Use DefaultValue to compute md --- llvm/lib/SYCLLowerIR/SpecConstants.cpp | 14 +++----------- 1 file changed, 3 insertions(+), 11 deletions(-) diff --git a/llvm/lib/SYCLLowerIR/SpecConstants.cpp b/llvm/lib/SYCLLowerIR/SpecConstants.cpp index 02ea4c7ca9234..83bb489486911 100644 --- a/llvm/lib/SYCLLowerIR/SpecConstants.cpp +++ b/llvm/lib/SYCLLowerIR/SpecConstants.cpp @@ -9,9 +9,7 @@ //===----------------------------------------------------------------------===// #include "llvm/SYCLLowerIR/SpecConstants.h" -#include "llvm/IR/DerivedTypes.h" #include "llvm/SYCLLowerIR/Support.h" -#include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h" #include "llvm/ADT/APInt.h" #include "llvm/ADT/StringMap.h" @@ -939,7 +937,6 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, unsigned Size = M.getDataLayout().getTypeStoreSize(SCTy); uint64_t Align = M.getDataLayout().getABITypeAlign(SCTy).value(); - // Ensure correct alignment if (CurrentOffset % Align != 0) { // Compute necessary padding to correctly align the constant. @@ -954,14 +951,9 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, updatePaddingInLastMDNode(Ctx, SCMetadata, Padding); } - if (sycl::utils::isSYCLNativeCPU(M) && isa(DefaultValue->getType())) { - auto STy = cast(DefaultValue->getType()); - SCMetadata[SymID] = generateSpecConstantMetadata( - M, SymID, STy, NextID, /* is native spec constant */ false); - } else { - SCMetadata[SymID] = generateSpecConstantMetadata( - M, SymID, SCTy, NextID, /* is native spec constant */ false); - } + auto *DefValTy = DefaultValue->getType(); + SCMetadata[SymID] = generateSpecConstantMetadata( + M, SymID, DefValTy, NextID, /* is native spec constant */ false); ++NextID.ID; NextOffset += Size; From d01810c8b308073b5b34a98bfc43a5702747f4c9 Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Wed, 14 Aug 2024 09:55:45 +0100 Subject: [PATCH 6/6] Update __SYCL_DEVICE_BINARY_TARGET_NATIVE_CPU --- .../PipelineSYCLNativeCPU.cpp | 3 +-- .../PrepareSYCLNativeCPU.cpp | 1 + sycl/cmake/modules/FetchUnifiedRuntime.cmake | 18 +++++++++--------- sycl/source/kernel_bundle.cpp | 2 +- 4 files changed, 12 insertions(+), 12 deletions(-) diff --git a/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp b/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp index e44afd21526d8..c78e0d9223ef9 100644 --- a/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp +++ b/llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp @@ -14,10 +14,9 @@ #include "llvm/SYCLLowerIR/ConvertToMuxBuiltinsSYCLNativeCPU.h" #include "llvm/SYCLLowerIR/PrepareSYCLNativeCPU.h" #include "llvm/SYCLLowerIR/RenameKernelSYCLNativeCPU.h" -#include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h" -#include "llvm/Support/CommandLine.h" #include "llvm/SYCLLowerIR/SpecConstants.h" #include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h" +#include "llvm/Support/CommandLine.h" #ifdef NATIVECPU_USE_OCK #include "compiler/utils/builtin_info.h" diff --git a/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp b/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp index c8b6ef67d5aba..b3888db8a7b50 100644 --- a/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp +++ b/llvm/lib/SYCLNativeCPUUtils/PrepareSYCLNativeCPU.cpp @@ -473,5 +473,6 @@ PreservedAnalyses PrepareSYCLNativeCPUPass::run(Module &M, ModuleChanged = true; } } + return ModuleChanged ? PreservedAnalyses::none() : PreservedAnalyses::all(); } diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 9ec486047ec67..5fa953a230d42 100644 --- 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 - # 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 + # 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 @@ -156,8 +156,8 @@ if(SYCL_UR_USE_FETCH_CONTENT) ) fetch_adapter_source(native_cpu - "https://github.com/PietroGhg/unified-runtime.git" - pietro/native_cpu_specconstants + ${UNIFIED_RUNTIME_REPO} + ${UNIFIED_RUNTIME_TAG} ) if(SYCL_UR_OVERRIDE_FETCH_CONTENT_REPO) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 4f6d27b83bd57..12ca87b6604f5 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -329,7 +329,7 @@ bool is_compatible(const std::vector &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_PI_DEVICE_BINARY_TARGET_NATIVE_CPU) == 0) { + } else if (strcmp(Target, __SYCL_DEVICE_BINARY_TARGET_NATIVE_CPU) == 0) { return BE == sycl::backend::ext_oneapi_native_cpu; }