From dd695f1315450aae32039f5870543afb20257977 Mon Sep 17 00:00:00 2001 From: "Sarnie, Nick" Date: Thu, 29 Sep 2022 10:41:22 -0400 Subject: [PATCH 1/5] [SYCL] Allow specification of double GRF mode for SYCL Signed-off-by: Sarnie, Nick --- .../llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h | 1 - .../llvm/SYCLLowerIR/LowerKernelProps.h | 28 +++++++++ llvm/lib/Passes/PassBuilder.cpp | 1 + llvm/lib/Passes/PassRegistry.def | 2 +- llvm/lib/SYCLLowerIR/CMakeLists.txt | 2 +- llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp | 1 + llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 4 +- ...MDKernelProps.cpp => LowerKernelProps.cpp} | 22 ++++--- .../SYCLLowerIR/ESIMD/lower_kernel_props.ll | 44 -------------- llvm/test/SYCLLowerIR/lower_kernel_props.ll | 44 ++++++++++++++ .../tools/sycl-post-link/sycl-double-grf.ll | 58 +++++++++++++++++++ .../{sycl-esimd => }/sycl-esimd-double-grf.ll | 10 ++-- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 10 ++-- llvm/tools/sycl-post-link/ModuleSplitter.h | 2 +- llvm/tools/sycl-post-link/sycl-post-link.cpp | 14 ++--- sycl/include/sycl/ext/intel/esimd.hpp | 1 - .../{esimd => }/detail/misc_intrin.hpp | 16 +++-- .../{esimd => }/kernel_properties.hpp | 15 +++-- .../program_manager/program_manager.cpp | 12 ++-- 19 files changed, 187 insertions(+), 100 deletions(-) create mode 100644 llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h rename llvm/lib/SYCLLowerIR/{ESIMD/LowerESIMDKernelProps.cpp => LowerKernelProps.cpp} (74%) delete mode 100644 llvm/test/SYCLLowerIR/ESIMD/lower_kernel_props.ll create mode 100644 llvm/test/SYCLLowerIR/lower_kernel_props.ll create mode 100644 llvm/test/tools/sycl-post-link/sycl-double-grf.ll rename llvm/test/tools/sycl-post-link/{sycl-esimd => }/sycl-esimd-double-grf.ll (86%) rename sycl/include/sycl/ext/intel/experimental/{esimd => }/detail/misc_intrin.hpp (59%) rename sycl/include/sycl/ext/intel/experimental/{esimd => }/kernel_properties.hpp (83%) diff --git a/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h b/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h index 4f3be898663ba..fe7fe6708b9da 100644 --- a/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h +++ b/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h @@ -16,7 +16,6 @@ namespace llvm { namespace esimd { -constexpr char ATTR_DOUBLE_GRF[] = "esimd-double-grf"; constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd"; using CallGraphNodeAction = std::function; diff --git a/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h b/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h new file mode 100644 index 0000000000000..68599c75ac195 --- /dev/null +++ b/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h @@ -0,0 +1,28 @@ +//===---- LowerKernelProps.h - lower kernel properties -----------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// Lowers SYCL kernel properties into attributes used by sycl-post-link +//===----------------------------------------------------------------------===// + +#pragma once + +#include "llvm/IR/PassManager.h" + +namespace llvm { + +namespace sycl_kernel_props { +constexpr char ATTR_DOUBLE_GRF[] = "double-grf"; +} + +// Lowers calls to __sycl_set_kernel_properties +class SYCLLowerKernelPropsPass + : public PassInfoMixin { +public: + PreservedAnalyses run(Module &M, ModuleAnalysisManager &); +}; + +} // namespace llvm diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp index a9ac958e54259..08e2303dba8b4 100644 --- a/llvm/lib/Passes/PassBuilder.cpp +++ b/llvm/lib/Passes/PassBuilder.cpp @@ -81,6 +81,7 @@ #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h" #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" +#include "llvm/SYCLLowerIR/LowerKernelProps.h" #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h" #include "llvm/SYCLLowerIR/LowerWGScope.h" #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h" diff --git a/llvm/lib/Passes/PassRegistry.def b/llvm/lib/Passes/PassRegistry.def index f8991b430f526..83aa4a98fa094 100644 --- a/llvm/lib/Passes/PassRegistry.def +++ b/llvm/lib/Passes/PassRegistry.def @@ -127,7 +127,7 @@ MODULE_PASS("memprof-module", ModuleMemProfilerPass()) MODULE_PASS("poison-checking", PoisonCheckingPass()) MODULE_PASS("pseudo-probe-update", PseudoProbeUpdatePass()) MODULE_PASS("LowerESIMD", SYCLLowerESIMDPass()) -MODULE_PASS("lower-esimd-kernel-props", SYCLLowerESIMDKernelPropsPass()) +MODULE_PASS("lower-kernel-props", SYCLLowerKernelPropsPass()) MODULE_PASS("ESIMDLowerVecArg", ESIMDLowerVecArgPass()) MODULE_PASS("esimd-opt-call-conv", ESIMDOptimizeVecArgCallConvPass()) MODULE_PASS("esimd-verifier", ESIMDVerifierPass()) diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index ff65c2e9d25e3..e58303d49aa05 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt @@ -48,7 +48,6 @@ set_property(GLOBAL PROPERTY LLVMGenXIntrinsics_BINARY_PROP ${LLVMGenXIntrinsics add_llvm_component_library(LLVMSYCLLowerIR ESIMD/LowerESIMD.cpp - ESIMD/LowerESIMDKernelProps.cpp ESIMD/LowerESIMDVLoadVStore.cpp ESIMD/LowerESIMDVecArg.cpp ESIMD/ESIMDUtils.cpp @@ -56,6 +55,7 @@ add_llvm_component_library(LLVMSYCLLowerIR ESIMD/LowerESIMDKernelAttrs.cpp ESIMD/ESIMDOptimizeVecArgCallConv.cpp LowerInvokeSimd.cpp + LowerKernelProps.cpp LowerWGScope.cpp LowerWGLocalMemory.cpp MutatePrintfAddrspace.cpp diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index ce24993760090..cc6d643dddec9 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp @@ -48,6 +48,7 @@ static const char *LegalSYCLFunctions[] = { "^sycl::_V1::exp<.+>", "^sycl::_V1::bit_cast<.+>", "^sycl::_V1::operator.+<.+>", + "^sycl::_V1::ext::intel::experimental::set_kernel_properties", "^sycl::_V1::ext::oneapi::sub_group::.+", "^sycl::_V1::ext::oneapi::experimental::spec_constant<.+>::.+", "^sycl::_V1::ext::oneapi::experimental::this_sub_group", diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index ea35c02a7530d..ccbbd7ad79f0a 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1771,8 +1771,8 @@ size_t SYCLLowerESIMDPass::runOnFunction(Function &F, ToErase.push_back(CI); continue; } - assert(!Name.startswith("__esimd_set_kernel_properties") && - "__esimd_set_kernel_properties must have been lowered"); + assert(!Name.startswith("__sycl_set_kernel_properties") && + "__sycl_set_kernel_properties must have been lowered"); if (Name.empty() || !Name.startswith(ESIMD_INTRIN_PREF1)) continue; diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelProps.cpp b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp similarity index 74% rename from llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelProps.cpp rename to llvm/lib/SYCLLowerIR/LowerKernelProps.cpp index 684f7c47ebc06..6bc4e379c8e41 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelProps.cpp +++ b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp @@ -1,16 +1,16 @@ -//===---- LowerESIMDKernelProps.h - lower __esimd_set_kernel_properties ---===// +//===---- LowerESIMDKernelProps.h - lower __sycl_set_kernel_properties ---===// // // 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 // //===----------------------------------------------------------------------===// -// Finds and lowers __esimd_set_kernel_properties calls: converts the call to +// Finds and lowers __sycl_set_kernel_properties calls: converts the call to // function attributes and adds those attributes to all kernels which can // potentially call this intrinsic. +#include "llvm/SYCLLowerIR/LowerKernelProps.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" -#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" #include "llvm/ADT/SmallPtrSet.h" #include "llvm/IR/Instructions.h" @@ -18,17 +18,17 @@ #include "llvm/IR/Operator.h" #include "llvm/Pass.h" -#define DEBUG_TYPE "LowerESIMDKernelProps" +#define DEBUG_TYPE "LowerKernelProps" using namespace llvm; namespace { constexpr char SET_KERNEL_PROPS_FUNC_NAME[] = - "_Z29__esimd_set_kernel_propertiesi"; + "_Z28__sycl_set_kernel_propertiesi"; // Kernel property identifiers. Should match ones in -// sycl/include/sycl/ext/intel/experimental/esimd/kernel_properties.hpp +// sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp enum property_ids { use_double_grf = 0 }; void processSetKernelPropertiesCall(CallInst &CI) { @@ -47,9 +47,7 @@ void processSetKernelPropertiesCall(CallInst &CI) { // TODO: Keep track of traversed functions to avoid repeating traversals // over same function. llvm::esimd::traverseCallgraphUp(F, [](Function *GraphNode) { - if (llvm::esimd::isESIMDKernel(*GraphNode)) { - GraphNode->addFnAttr(llvm::esimd::ATTR_DOUBLE_GRF); - } + GraphNode->addFnAttr(llvm::sycl_kernel_props::ATTR_DOUBLE_GRF); }); break; default: @@ -60,8 +58,8 @@ void processSetKernelPropertiesCall(CallInst &CI) { } // namespace namespace llvm { -PreservedAnalyses -SYCLLowerESIMDKernelPropsPass::run(Module &M, ModuleAnalysisManager &MAM) { +PreservedAnalyses SYCLLowerKernelPropsPass::run(Module &M, + ModuleAnalysisManager &MAM) { Function *F = M.getFunction(SET_KERNEL_PROPS_FUNC_NAME); if (!F) { @@ -71,7 +69,7 @@ SYCLLowerESIMDKernelPropsPass::run(Module &M, ModuleAnalysisManager &MAM) { SmallVector Users(F->users()); for (User *Usr : Users) { - // a call can be the only use of the __esimd_set_kernel_properties built-in + // a call can be the only use of the __sycl_set_kernel_properties built-in CallInst *CI = cast(Usr); processSetKernelPropertiesCall(*CI); CI->eraseFromParent(); diff --git a/llvm/test/SYCLLowerIR/ESIMD/lower_kernel_props.ll b/llvm/test/SYCLLowerIR/ESIMD/lower_kernel_props.ll deleted file mode 100644 index d9b2fe06cde20..0000000000000 --- a/llvm/test/SYCLLowerIR/ESIMD/lower_kernel_props.ll +++ /dev/null @@ -1,44 +0,0 @@ -; This test checks handling of the -; __esimd_set_kernel_properties(...); -; intrinsic by LowerESIMDKernelProps pass - it should: -; - determine kernels calling this intrinsic (walk up the call graph) -; - remove the intrinsic call -; - mark the kernel with corresponding attribute (only "esimd-double-grf" for now) - -; RUN: opt -passes=lower-esimd-kernel-props -S %s -o - | FileCheck %s - -; ModuleID = 'double_grf.bc' -source_filename = "llvm-link" -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" - -define dso_local spir_func void @_Z17double_grf_markerv() { -; CHECK: define dso_local spir_func void @_Z17double_grf_markerv() -; -- '0' constant argument means "double GRF" property: - call spir_func void @_Z29__esimd_set_kernel_propertiesi(i32 noundef 0) -; -- Check that ESIMD lowering removed the marker call above: -; CHECK-NOT: {{.*}} @_Z29__esimd_set_kernel_propertiesi - ret void -; CHECK-NEXT: ret void -} - -declare dso_local spir_func void @_Z29__esimd_set_kernel_propertiesi(i32 noundef) - -; -- This kernel calls the marker function indirectly -define weak_odr dso_local spir_kernel void @__ESIMD_double_grf_kernel1() !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { -; CHECK: {{.*}} spir_kernel void @__ESIMD_double_grf_kernel1() #0 - call spir_func void @_Z17double_grf_markerv() - ret void -} - -; -- This kernel calls the marker function directly -define weak_odr dso_local spir_kernel void @__ESIMD_double_grf_kernel2() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { -; CHECK: {{.*}} spir_kernel void @__ESIMD_double_grf_kernel2() #0 - call spir_func void @_Z29__esimd_set_kernel_propertiesi(i32 noundef 0) - ret void -} - -attributes #0 = { "esimd-double-grf" } - -!0 = !{} -!1 = !{i32 1} diff --git a/llvm/test/SYCLLowerIR/lower_kernel_props.ll b/llvm/test/SYCLLowerIR/lower_kernel_props.ll new file mode 100644 index 0000000000000..aa36cd8c94c06 --- /dev/null +++ b/llvm/test/SYCLLowerIR/lower_kernel_props.ll @@ -0,0 +1,44 @@ +; This test checks handling of the +; __sycl_set_kernel_properties(...); +; intrinsic by LowerKernelProps pass - it should: +; - determine kernels calling this intrinsic (walk up the call graph) +; - remove the intrinsic call +; - mark the kernel with corresponding attribute (only "double-grf" for now) + +; RUN: opt -passes=lower-kernel-props -S %s -o - | FileCheck %s + +; ModuleID = 'double_grf.bc' +source_filename = "llvm-link" +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" + +define dso_local spir_func void @_Z17double_grf_markerv() { +; CHECK: define dso_local spir_func void @_Z17double_grf_markerv() +; -- '0' constant argument means "double GRF" property: + call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0) +; -- Check that LowerKernelProps removed the marker call above: +; CHECK-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi + ret void +; CHECK-NEXT: ret void +} + +declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef) + +; -- This kernel calls the marker function indirectly +define weak_odr dso_local spir_kernel void @__double_grf_kernel1() !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { +; CHECK: {{.*}} spir_kernel void @__double_grf_kernel1() #0 + call spir_func void @_Z17double_grf_markerv() + ret void +} + +; -- This kernel calls the marker function directly +define weak_odr dso_local spir_kernel void @__double_grf_kernel2() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { +; CHECK: {{.*}} spir_kernel void @__double_grf_kernel2() #0 + call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0) + ret void +} + +attributes #0 = { "double-grf" } + +!0 = !{} +!1 = !{i32 1} diff --git a/llvm/test/tools/sycl-post-link/sycl-double-grf.ll b/llvm/test/tools/sycl-post-link/sycl-double-grf.ll new file mode 100644 index 0000000000000..dcc5695ce48e9 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/sycl-double-grf.ll @@ -0,0 +1,58 @@ +; This test checks handling of the +; set_kernel_properties(kernel_properties::use_double_grf); +; by the post-link-tool: +; - ESIMD/SYCL splitting happens as usual +; - ESIMD module is further split into callgraphs for entry points requesting +; "double GRF" and callgraphs for entry points which are not +; - Compiler adds 'isDoubleGRF' property to the device binary +; images requesting "double GRF" + +; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S %s -o %t.table +; RUN: FileCheck %s -input-file=%t.table +; RUN: FileCheck %s -input-file=%t_x2grf_0.ll --check-prefixes CHECK-2xGRF-IR +; RUN: FileCheck %s -input-file=%t_x2grf_0.prop --check-prefixes CHECK-2xGRF-PROP +; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-SYM +; RUN: FileCheck %s -input-file=%t_x2grf_0.sym --check-prefixes CHECK-2xGRF-SYM + +; CHECK: [Code|Properties|Symbols] +; CHECK: {{.*}}_x2grf_0.ll|{{.*}}_x2grf_0.prop|{{.*}}_x2grf_0.sym +; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym + +; CHECK-2xGRF-PROP: isDoubleGRF=1|1 + +; CHECK-SYCL-SYM: __SYCL_kernel +; CHECK-SYCL-SYM-EMPTY: + +; CHECK-2xGRF-SYM: __double_grf_kernel +; CHECK-2xGRF-SYM-EMPTY: + +; ModuleID = 'double_grf.bc' +source_filename = "llvm-link" +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" + +define weak_odr dso_local spir_kernel void @__SYCL_kernel() #0 { +entry: + ret void +} + +define dso_local spir_func void @_Z17double_grf_markerv() { +entry: + call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0) +; -- Check that LowerKernelProps lowering removed the marker call above: +; CHECK-2xGRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi + ret void +} + +declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef) + +define weak_odr dso_local spir_kernel void @__double_grf_kernel() #0 { +entry: + call spir_func void @_Z17double_grf_markerv() + ret void +} + +attributes #0 = { "sycl-module-id"="a.cpp" } + +!0 = !{} +!1 = !{i32 1} diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-double-grf.ll b/llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll similarity index 86% rename from llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-double-grf.ll rename to llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll index 0abc1150b9634..5d6104a40505f 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-double-grf.ll +++ b/llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll @@ -4,7 +4,7 @@ ; - ESIMD/SYCL splitting happens as usual ; - ESIMD module is further split into callgraphs for entry points requesting ; "double GRF" and callgraphs for entry points which are not -; - Compiler adds 'isDoubleGRFEsimdImage' property to the ESIMD device binary +; - Compiler adds 'isDoubleGRF' property to the ESIMD device binary ; images requesting "double GRF" ; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S %s -o %t.table @@ -21,7 +21,7 @@ ; CHECK: {{.*}}esimd_0.ll|{{.*}}esimd_0.prop|{{.*}}esimd_0.sym ; CHECK-ESIMD-2xGRF-PROP: isEsimdImage=1|1 -; CHECK-ESIMD-2xGRF-PROP: isDoubleGRFEsimdImage=1|1 +; CHECK-ESIMD-2xGRF-PROP: isDoubleGRF=1|1 ; CHECK-SYCL-SYM: __SYCL_kernel ; CHECK-SYCL-SYM-EMPTY: @@ -49,13 +49,13 @@ entry: define dso_local spir_func void @_Z17double_grf_markerv() { entry: - call spir_func void @_Z29__esimd_set_kernel_propertiesi(i32 noundef 0) + call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0) ; -- Check that ESIMD lowering removed the marker call above: -; CHECK-ESIMD-2xGRF-IR-NOT: {{.*}} @_Z29__esimd_set_kernel_propertiesi +; CHECK-ESIMD-2xGRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi ret void } -declare dso_local spir_func void @_Z29__esimd_set_kernel_propertiesi(i32 noundef) +declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef) define weak_odr dso_local spir_kernel void @__ESIMD_double_grf_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { entry: diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 5f10a9436a255..a8a3e37888917 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -19,6 +19,7 @@ #include "llvm/IR/LegacyPassManager.h" #include "llvm/IR/Module.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" +#include "llvm/SYCLLowerIR/LowerKernelProps.h" #include "llvm/Transforms/IPO.h" #include "llvm/Transforms/IPO/GlobalDCE.h" #include "llvm/Transforms/IPO/StripDeadPrototypes.h" @@ -41,7 +42,6 @@ constexpr char ESIMD_SCOPE_NAME[] = ""; constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd"; constexpr char ATTR_SYCL_MODULE_ID[] = "sycl-module-id"; -constexpr char ATTR_DOUBLE_GRF[] = "esimd-double-grf"; bool hasIndirectFunctionsOrCalls(const Module &M) { for (const auto &F : M.functions()) { @@ -726,11 +726,11 @@ void EntryPointGroup::rebuildFromNames(const std::vector &Names, } std::unique_ptr -getESIMDDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) { +getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) { EntryPointGroupVec Groups = groupEntryPointsByAttribute( - MD, ATTR_DOUBLE_GRF, EmitOnlyKernelsAsEntryPoints, - [](EntryPointGroup &G) { - if (G.GroupId == ATTR_DOUBLE_GRF) { + MD, llvm::sycl_kernel_props::ATTR_DOUBLE_GRF, + EmitOnlyKernelsAsEntryPoints, [](EntryPointGroup &G) { + if (G.GroupId == llvm::sycl_kernel_props::ATTR_DOUBLE_GRF) { G.Props.UsesDoubleGRF = true; } }); diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index b5847b53b2777..f362c2c1973da 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -251,7 +251,7 @@ getSplitterByMode(ModuleDesc &&MD, IRSplitMode Mode, bool EmitOnlyKernelsAsEntryPoints); std::unique_ptr -getESIMDDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); +getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); #ifndef NDEBUG void dumpEntryPoints(const EntryPointSet &C, const char *msg = "", int Tab = 0); diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index f444e4b17fa69..345e6b98763b4 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -38,6 +38,7 @@ #include "llvm/Passes/PassBuilder.h" #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" +#include "llvm/SYCLLowerIR/LowerKernelProps.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/InitLLVM.h" @@ -436,8 +437,7 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isEsimdImage", true}); } if (MD.isDoubleGRF()) - PropSet[PropSetRegTy::SYCL_MISC_PROP].insert( - {"isDoubleGRFEsimdImage", true}); + PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isDoubleGRF", true}); { std::vector FuncNames = getKernelNamesUsingAssert(M); for (const StringRef &FName : FuncNames) @@ -561,8 +561,8 @@ bool lowerEsimdConstructs(module_split::ModuleDesc &MD) { IrPropSymFilenameTriple saveModule(module_split::ModuleDesc &MD, int I, StringRef IRFilename = "") { IrPropSymFilenameTriple Res; - StringRef Suffix = - MD.isDoubleGRF() ? "_esimd_x2grf" : (MD.isESIMD() ? "_esimd" : ""); + StringRef Suffix = MD.isDoubleGRF() ? MD.isESIMD() ? "_esimd_x2grf" : "_x2grf" + : (MD.isESIMD() ? "_esimd" : ""); if (!IRFilename.empty()) { // don't save IR, just record the filename @@ -725,7 +725,7 @@ processInputModule(std::unique_ptr M) { // Lower kernel properties setting APIs before "double GRF" splitting, as: // - the latter uses the result of the former // - saves processing time - Modified |= runModulePass(*M); + Modified |= runModulePass(*M); DUMP_ENTRY_POINTS(*M, EmitOnlyKernelsAsEntryPoints, "Input"); @@ -762,8 +762,8 @@ processInputModule(std::unique_ptr M) { DUMP_ENTRY_POINTS(MDesc.entries(), MDesc.Name.c_str(), 1); std::unique_ptr DoubleGRFSplitter = - module_split::getESIMDDoubleGRFSplitter(std::move(MDesc), - EmitOnlyKernelsAsEntryPoints); + module_split::getDoubleGRFSplitter(std::move(MDesc), + EmitOnlyKernelsAsEntryPoints); const bool SplitByDoubleGRF = DoubleGRFSplitter->totalSplits() > 1; Modified |= SplitByDoubleGRF; diff --git a/sycl/include/sycl/ext/intel/esimd.hpp b/sycl/include/sycl/ext/intel/esimd.hpp index 36bd7f7662d7f..58758c109178a 100644 --- a/sycl/include/sycl/ext/intel/esimd.hpp +++ b/sycl/include/sycl/ext/intel/esimd.hpp @@ -87,7 +87,6 @@ #include #include #include -#include #include #include diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/detail/misc_intrin.hpp b/sycl/include/sycl/ext/intel/experimental/detail/misc_intrin.hpp similarity index 59% rename from sycl/include/sycl/ext/intel/experimental/esimd/detail/misc_intrin.hpp rename to sycl/include/sycl/ext/intel/experimental/detail/misc_intrin.hpp index 608005c425b86..fbae9267158ed 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/detail/misc_intrin.hpp +++ b/sycl/include/sycl/ext/intel/experimental/detail/misc_intrin.hpp @@ -1,18 +1,24 @@ -//==------------ misc_intrin.hpp - DPC++ Explicit SIMD API -----------------==// +//==------------ misc_intrin.hpp - SYCL Kernel Properties -----------------==// // // 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 // //===----------------------------------------------------------------------===// -// Declares miscellaneous Explicit SIMD intrinsics. +// Declares miscellaneous SYCL intrinsics. //===----------------------------------------------------------------------===// #pragma once -/// @cond ESIMD_DETAIL +/// @cond SYCL_DETAIL -__ESIMD_INTRIN void __esimd_set_kernel_properties(int prop_mask) +#ifdef __SYCL_DEVICE_ONLY__ +#define __SYCL_INTRIN SYCL_EXTERNAL +#else +#define __SYCL_INTRIN inline +#endif // __SYCL_DEVICE_ONLY__ + +__SYCL_INTRIN void __sycl_set_kernel_properties(int prop_mask) #ifdef __SYCL_DEVICE_ONLY__ ; #else @@ -20,4 +26,4 @@ __ESIMD_INTRIN void __esimd_set_kernel_properties(int prop_mask) } // Only "double GRF" property is supported for now, safe to ignore on host. #endif // __SYCL_DEVICE_ONLY__ -/// @endcond ESIMD_DETAIL +/// @endcond SYCL_DETAIL diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/kernel_properties.hpp b/sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp similarity index 83% rename from sycl/include/sycl/ext/intel/experimental/esimd/kernel_properties.hpp rename to sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp index 445a7ec4e70ac..786a396e921a6 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/kernel_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp @@ -1,4 +1,4 @@ -//==---------------- kernel_properties.hpp - DPC++ Explicit SIMD API -------==// +//==---------------- kernel_properties.hpp - SYCL Kernel Properties -------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -10,14 +10,13 @@ #pragma once -#include -#include +#include #include namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { -namespace ext::intel::experimental::esimd { +namespace ext::intel::experimental { namespace kernel_properties { @@ -26,7 +25,7 @@ namespace kernel_properties { /// // Implementation note: ::value fields should match property IDs -// specified in llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +// specified in llvm/lib/SYCLLowerIR/LowerKernelProps.cpp namespace detail { // Proxy to access private property classes' fields from the API code. @@ -65,16 +64,16 @@ void set_kernel_properties(KernelProps... props) { constexpr bool IsDoubleGRF = std::is_same_v; if constexpr (IsDoubleGRF) { - __esimd_set_kernel_properties( + __sycl_set_kernel_properties( kernel_properties::detail::proxy< kernel_properties::use_double_grf_tag>::value); } else { static_assert(IsDoubleGRF && - "set_kernel_properties: invalid ESIMD kernel property"); + "set_kernel_properties: invalid kernel property"); } }); } -} // namespace ext::intel::experimental::esimd +} // namespace ext::intel::experimental } // __SYCL_INLINE_VER_NAMESPACE(_V1) } // namespace sycl diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 7ad1c1b649a63..0d9d9fb58ea9e 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -403,10 +403,7 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, CompileOpts += std::string(TemporaryStr); } bool isEsimdImage = getUint32PropAsBool(Img, "isEsimdImage"); - bool isDoubleGRFEsimdImage = - getUint32PropAsBool(Img, "isDoubleGRFEsimdImage"); - assert((!isDoubleGRFEsimdImage || isEsimdImage) && - "doubleGRF applies only to ESIMD binary images"); + bool isDoubleGRF = getUint32PropAsBool(Img, "isDoubleGRF"); // The -vc-codegen option is always preserved for ESIMD kernels, regardless // of the contents SYCL_PROGRAM_COMPILE_OPTIONS environment variable. if (isEsimdImage) { @@ -418,9 +415,10 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, if (detail::SYCLConfig::get() == 0) CompileOpts += " -disable-finalizer-msg"; } - if (isDoubleGRFEsimdImage) { - assert(!CompileOpts.empty()); // -vc-codegen must be present - CompileOpts += " -doubleGRF"; + if (isDoubleGRF) { + if (!CompileOpts.empty()) + CompileOpts += " "; + CompileOpts += "-ze-opt-large-register-file"; } } From 53414ebfcc036cd5429fa8027a4e7044902319d3 Mon Sep 17 00:00:00 2001 From: "Sarnie, Nick" Date: Fri, 30 Sep 2022 09:59:07 -0400 Subject: [PATCH 2/5] address review feedback Signed-off-by: Sarnie, Nick --- .../include/llvm/SYCLLowerIR/CallgraphUtils.h | 48 +++++++++++++ .../llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h | 31 -------- llvm/lib/SYCLLowerIR/CMakeLists.txt | 1 + llvm/lib/SYCLLowerIR/CallgraphUtils.cpp | 71 +++++++++++++++++++ llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp | 55 -------------- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 5 +- .../ESIMD/LowerESIMDKernelAttrs.cpp | 3 +- llvm/lib/SYCLLowerIR/LowerKernelProps.cpp | 4 +- llvm/tools/sycl-post-link/sycl-post-link.cpp | 11 ++- 9 files changed, 136 insertions(+), 93 deletions(-) create mode 100644 llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h create mode 100644 llvm/lib/SYCLLowerIR/CallgraphUtils.cpp diff --git a/llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h b/llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h new file mode 100644 index 0000000000000..6eefe759475b1 --- /dev/null +++ b/llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h @@ -0,0 +1,48 @@ +//===------------ CallgraphUtils.h - Callgraph utility functions +//------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// Utility functions for traversing callgraphs. +//===----------------------------------------------------------------------===// +#pragma once + +#include "llvm/ADT/SmallPtrSet.h" +#include "llvm/IR/Function.h" + +#include +namespace llvm { +namespace CallgraphUtils { +using CallGraphNodeAction = std::function; + +// Traverses call graph starting from given function up the call chain applying +// given action to each function met on the way. If \c ErrorOnNonCallUse +// parameter is true, then no functions' uses are allowed except calls. +// Otherwise, any function where use of the current one happened is added to the +// call graph as if the use was a call. +// Functions which are part of the visited set ('Visited' parameter) are not +// traversed. +void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction NodeF, + SmallPtrSetImpl &Visited, + bool ErrorOnNonCallUse); + +template +void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF, + SmallPtrSetImpl &Visited, + bool ErrorOnNonCallUse) { + traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited, + ErrorOnNonCallUse); +} + +template +void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF, + bool ErrorOnNonCallUse = true) { + SmallPtrSet Visited; + traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited, + ErrorOnNonCallUse); +} +} // namespace CallgraphUtils +} // namespace llvm diff --git a/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h b/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h index fe7fe6708b9da..d8ab03e37c746 100644 --- a/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h +++ b/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h @@ -11,42 +11,11 @@ #include "llvm/ADT/SmallPtrSet.h" #include "llvm/IR/Function.h" -#include - namespace llvm { namespace esimd { constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd"; -using CallGraphNodeAction = std::function; - -// Traverses call graph starting from given function up the call chain applying -// given action to each function met on the way. If \c ErrorOnNonCallUse -// parameter is true, then no functions' uses are allowed except calls. -// Otherwise, any function where use of the current one happened is added to the -// call graph as if the use was a call. -// Functions which are part of the visited set ('Visited' parameter) are not -// traversed. -void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction NodeF, - SmallPtrSetImpl &Visited, - bool ErrorOnNonCallUse); - -template -void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF, - SmallPtrSetImpl &Visited, - bool ErrorOnNonCallUse) { - traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited, - ErrorOnNonCallUse); -} - -template -void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF, - bool ErrorOnNonCallUse = true) { - SmallPtrSet Visited; - traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited, - ErrorOnNonCallUse); -} - // Tells whether given function is a ESIMD kernel. bool isESIMDKernel(const Function &F); // Tells whether given function is a ESIMD function. diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index e58303d49aa05..6fb59d7d323fd 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt @@ -54,6 +54,7 @@ add_llvm_component_library(LLVMSYCLLowerIR ESIMD/ESIMDVerifier.cpp ESIMD/LowerESIMDKernelAttrs.cpp ESIMD/ESIMDOptimizeVecArgCallConv.cpp + CallgraphUtils.cpp LowerInvokeSimd.cpp LowerKernelProps.cpp LowerWGScope.cpp diff --git a/llvm/lib/SYCLLowerIR/CallgraphUtils.cpp b/llvm/lib/SYCLLowerIR/CallgraphUtils.cpp new file mode 100644 index 0000000000000..d5719c1b2a2f4 --- /dev/null +++ b/llvm/lib/SYCLLowerIR/CallgraphUtils.cpp @@ -0,0 +1,71 @@ +//===------------ CallgraphUtils.cpp - Callgraph utility functions +//------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// Utility functions for traversing callgraphs. +//===----------------------------------------------------------------------===// +#include "llvm/SYCLLowerIR/CallgraphUtils.h" +#include "llvm/IR/Instructions.h" + +namespace llvm { +namespace CallgraphUtils { +void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF, + SmallPtrSetImpl &FunctionsVisited, + bool ErrorOnNonCallUse) { + SmallVector Worklist; + + if (FunctionsVisited.count(F) == 0) + Worklist.push_back(F); + + while (!Worklist.empty()) { + Function *CurF = Worklist.pop_back_val(); + FunctionsVisited.insert(CurF); + // Apply the action function. + ActionF(CurF); + + // Update all callers as well. + for (auto It = CurF->use_begin(); It != CurF->use_end(); It++) { + auto FCall = It->getUser(); + auto ErrMsg = + llvm::Twine(__FILE__ " ") + + "Function use other than call detected while traversing call\n" + "graph up to a kernel"; + if (!isa(FCall)) { + // A use other than a call is met... + if (ErrorOnNonCallUse) { + // ... non-call is an error - report + llvm::report_fatal_error(ErrMsg); + } else { + // ... non-call is OK - add using function to the worklist + if (auto *I = dyn_cast(FCall)) { + auto UseF = I->getFunction(); + + if (FunctionsVisited.count(UseF) == 0) { + Worklist.push_back(UseF); + } + } + } + } else { + auto *CI = cast(FCall); + + if ((CI->getCalledFunction() != CurF)) { + // CurF is used in a call, but not as the callee. + if (ErrorOnNonCallUse) + llvm::report_fatal_error(ErrMsg); + } else { + auto FCaller = CI->getFunction(); + + if (!FunctionsVisited.count(FCaller)) { + Worklist.push_back(FCaller); + } + } + } + } + } +} +} // namespace CallgraphUtils +} // namespace llvm diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp index 0c43697b4396b..1a150f8dca551 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp @@ -19,61 +19,6 @@ namespace llvm { namespace esimd { -void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF, - SmallPtrSetImpl &FunctionsVisited, - bool ErrorOnNonCallUse) { - SmallVector Worklist; - - if (FunctionsVisited.count(F) == 0) - Worklist.push_back(F); - - while (!Worklist.empty()) { - Function *CurF = Worklist.pop_back_val(); - FunctionsVisited.insert(CurF); - // Apply the action function. - ActionF(CurF); - - // Update all callers as well. - for (auto It = CurF->use_begin(); It != CurF->use_end(); It++) { - auto FCall = It->getUser(); - auto ErrMsg = - llvm::Twine(__FILE__ " ") + - "Function use other than call detected while traversing call\n" - "graph up to a kernel"; - if (!isa(FCall)) { - // A use other than a call is met... - if (ErrorOnNonCallUse) { - // ... non-call is an error - report - llvm::report_fatal_error(ErrMsg); - } else { - // ... non-call is OK - add using function to the worklist - if (auto *I = dyn_cast(FCall)) { - auto UseF = I->getFunction(); - - if (FunctionsVisited.count(UseF) == 0) { - Worklist.push_back(UseF); - } - } - } - } else { - auto *CI = cast(FCall); - - if ((CI->getCalledFunction() != CurF)) { - // CurF is used in a call, but not as the callee. - if (ErrorOnNonCallUse) - llvm::report_fatal_error(ErrMsg); - } else { - auto FCaller = CI->getFunction(); - - if (!FunctionsVisited.count(FCaller)) { - Worklist.push_back(FCaller); - } - } - } - } - } -} - bool isESIMD(const Function &F) { return F.getMetadata(ESIMD_MARKER_MD) != nullptr; } diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index ccbbd7ad79f0a..1c440dec184b7 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -14,6 +14,7 @@ //===----------------------------------------------------------------------===// #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" +#include "llvm/SYCLLowerIR/CallgraphUtils.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" #include "llvm/ADT/DenseMap.h" @@ -977,7 +978,7 @@ static void translateSLMInit(CallInst &CI) { *F->getParent(), genx::KernelMDOp::SLMSize, NewVal}; // TODO: Keep track of traversed functions (use 4-argument version of // traverseCallgraphUp) to avoid repeating traversals over same function. - esimd::traverseCallgraphUp(F, SetMaxSLMSize); + CallgraphUtils::traverseCallgraphUp(F, SetMaxSLMSize); } // This function sets/updates VCNamedBarrierCount attribute to the kernels @@ -995,7 +996,7 @@ static void translateNbarrierInit(CallInst &CI) { *F->getParent(), genx::KernelMDOp::NBarrierCnt, NewVal}; // TODO: Keep track of traversed functions to avoid repeating traversals // over same function. - esimd::traverseCallgraphUp(F, SetMaxNBarrierCnt); + CallgraphUtils::traverseCallgraphUp(F, SetMaxNBarrierCnt); } static void translatePackMask(CallInst &CI) { diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp index f74544133e82e..c4b4490959340 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp @@ -8,6 +8,7 @@ // Finds and adds sycl_explicit_simd attributes to wrapper functions that wrap // ESIMD kernel functions +#include "llvm/SYCLLowerIR/CallgraphUtils.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" @@ -26,7 +27,7 @@ SYCLFixupESIMDKernelWrapperMDPass::run(Module &M, ModuleAnalysisManager &MAM) { if (llvm::esimd::isESIMD(F)) { // TODO: Keep track of traversed functions to avoid repeating traversals // over same function. - llvm::esimd::traverseCallgraphUp( + llvm::CallgraphUtils::traverseCallgraphUp( &F, [&](Function *GraphNode) { if (!llvm::esimd::isESIMD(*GraphNode)) { diff --git a/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp index 6bc4e379c8e41..790b79a47125a 100644 --- a/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp +++ b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp @@ -10,7 +10,7 @@ // potentially call this intrinsic. #include "llvm/SYCLLowerIR/LowerKernelProps.h" -#include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" +#include "llvm/SYCLLowerIR/CallgraphUtils.h" #include "llvm/ADT/SmallPtrSet.h" #include "llvm/IR/Instructions.h" @@ -46,7 +46,7 @@ void processSetKernelPropertiesCall(CallInst &CI) { case property_ids::use_double_grf: // TODO: Keep track of traversed functions to avoid repeating traversals // over same function. - llvm::esimd::traverseCallgraphUp(F, [](Function *GraphNode) { + llvm::CallgraphUtils::traverseCallgraphUp(F, [](Function *GraphNode) { GraphNode->addFnAttr(llvm::sycl_kernel_props::ATTR_DOUBLE_GRF); }); break; diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 345e6b98763b4..c09849a933af2 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -552,6 +552,14 @@ bool lowerEsimdConstructs(module_split::ModuleDesc &MD) { return !Res.areAllPreserved(); } +// Compute the filename suffix for the module +StringRef getModuleSuffix(const module_split::ModuleDesc &MD) { + if (MD.isDoubleGRF()) { + return MD.isESIMD() ? "_esimd_x2grf" : "_x2grf"; + } + return MD.isESIMD() ? "_esimd" : ""; +} + // @param MD Module descriptor to save // @param IRFilename filename of already available IR component. If not empty, // IR component saving is skipped, and this file name is recorded as such in @@ -561,8 +569,7 @@ bool lowerEsimdConstructs(module_split::ModuleDesc &MD) { IrPropSymFilenameTriple saveModule(module_split::ModuleDesc &MD, int I, StringRef IRFilename = "") { IrPropSymFilenameTriple Res; - StringRef Suffix = MD.isDoubleGRF() ? MD.isESIMD() ? "_esimd_x2grf" : "_x2grf" - : (MD.isESIMD() ? "_esimd" : ""); + StringRef Suffix = getModuleSuffix(MD); if (!IRFilename.empty()) { // don't save IR, just record the filename From 9497c683acaea66e91c75496aceaa7f4c5e2f123 Mon Sep 17 00:00:00 2001 From: "Sarnie, Nick" Date: Mon, 3 Oct 2022 09:18:37 -0400 Subject: [PATCH 3/5] address feedback 2 Signed-off-by: Sarnie, Nick --- llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h | 6 ++++-- llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h | 9 ++++----- llvm/lib/SYCLLowerIR/CallgraphUtils.cpp | 6 ++++-- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 4 ++-- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp | 2 +- llvm/lib/SYCLLowerIR/LowerKernelProps.cpp | 6 +++--- llvm/tools/sycl-post-link/ModuleSplitter.cpp | 6 +++--- 7 files changed, 21 insertions(+), 18 deletions(-) diff --git a/llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h b/llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h index 6eefe759475b1..60a9fce9a0a8d 100644 --- a/llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h +++ b/llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h @@ -15,7 +15,8 @@ #include namespace llvm { -namespace CallgraphUtils { +namespace sycl { +namespace utils { using CallGraphNodeAction = std::function; // Traverses call graph starting from given function up the call chain applying @@ -44,5 +45,6 @@ void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF, traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited, ErrorOnNonCallUse); } -} // namespace CallgraphUtils +} // namespace utils +} // namespace sycl } // namespace llvm diff --git a/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h b/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h index 68599c75ac195..81e12f9c93d9e 100644 --- a/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h +++ b/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h @@ -12,17 +12,16 @@ #include "llvm/IR/PassManager.h" -namespace llvm { - -namespace sycl_kernel_props { +namespace sycl { +namespace kernel_props { constexpr char ATTR_DOUBLE_GRF[] = "double-grf"; } - +} // namespace sycl +namespace llvm { // Lowers calls to __sycl_set_kernel_properties class SYCLLowerKernelPropsPass : public PassInfoMixin { public: PreservedAnalyses run(Module &M, ModuleAnalysisManager &); }; - } // namespace llvm diff --git a/llvm/lib/SYCLLowerIR/CallgraphUtils.cpp b/llvm/lib/SYCLLowerIR/CallgraphUtils.cpp index d5719c1b2a2f4..63a5fc0062ab0 100644 --- a/llvm/lib/SYCLLowerIR/CallgraphUtils.cpp +++ b/llvm/lib/SYCLLowerIR/CallgraphUtils.cpp @@ -12,7 +12,8 @@ #include "llvm/IR/Instructions.h" namespace llvm { -namespace CallgraphUtils { +namespace sycl { +namespace utils { void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF, SmallPtrSetImpl &FunctionsVisited, bool ErrorOnNonCallUse) { @@ -67,5 +68,6 @@ void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF, } } } -} // namespace CallgraphUtils +} // namespace utils +} // namespace sycl } // namespace llvm diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index 1c440dec184b7..0c705a94893aa 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -978,7 +978,7 @@ static void translateSLMInit(CallInst &CI) { *F->getParent(), genx::KernelMDOp::SLMSize, NewVal}; // TODO: Keep track of traversed functions (use 4-argument version of // traverseCallgraphUp) to avoid repeating traversals over same function. - CallgraphUtils::traverseCallgraphUp(F, SetMaxSLMSize); + sycl::utils::traverseCallgraphUp(F, SetMaxSLMSize); } // This function sets/updates VCNamedBarrierCount attribute to the kernels @@ -996,7 +996,7 @@ static void translateNbarrierInit(CallInst &CI) { *F->getParent(), genx::KernelMDOp::NBarrierCnt, NewVal}; // TODO: Keep track of traversed functions to avoid repeating traversals // over same function. - CallgraphUtils::traverseCallgraphUp(F, SetMaxNBarrierCnt); + sycl::utils::traverseCallgraphUp(F, SetMaxNBarrierCnt); } static void translatePackMask(CallInst &CI) { diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp index c4b4490959340..b59a26df7419e 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp @@ -27,7 +27,7 @@ SYCLFixupESIMDKernelWrapperMDPass::run(Module &M, ModuleAnalysisManager &MAM) { if (llvm::esimd::isESIMD(F)) { // TODO: Keep track of traversed functions to avoid repeating traversals // over same function. - llvm::CallgraphUtils::traverseCallgraphUp( + sycl::utils::traverseCallgraphUp( &F, [&](Function *GraphNode) { if (!llvm::esimd::isESIMD(*GraphNode)) { diff --git a/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp index 790b79a47125a..dd42af256556a 100644 --- a/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp +++ b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp @@ -1,4 +1,4 @@ -//===---- LowerESIMDKernelProps.h - lower __sycl_set_kernel_properties ---===// +//===---- LowerKernelProps.h - lower __sycl_set_kernel_properties ---===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -46,8 +46,8 @@ void processSetKernelPropertiesCall(CallInst &CI) { case property_ids::use_double_grf: // TODO: Keep track of traversed functions to avoid repeating traversals // over same function. - llvm::CallgraphUtils::traverseCallgraphUp(F, [](Function *GraphNode) { - GraphNode->addFnAttr(llvm::sycl_kernel_props::ATTR_DOUBLE_GRF); + llvm::sycl::utils::traverseCallgraphUp(F, [](Function *GraphNode) { + GraphNode->addFnAttr(::sycl::kernel_props::ATTR_DOUBLE_GRF); }); break; default: diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index a8a3e37888917..e74b91c0750c8 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -728,9 +728,9 @@ void EntryPointGroup::rebuildFromNames(const std::vector &Names, std::unique_ptr getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) { EntryPointGroupVec Groups = groupEntryPointsByAttribute( - MD, llvm::sycl_kernel_props::ATTR_DOUBLE_GRF, - EmitOnlyKernelsAsEntryPoints, [](EntryPointGroup &G) { - if (G.GroupId == llvm::sycl_kernel_props::ATTR_DOUBLE_GRF) { + MD, sycl::kernel_props::ATTR_DOUBLE_GRF, EmitOnlyKernelsAsEntryPoints, + [](EntryPointGroup &G) { + if (G.GroupId == sycl::kernel_props::ATTR_DOUBLE_GRF) { G.Props.UsesDoubleGRF = true; } }); From a5d466f70309d706f9ed30f6157c13a64d12dd59 Mon Sep 17 00:00:00 2001 From: "Sarnie, Nick" Date: Mon, 3 Oct 2022 16:05:52 -0400 Subject: [PATCH 4/5] address feedback 3 Signed-off-by: Sarnie, Nick --- .../llvm/SYCLLowerIR/{CallgraphUtils.h => SYCLUtils.h} | 4 ++-- llvm/lib/SYCLLowerIR/CMakeLists.txt | 2 +- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 2 +- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp | 2 +- llvm/lib/SYCLLowerIR/LowerKernelProps.cpp | 2 +- llvm/lib/SYCLLowerIR/{CallgraphUtils.cpp => SYCLUtils.cpp} | 6 +++--- 6 files changed, 9 insertions(+), 9 deletions(-) rename llvm/include/llvm/SYCLLowerIR/{CallgraphUtils.h => SYCLUtils.h} (94%) rename llvm/lib/SYCLLowerIR/{CallgraphUtils.cpp => SYCLUtils.cpp} (93%) diff --git a/llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h similarity index 94% rename from llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h rename to llvm/include/llvm/SYCLLowerIR/SYCLUtils.h index 60a9fce9a0a8d..65b089a6233a2 100644 --- a/llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h +++ b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h @@ -1,4 +1,4 @@ -//===------------ CallgraphUtils.h - Callgraph utility functions +//===------------ SYCLUtils.h - SYCL utility functions //------------------===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. @@ -6,7 +6,7 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// Utility functions for traversing callgraphs. +// Utility functions for SYCL. //===----------------------------------------------------------------------===// #pragma once diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index 6fb59d7d323fd..b3b14afe77931 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt @@ -54,13 +54,13 @@ add_llvm_component_library(LLVMSYCLLowerIR ESIMD/ESIMDVerifier.cpp ESIMD/LowerESIMDKernelAttrs.cpp ESIMD/ESIMDOptimizeVecArgCallConv.cpp - CallgraphUtils.cpp LowerInvokeSimd.cpp LowerKernelProps.cpp LowerWGScope.cpp LowerWGLocalMemory.cpp MutatePrintfAddrspace.cpp SYCLPropagateAspectsUsage.cpp + SYCLUtils.cpp LocalAccessorToSharedMemory.cpp GlobalOffset.cpp diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index 0c705a94893aa..bcaa462a3dde5 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -14,8 +14,8 @@ //===----------------------------------------------------------------------===// #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" -#include "llvm/SYCLLowerIR/CallgraphUtils.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" +#include "llvm/SYCLLowerIR/SYCLUtils.h" #include "llvm/ADT/DenseMap.h" #include "llvm/ADT/DenseSet.h" diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp index b59a26df7419e..ccf7da4ed8a53 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp @@ -8,9 +8,9 @@ // Finds and adds sycl_explicit_simd attributes to wrapper functions that wrap // ESIMD kernel functions -#include "llvm/SYCLLowerIR/CallgraphUtils.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" +#include "llvm/SYCLLowerIR/SYCLUtils.h" #include "llvm/IR/Module.h" #include "llvm/Pass.h" diff --git a/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp index dd42af256556a..9702ff522332c 100644 --- a/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp +++ b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp @@ -10,7 +10,7 @@ // potentially call this intrinsic. #include "llvm/SYCLLowerIR/LowerKernelProps.h" -#include "llvm/SYCLLowerIR/CallgraphUtils.h" +#include "llvm/SYCLLowerIR/SYCLUtils.h" #include "llvm/ADT/SmallPtrSet.h" #include "llvm/IR/Instructions.h" diff --git a/llvm/lib/SYCLLowerIR/CallgraphUtils.cpp b/llvm/lib/SYCLLowerIR/SYCLUtils.cpp similarity index 93% rename from llvm/lib/SYCLLowerIR/CallgraphUtils.cpp rename to llvm/lib/SYCLLowerIR/SYCLUtils.cpp index 63a5fc0062ab0..5eaf6a9a02131 100644 --- a/llvm/lib/SYCLLowerIR/CallgraphUtils.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLUtils.cpp @@ -1,4 +1,4 @@ -//===------------ CallgraphUtils.cpp - Callgraph utility functions +//===------------ SYCLUtils.cpp - SYCL utility functions //------------------===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. @@ -6,9 +6,9 @@ // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -// Utility functions for traversing callgraphs. +// Utility functions for SYCL. //===----------------------------------------------------------------------===// -#include "llvm/SYCLLowerIR/CallgraphUtils.h" +#include "llvm/SYCLLowerIR/SYCLUtils.h" #include "llvm/IR/Instructions.h" namespace llvm { From 95377cd09a0f8696cf55bbcb16d2340b4aaa27a8 Mon Sep 17 00:00:00 2001 From: "Sarnie, Nick" Date: Tue, 4 Oct 2022 14:33:42 -0400 Subject: [PATCH 5/5] address feedback 4 Signed-off-by: Sarnie, Nick --- llvm/lib/SYCLLowerIR/LowerKernelProps.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp index 9702ff522332c..2eb227b29e34c 100644 --- a/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp +++ b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp @@ -1,4 +1,4 @@ -//===---- LowerKernelProps.h - lower __sycl_set_kernel_properties ---===// +//===---- LowerKernelProps.cpp - lower __sycl_set_kernel_properties ---===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information.