From a2f51f5207812ad924df888be3c0c07c97740b0c Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Thu, 20 Jan 2022 13:23:16 +0300 Subject: [PATCH 01/17] [sycl-post-link] Add the DeviceGlobalsMet property to GlobalBinImageProps --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index c375da0121ea5..d7083d2076583 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -203,7 +203,7 @@ struct GlobalBinImageProps { bool EmitProgramMetadata; bool EmitExportedSymbols; bool IsEsimdKernel; - bool EmitDeviceGlobalPropSet; + bool DeviceGlobalsMet; }; void error(const Twine &Msg) { @@ -631,7 +631,7 @@ void saveModuleProperties(Module &M, const EntryPointGroup &ModuleEntryPoints, PropSet[PropSetRegTy::SYCL_ASSERT_USED].insert({FName, true}); } - if (ImgPSInfo.EmitDeviceGlobalPropSet) { + if (ImgPSInfo.DeviceGlobalsMet) { // Extract device global maps per module auto DevGlobalPropertyMap = collectDeviceGlobalProperties(M); if (!DevGlobalPropertyMap.empty()) From fbeb3941f897766b1802e4d586b50752136122af Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Mon, 7 Feb 2022 17:05:32 +0300 Subject: [PATCH 02/17] Revert "[sycl-post-link] Add the DeviceGlobalsMet property to GlobalBinImageProps" This reverts commit a2f51f5207812ad924df888be3c0c07c97740b0c. --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index d7083d2076583..c375da0121ea5 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -203,7 +203,7 @@ struct GlobalBinImageProps { bool EmitProgramMetadata; bool EmitExportedSymbols; bool IsEsimdKernel; - bool DeviceGlobalsMet; + bool EmitDeviceGlobalPropSet; }; void error(const Twine &Msg) { @@ -631,7 +631,7 @@ void saveModuleProperties(Module &M, const EntryPointGroup &ModuleEntryPoints, PropSet[PropSetRegTy::SYCL_ASSERT_USED].insert({FName, true}); } - if (ImgPSInfo.DeviceGlobalsMet) { + if (ImgPSInfo.EmitDeviceGlobalPropSet) { // Extract device global maps per module auto DevGlobalPropertyMap = collectDeviceGlobalProperties(M); if (!DevGlobalPropertyMap.empty()) From bcf895015c7172e8700ef01f12bba3fc628b2137 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Tue, 8 Feb 2022 14:18:40 +0300 Subject: [PATCH 03/17] [sycl-post-link] Extract the hasDeviceImageScopeProperty function --- llvm/tools/sycl-post-link/DeviceGlobals.cpp | 16 ++++++++++++++-- llvm/tools/sycl-post-link/DeviceGlobals.h | 8 ++++++++ 2 files changed, 22 insertions(+), 2 deletions(-) diff --git a/llvm/tools/sycl-post-link/DeviceGlobals.cpp b/llvm/tools/sycl-post-link/DeviceGlobals.cpp index 95613000112ac..016509c18119d 100644 --- a/llvm/tools/sycl-post-link/DeviceGlobals.cpp +++ b/llvm/tools/sycl-post-link/DeviceGlobals.cpp @@ -59,6 +59,19 @@ bool isDeviceGlobalVariable(const GlobalVariable &GV) { return GV.hasAttribute(SYCL_DEVICE_GLOBAL_SIZE_ATTR); } +/// Return \c true if the variable @GV has the "device_image_scope" property. +/// +/// The function checks whether the variable has the LLVM IR attribute \c +/// device_image_scope and the attribute is not set to \c "false" +/// +/// @param GV [in] A variable to test. +/// +/// @return \c true if the variable has the "device_image_scope" property, +/// \c false otherwise. +bool hasDeviceImageScopeProperty(const GlobalVariable &GV) { + return hasProperty(GV, SYCL_DEVICE_IMAGE_SCOPE_ATTR); +} + /// Returns the unique id for the device global variable. /// /// The function gets this value from the LLVM IR attribute \c @@ -88,8 +101,7 @@ DeviceGlobalPropertyMapTy collectDeviceGlobalProperties(const Module &M) { continue; DGM[getGlobalVariableUniqueId(GV)] = { - {{getUnderlyingTypeSize(GV), - hasProperty(GV, SYCL_DEVICE_IMAGE_SCOPE_ATTR)}}}; + {{getUnderlyingTypeSize(GV), hasDeviceImageScopeProperty(GV)}}}; } return DGM; diff --git a/llvm/tools/sycl-post-link/DeviceGlobals.h b/llvm/tools/sycl-post-link/DeviceGlobals.h index c8c6d253b33b4..0a75c0c61752e 100644 --- a/llvm/tools/sycl-post-link/DeviceGlobals.h +++ b/llvm/tools/sycl-post-link/DeviceGlobals.h @@ -59,6 +59,14 @@ DeviceGlobalPropertyMapTy collectDeviceGlobalProperties(const Module &M); /// otherwise. bool isDeviceGlobalVariable(const GlobalVariable &GV); +/// Return \c true if the variable @GV has the "device_image_scope" property. +/// +/// @param GV [in] A variable to test. +/// +/// @return \c true if the variable has the "device_image_scope" property, +/// \c false otherwise. +bool hasDeviceImageScopeProperty(const GlobalVariable &GV); + /// Returns the unique id for the device global variable. /// /// @param GV [in] Device Global variable. From 66ab81c0971991ea5587aec92ae5af5b6e2c7f40 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Tue, 8 Feb 2022 19:25:42 +0300 Subject: [PATCH 04/17] [sycl-post-link] Add a check for device globals with device_image_scope For device global variables with the 'device_image_scope' property, the check that there are no usages of a single device global variable from kernels grouped to different modules is required as it is described in the design document [1]. [1] https://github.com/intel/llvm/blob/sycl/sycl/doc/DeviceGlobal.md#changes-to-the-sycl-post-link-tool --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 93 ++++++++++++++++++-- 1 file changed, 88 insertions(+), 5 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index c375da0121ea5..1c7031371799a 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -365,6 +365,80 @@ void groupEntryPoints(const Module &M, EntryPointGroupMap &EntryPointsGroups, EntryPointsGroups[GLOBAL_SCOPE_NAME] = {}; } +// For device global variables with the 'device_image_scope' property, +// the function checks that there are no usages of a single device global +// variable from kernels grouped to different modules. +void checkImageScopedDeviceGlobals(const Module &M, + const EntryPointGroupMap &GMap) { + // Early exit if there is only one group + if (GMap.size() < 2) + return; + + // Reverse the EntryPointGroupMap to get a map of entry point -> module's name + unsigned EntryPointNumber = 0; + for (const auto &Group : GMap) { + EntryPointNumber += static_cast(Group.second.size()); + } + DenseMap EntryPointModules( + EntryPointNumber); + for (const auto &Group : GMap) { + auto *ModuleName = &Group.first; + for (const auto *F : Group.second) { + EntryPointModules.insert({F, ModuleName}); + } + } + + // Processing device global variables with the "device_image_scope" property + for (auto &GV : M.globals()) { + if (!isDeviceGlobalVariable(GV) || !hasDeviceImageScopeProperty(GV)) + continue; + + const StringRef *VarEntryPointModule = nullptr; + auto CheckEntryPointModule = [&VarEntryPointModule, &EntryPointModules, + &GV](const auto *F) { + auto EntryPointModulesIt = EntryPointModules.find(F); + assert(EntryPointModulesIt != EntryPointModules.end() + && "There is no group for an entry point"); + if (VarEntryPointModule == nullptr) { + VarEntryPointModule = EntryPointModulesIt->second; + return; + } + if (EntryPointModulesIt->second != VarEntryPointModule) { + error("device_global variable '" + Twine(GV.getName()) + + "' with property \"device_image_scope\" is contained in more " + "than one device image."); + } + }; + + SmallVector Workqueue; + SmallPtrSet Visited; + for (auto *U : GV.users()) + Workqueue.push_back(U); + + while (!Workqueue.empty()) { + const User *U = Workqueue.back(); + Workqueue.pop_back(); + if (!Visited.insert(U).second) + continue; + if (auto *I = dyn_cast(U)) { + auto *F = I->getFunction(); + if (!Visited.contains(F)) + Workqueue.push_back(F); + } else { + if (auto *F = dyn_cast(U)) { + if (isEntryPoint(*F)) { + CheckEntryPointModule(F); + } + } + for (auto *UU : U->users()) { + if (!Visited.contains(UU)) + Workqueue.push_back(UU); + } + } + } + } +} + // This function traverses over reversed call graph by BFS algorithm. // It means that an edge links some function @func with functions // which contain call of function @func. It starts from @@ -415,7 +489,7 @@ TraverseCGToFindSPIRKernels(const Function *StartingFunction) { } std::vector getKernelNamesUsingAssert(const Module &M) { - auto DevicelibAssertFailFunction = M.getFunction("__devicelib_assert_fail"); + auto *DevicelibAssertFailFunction = M.getFunction("__devicelib_assert_fail"); if (!DevicelibAssertFailFunction) return {}; @@ -478,18 +552,25 @@ extractCallGraph(const Module &M, const EntryPointGroup &ModuleEntryPoints) { const Function *F = &*Workqueue.back(); Workqueue.pop_back(); for (const auto &I : instructions(F)) { - if (const CallBase *CB = dyn_cast(&I)) - if (const Function *CF = CB->getCalledFunction()) + if (const CallBase *CB = dyn_cast(&I)) { + if (const Function *CF = CB->getCalledFunction()) { if (!CF->isDeclaration() && !GVs.count(CF)) { GVs.insert(CF); Workqueue.push_back(CF); } + } + } } } // It's not easy to trace global variable's uses inside needed functions // because global variable can be used inside a combination of operators, so // mark all global variables as needed and remove dead ones after cloning. + // Notice. For device global variables with the 'device_image_scope' property, + // removing dead ones is a must, the 'checkImageScopedDeviceGlobals' function + // checks that there are no usages of a single device global variable with the + // 'device_image_scope' property from multiple modules and the splitter must + // not add such usages after the check. for (const auto &G : M.globals()) { GVs.insert(&G); } @@ -731,6 +812,8 @@ class ModuleSplitter { EntryPointsGroupScope Scope) : InputModule(std::move(M)), IsSplit(Split) { groupEntryPoints(*InputModule, GMap, Scope); + if (DeviceGlobals) + checkImageScopedDeviceGlobals(*InputModule, GMap); assert(!GMap.empty() && "Entry points group map is empty!"); GMapIt = GMap.cbegin(); } @@ -745,9 +828,9 @@ class ModuleSplitter { ++GMapIt; std::unique_ptr SplitModule{nullptr}; - if (IsSplit && !SplitModuleEntryPoints.empty()) + if (IsSplit && !SplitModuleEntryPoints.empty()) { SplitModule = extractCallGraph(*InputModule, SplitModuleEntryPoints); - else { + } else { assert(GMap.size() == 1 && "Too many entry points groups in map!"); SplitModule = std::move(InputModule); } From 19fbd7dcbfb8b14cbcecda3ab08e715f27384cec Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Wed, 9 Feb 2022 15:08:41 +0300 Subject: [PATCH 05/17] [sycl-post-link] Add lit tests for a multi-module application --- ...bal_variable_many_kernels_in_one_module.ll | 189 ++++++++++++++++++ ...bal_variable_many_modules_no_dev_global.ll | 178 +++++++++++++++++ ..._variable_many_modules_no_dev_img_scope.ll | 182 +++++++++++++++++ ...bal_variable_many_modules_one_var_error.ll | 188 +++++++++++++++++ ...lobal_variable_many_modules_two_vars_ok.ll | 179 +++++++++++++++++ 5 files changed, 916 insertions(+) create mode 100644 llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll create mode 100644 llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll create mode 100644 llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll create mode 100644 llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll create mode 100644 llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll new file mode 100644 index 0000000000000..704d3d7aa4b3d --- /dev/null +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll @@ -0,0 +1,189 @@ +; RUN: sycl-post-link --device-globals --split=source -S %s -o %t.files.table +; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD0 +; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD1 + +; This test is intended to check that sycl-post-link generates no errors +; when a device global variable with the 'device_image_scope' property +; is used in many kernels grouped to one module. + +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" + +%"class.cl::sycl::ext::oneapi::device_global" = type { i32 } +%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } +%"class.cl::sycl::detail::array" = type { [1 x i64] } +%"class.cl::sycl::detail::accessor_common" = type { i8 } + +$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZN2cl4sycl2idILi1EEC2Ev = comdat any +$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any +$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel4RN2cl4sycl5queueEEUlvE_ = comdat any + +$dg_int2 = comdat any +@dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 #0 +; CHECK-MOD0: @dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 +; CHECK-MOD1-NOT: @dg_int2 + +; Third kernel that uses no device-global variables +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)* + call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #7 + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent mustprogress noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #2 comdat align 2 { +entry: + call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #7 + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel4RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #7 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #7 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + call spir_func void @_Z14kernel1_level1v() #7 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define dso_local spir_func void @_Z14kernel1_level1v() #4 { +entry: + %dg_int_ptr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %dg_int_ptr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %dg_int_ptr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*), %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %dg_int_ptr.ascast, align 8 + %0 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %dg_int_ptr.ascast, align 8 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %0) #9 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +entry: + %retval = alloca i32 addrspace(4)*, align 8 + %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)* + %this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0 + ret i32 addrspace(4)* %val +} + +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #7 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +entry: + %retval = alloca i32 addrspace(4)*, align 8 + %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)* + %this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0 + ret i32 addrspace(4)* %val +} + +attributes #0 = { "sycl-unique-id"="dg_int2" "device_image_scope"="true" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" } +attributes #1 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #3 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #5 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } +attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } +attributes #7 = { convergent } +attributes #8 = { convergent nounwind } +attributes #9 = { nobuiltin allocsize(0) "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" } + +!llvm.dependent-libraries = !{!0} +!opencl.spir.version = !{!1} +!spirv.Source = !{!2} +!llvm.ident = !{!3} +!llvm.module.flags = !{!4, !5} + +!0 = !{!"libcpmt"} +!1 = !{i32 1, i32 2} +!2 = !{i32 4, i32 100000} +!3 = !{!"clang version 14.0.0"} +!4 = !{i32 1, !"wchar_size", i32 2} +!5 = !{i32 7, !"frame-pointer", i32 2} +!6 = !{} diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll new file mode 100644 index 0000000000000..17fbf87c7b6e5 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll @@ -0,0 +1,178 @@ +; RUN: sycl-post-link --device-globals --split=source -S %s -o %t.files.table +; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD0 +; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD1 +; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD2 + +; This test is intended to check that sycl-post-link generates no error if the +; 'device_image_scope' property is attached to not a device global variable. + +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" + +%"class.cl::sycl::ext::oneapi::device_global" = type { i32 } +%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } +%"class.cl::sycl::detail::array" = type { [1 x i64] } +%"class.cl::sycl::detail::accessor_common" = type { i8 } + +$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZN2cl4sycl2idILi1EEC2Ev = comdat any +$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any +$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any + +$dg_int2 = comdat any +@dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 #0 +; CHECK-MOD0: @dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 +; CHECK-MOD1: @dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 +; CHECK-MOD2-NOT: @dg_int2 + +; Third kernel that uses no device-global variables +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #7 comdat !kernel_arg_buffer_location !6 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)* + call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #8 + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent mustprogress noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #2 comdat align 2 { +entry: + call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #8 + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + call spir_func void @_Z14kernel1_level1v() #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define dso_local spir_func void @_Z14kernel1_level1v() #4 { +entry: + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #9 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +entry: + %retval = alloca i32 addrspace(4)*, align 8 + %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)* + %this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0 + ret i32 addrspace(4)* %val +} + +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #9 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +entry: + %retval = alloca i32 addrspace(4)*, align 8 + %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)* + %this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0 + ret i32 addrspace(4)* %val +} + +; This is not a device_global variable: no "sycl-device-global-size" property +attributes #0 = { "sycl-unique-id"="dg_int2" "device_image_scope" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" } +attributes #1 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #3 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #5 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } +attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } +attributes #7 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } +attributes #8 = { convergent } +attributes #9 = { convergent nounwind } + +!llvm.dependent-libraries = !{!0} +!opencl.spir.version = !{!1} +!spirv.Source = !{!2} +!llvm.ident = !{!3} +!llvm.module.flags = !{!4, !5} + +!0 = !{!"libcpmt"} +!1 = !{i32 1, i32 2} +!2 = !{i32 4, i32 100000} +!3 = !{!"clang version 14.0.0"} +!4 = !{i32 1, !"wchar_size", i32 2} +!5 = !{i32 7, !"frame-pointer", i32 2} +!6 = !{} diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll new file mode 100644 index 0000000000000..ab71595e62b80 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll @@ -0,0 +1,182 @@ +; RUN: sycl-post-link --device-globals --split=source -S %s -o %t.files.table +; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD0 +; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD1 +; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD2 + +; This test is intended to check that sycl-post-link generates no errors even +; when a single device global variable but without the 'device_image_scope' +; property is used from multiple device images. + +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" + +%"class.cl::sycl::ext::oneapi::device_global" = type { i32 } +%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } +%"class.cl::sycl::detail::array" = type { [1 x i64] } +%"class.cl::sycl::detail::accessor_common" = type { i8 } + +$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZN2cl4sycl2idILi1EEC2Ev = comdat any +$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any +$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any + +$dg_int2 = comdat any +@dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 #0 +; CHECK-MOD0: @dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 +; CHECK-MOD1: @dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 +; CHECK-MOD2-NOT: @dg_int2 + +; Third kernel that uses no device-global variables +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #7 comdat !kernel_arg_buffer_location !6 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)* + call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #8 + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent mustprogress noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #2 comdat align 2 { +entry: + call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #8 + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + call spir_func void @_Z14kernel1_level1v() #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define dso_local spir_func void @_Z14kernel1_level1v() #4 { +entry: + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #9 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +entry: + %retval = alloca i32 addrspace(4)*, align 8 + %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)* + %this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0 + ret i32 addrspace(4)* %val +} + +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #9 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +entry: + %retval = alloca i32 addrspace(4)*, align 8 + %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)* + %this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0 + ret i32 addrspace(4)* %val +} + +; This device_global variable has no "device_image_scope" property +attributes #0 = { "sycl-unique-id"="dg_int2" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" } +attributes #1 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #3 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #5 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } +attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } +attributes #7 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } +attributes #8 = { convergent } +attributes #9 = { convergent nounwind } + +!llvm.dependent-libraries = !{!0} +!opencl.spir.version = !{!1} +!spirv.Source = !{!2} +!llvm.ident = !{!3} +!llvm.module.flags = !{!4, !5} + +!0 = !{!"libcpmt"} +!1 = !{i32 1, i32 2} +!2 = !{i32 4, i32 100000} +!3 = !{!"clang version 14.0.0"} +!4 = !{i32 1, !"wchar_size", i32 2} +!5 = !{i32 7, !"frame-pointer", i32 2} +!6 = !{} +; CHECK-MOD0: !{i32 6147, i32 1, !"dg_int2"} +; CHECK-MOD1: !{i32 6147, i32 1, !"dg_int2"} +; CHECK-MOD2-NOT: !{i32 6147, i32 1, !"dg_int2"} diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll new file mode 100644 index 0000000000000..63651e9588dbd --- /dev/null +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll @@ -0,0 +1,188 @@ +; RUN: not sycl-post-link --device-globals --split=source %s -o %t.files.table 2>&1 | FileCheck %s + +; This test is intended to check that sycl-post-link does not allow to use a +; single device global variable with the 'device_image_scope' property from +; multiple device images. + +; CHECK: sycl-post-link: device_global variable 'dg_int2' with property "device_image_scope" is contained in more than one device image. + +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" + +%"class.cl::sycl::ext::oneapi::device_global" = type { i32 } +%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } +%"class.cl::sycl::detail::array" = type { [1 x i64] } +%"class.cl::sycl::detail::accessor_common" = type { i8 } + +$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZN2cl4sycl2idILi1EEC2Ev = comdat any +$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any +$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel4RN2cl4sycl5queueEEUlvE_ = comdat any + +$dg_int2 = comdat any +@dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 #0 + +; Third kernel that uses no device-global variables +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #7 comdat !kernel_arg_buffer_location !6 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)* + call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #8 + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent mustprogress noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #2 comdat align 2 { +entry: + call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #8 + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +entry: + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel4RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + call spir_func void @_Z14kernel1_level1v() #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define dso_local spir_func void @_Z14kernel1_level1v() #4 { +entry: + %dg_int_ptr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %dg_int_ptr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %dg_int_ptr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*), %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %dg_int_ptr.ascast, align 8 + %0 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %dg_int_ptr.ascast, align 8 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %0) #10 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +entry: + %retval = alloca i32 addrspace(4)*, align 8 + %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)* + %this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0 + ret i32 addrspace(4)* %val +} + +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #9 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +entry: + %retval = alloca i32 addrspace(4)*, align 8 + %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)* + %this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0 + ret i32 addrspace(4)* %val +} + +attributes #0 = { "sycl-unique-id"="dg_int2" "device_image_scope"="true" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" } +attributes #1 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #3 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #5 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } +attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } +attributes #7 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } +attributes #8 = { convergent } +attributes #9 = { convergent nounwind } +attributes #10 = { nobuiltin allocsize(0) "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" } + +!llvm.dependent-libraries = !{!0} +!opencl.spir.version = !{!1} +!spirv.Source = !{!2} +!llvm.ident = !{!3} +!llvm.module.flags = !{!4, !5} + +!0 = !{!"libcpmt"} +!1 = !{i32 1, i32 2} +!2 = !{i32 4, i32 100000} +!3 = !{!"clang version 14.0.0"} +!4 = !{i32 1, !"wchar_size", i32 2} +!5 = !{i32 7, !"frame-pointer", i32 2} +!6 = !{} diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll new file mode 100644 index 0000000000000..d2ce2ee7f9387 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll @@ -0,0 +1,179 @@ +; RUN: sycl-post-link --device-globals --split=source -S %s -o %t.files.table +; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD0 +; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD1 + +; This test is intended to check that sycl-post-link generates no errors +; when each device global variable with the 'device_image_scope' property +; is used in a single module only. + +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" + +%"class.cl::sycl::ext::oneapi::device_global" = type { i32 } +%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } +%"class.cl::sycl::detail::array" = type { [1 x i64] } +%"class.cl::sycl::detail::accessor_common" = type { i8 } + +$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZN2cl4sycl2idILi1EEC2Ev = comdat any +$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any +$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any +$_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any +$_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any + +@dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 4 #0 +; CHECK-MOD0: @dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 4 +; CHECK-MOD1-NOT: @dg_int2 +@dg_int3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 4 #1 +; CHECK-MOD1: @dg_int3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 4 +; CHECK-MOD0-NOT: @dg_int3 + +; Third kernel that uses no device-global variables +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #8 comdat !kernel_arg_buffer_location !6 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #4 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #4 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)* + call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #9 + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent mustprogress noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #3 comdat align 2 { +entry: + call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #9 + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #4 comdat align 2 { +entry: + ret void +} + +; Function is required to increase the function list only +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #9 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #5 align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + call spir_func void @_Z14kernel1_level1v() #9 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define dso_local spir_func void @_Z14kernel1_level1v() #5 { +entry: + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #10 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #5 align 2 { +entry: + %retval = alloca i32 addrspace(4)*, align 8 + %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)* + %this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0 + ret i32 addrspace(4)* %val +} + +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #7 comdat !kernel_arg_buffer_location !6 { +entry: + %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 + %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* + call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #9 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #5 align 2 { +entry: + %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int3 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #10 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #5 align 2 { +entry: + %retval = alloca i32 addrspace(4)*, align 8 + %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 + %retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)* + %this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* + store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0 + ret i32 addrspace(4)* %val +} + +attributes #0 = { "sycl-unique-id"="dg_int2" "device_image_scope"="true" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" } +attributes #1 = { "sycl-unique-id"="dg_int3" "device_image_scope"="true" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" } +attributes #2 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #3 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #4 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #5 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } +attributes #7 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } +attributes #8 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } +attributes #9 = { convergent } +attributes #10 = { convergent nounwind } + +!llvm.dependent-libraries = !{!0} +!opencl.spir.version = !{!1} +!spirv.Source = !{!2} +!llvm.ident = !{!3} +!llvm.module.flags = !{!4, !5} + +!0 = !{!"libcpmt"} +!1 = !{i32 1, i32 2} +!2 = !{i32 4, i32 100000} +!3 = !{!"clang version 14.0.0"} +!4 = !{i32 1, !"wchar_size", i32 2} +!5 = !{i32 7, !"frame-pointer", i32 2} +!6 = !{} From 2cf821182636d7d79e498667503d1069a03d6183 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Wed, 9 Feb 2022 15:27:28 +0300 Subject: [PATCH 06/17] [sycl-post-link][NFC] Fix formatting issues --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 1c7031371799a..2e667428beef3 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -397,8 +397,8 @@ void checkImageScopedDeviceGlobals(const Module &M, auto CheckEntryPointModule = [&VarEntryPointModule, &EntryPointModules, &GV](const auto *F) { auto EntryPointModulesIt = EntryPointModules.find(F); - assert(EntryPointModulesIt != EntryPointModules.end() - && "There is no group for an entry point"); + assert(EntryPointModulesIt != EntryPointModules.end() && + "There is no group for an entry point"); if (VarEntryPointModule == nullptr) { VarEntryPointModule = EntryPointModulesIt->second; return; From 2fb04a6e62936dccd868ff82bf99d4d816951ad4 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Fri, 11 Feb 2022 12:00:12 +0300 Subject: [PATCH 07/17] Use a clearer wording in a comment Co-authored-by: Alexey Sachkov --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 2e667428beef3..5943dd73a9add 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -405,7 +405,7 @@ void checkImageScopedDeviceGlobals(const Module &M, } if (EntryPointModulesIt->second != VarEntryPointModule) { error("device_global variable '" + Twine(GV.getName()) + - "' with property \"device_image_scope\" is contained in more " + "' with property \"device_image_scope\" is used in more " "than one device image."); } }; From 625166cb421bd7a1a37a3b2c015d86977ff6328c Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Fri, 11 Feb 2022 12:55:12 +0300 Subject: [PATCH 08/17] [sycl-post-link] Use Optional instead of StringRef* --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 5943dd73a9add..952c666ed1477 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -379,10 +379,9 @@ void checkImageScopedDeviceGlobals(const Module &M, for (const auto &Group : GMap) { EntryPointNumber += static_cast(Group.second.size()); } - DenseMap EntryPointModules( - EntryPointNumber); + DenseMap EntryPointModules(EntryPointNumber); for (const auto &Group : GMap) { - auto *ModuleName = &Group.first; + auto ModuleName = Group.first; for (const auto *F : Group.second) { EntryPointModules.insert({F, ModuleName}); } @@ -393,13 +392,13 @@ void checkImageScopedDeviceGlobals(const Module &M, if (!isDeviceGlobalVariable(GV) || !hasDeviceImageScopeProperty(GV)) continue; - const StringRef *VarEntryPointModule = nullptr; + Optional VarEntryPointModule{}; auto CheckEntryPointModule = [&VarEntryPointModule, &EntryPointModules, &GV](const auto *F) { auto EntryPointModulesIt = EntryPointModules.find(F); assert(EntryPointModulesIt != EntryPointModules.end() && "There is no group for an entry point"); - if (VarEntryPointModule == nullptr) { + if (!VarEntryPointModule.hasValue()) { VarEntryPointModule = EntryPointModulesIt->second; return; } From d8e9ea0c5bcd776993c0ba559e1fc38f8e58429c Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Fri, 11 Feb 2022 12:55:19 +0300 Subject: [PATCH 09/17] Revert "Use a clearer wording in a comment" This reverts commit 2fb04a6e62936dccd868ff82bf99d4d816951ad4. --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 952c666ed1477..006ff9d01cfd6 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -404,7 +404,7 @@ void checkImageScopedDeviceGlobals(const Module &M, } if (EntryPointModulesIt->second != VarEntryPointModule) { error("device_global variable '" + Twine(GV.getName()) + - "' with property \"device_image_scope\" is used in more " + "' with property \"device_image_scope\" is contained in more " "than one device image."); } }; From 98bbe22e724919bc81bae8ea8fb4ecf96158cf6e Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Fri, 11 Feb 2022 13:19:23 +0300 Subject: [PATCH 10/17] [sycl-post-link] Use SmallSetVector for graph traversing --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 19 ++++++------------- 1 file changed, 6 insertions(+), 13 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 006ff9d01cfd6..3606a208c9d5c 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -409,30 +409,23 @@ void checkImageScopedDeviceGlobals(const Module &M, } }; - SmallVector Workqueue; - SmallPtrSet Visited; + SmallSetVector Workqueue; for (auto *U : GV.users()) - Workqueue.push_back(U); + Workqueue.insert(U); while (!Workqueue.empty()) { const User *U = Workqueue.back(); Workqueue.pop_back(); - if (!Visited.insert(U).second) - continue; if (auto *I = dyn_cast(U)) { auto *F = I->getFunction(); - if (!Visited.contains(F)) - Workqueue.push_back(F); + Workqueue.insert(F); } else { if (auto *F = dyn_cast(U)) { - if (isEntryPoint(*F)) { + if (isEntryPoint(*F)) CheckEntryPointModule(F); - } - } - for (auto *UU : U->users()) { - if (!Visited.contains(UU)) - Workqueue.push_back(UU); } + for (auto *UU : U->users()) + Workqueue.insert(UU); } } } From f278edd05e8669a982272d7161c225f6f52d3cf7 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Fri, 11 Feb 2022 13:24:14 +0300 Subject: [PATCH 11/17] [sycl-post-link] Use continue for early exit from graph traversing Signed-off-by: Pavel Samolysov --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 3606a208c9d5c..70407369f3655 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -419,14 +419,14 @@ void checkImageScopedDeviceGlobals(const Module &M, if (auto *I = dyn_cast(U)) { auto *F = I->getFunction(); Workqueue.insert(F); - } else { - if (auto *F = dyn_cast(U)) { - if (isEntryPoint(*F)) - CheckEntryPointModule(F); - } - for (auto *UU : U->users()) - Workqueue.insert(UU); + continue; + } + if (auto *F = dyn_cast(U)) { + if (isEntryPoint(*F)) + CheckEntryPointModule(F); } + for (auto *UU : U->users()) + Workqueue.insert(UU); } } } From 5d09f1986d9208d30a708f075745beb54540e860 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Fri, 11 Feb 2022 13:49:43 +0300 Subject: [PATCH 12/17] [sycl-post-link] Remove unused functions from the lit tests Signed-off-by: Pavel Samolysov --- ...bal_variable_many_kernels_in_one_module.ll | 94 ++++-------------- ...bal_variable_many_modules_no_dev_global.ll | 91 ++++-------------- ..._variable_many_modules_no_dev_img_scope.ll | 91 ++++-------------- ...bal_variable_many_modules_one_var_error.ll | 96 +++++-------------- ...lobal_variable_many_modules_two_vars_ok.ll | 91 ++++-------------- 5 files changed, 100 insertions(+), 363 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll index 704d3d7aa4b3d..46f26b0e70ffd 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll @@ -10,15 +10,8 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256: target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global" = type { i32 } -%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } -%"class.cl::sycl::detail::array" = type { [1 x i64] } %"class.cl::sycl::detail::accessor_common" = type { i8 } -$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any -$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any -$_ZN2cl4sycl2idILi1EEC2Ev = comdat any -$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any -$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any $_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any $_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any $_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any @@ -30,95 +23,53 @@ $dg_int2 = comdat any ; CHECK-MOD1-NOT: @dg_int2 ; Third kernel that uses no device-global variables -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 { -entry: - %this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8 - %this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)* - store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)* - call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #7 - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent mustprogress noinline norecurse optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #2 comdat align 2 { -entry: - call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #7 - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #3 comdat !kernel_arg_buffer_location !6 { entry: ret void } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel4RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel4RN2cl4sycl5queueEEUlvE_() #2 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #7 + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #4 ret void } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #2 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #7 + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #4 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - call spir_func void @_Z14kernel1_level1v() #7 + call spir_func void @_Z14kernel1_level1v() #4 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define dso_local spir_func void @_Z14kernel1_level1v() #4 { +define dso_local spir_func void @_Z14kernel1_level1v() #1 { entry: %dg_int_ptr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 %dg_int_ptr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %dg_int_ptr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*), %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %dg_int_ptr.ascast, align 8 %0 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %dg_int_ptr.ascast, align 8 - %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %0) #9 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %0) #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #1 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -131,27 +82,27 @@ entry: } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #2 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #7 + call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #4 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #8 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #1 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -164,15 +115,12 @@ entry: } attributes #0 = { "sycl-unique-id"="dg_int2" "device_image_scope"="true" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" } -attributes #1 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #3 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #5 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } -attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } -attributes #7 = { convergent } -attributes #8 = { convergent nounwind } -attributes #9 = { nobuiltin allocsize(0) "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" } +attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } +attributes #3 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } +attributes #4 = { convergent } +attributes #5 = { convergent nounwind } +attributes #6 = { nobuiltin allocsize(0) "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" } !llvm.dependent-libraries = !{!0} !opencl.spir.version = !{!1} diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll index 17fbf87c7b6e5..ee645f23f3f6f 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_global.ll @@ -10,15 +10,8 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256: target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global" = type { i32 } -%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } -%"class.cl::sycl::detail::array" = type { [1 x i64] } %"class.cl::sycl::detail::accessor_common" = type { i8 } -$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any -$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any -$_ZN2cl4sycl2idILi1EEC2Ev = comdat any -$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any -$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any $_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any $_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any $_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any @@ -30,83 +23,40 @@ $dg_int2 = comdat any ; CHECK-MOD2-NOT: @dg_int2 ; Third kernel that uses no device-global variables -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #7 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #4 comdat !kernel_arg_buffer_location !6 { entry: ret void } -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 { -entry: - %this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8 - %this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)* - store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)* - call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #8 - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent mustprogress noinline norecurse optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #2 comdat align 2 { -entry: - call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #8 - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #2 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - call spir_func void @_Z14kernel1_level1v() #8 + call spir_func void @_Z14kernel1_level1v() #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define dso_local spir_func void @_Z14kernel1_level1v() #4 { +define dso_local spir_func void @_Z14kernel1_level1v() #1 { entry: - %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #9 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #1 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -119,27 +69,27 @@ entry: } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #3 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #9 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #1 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -153,15 +103,12 @@ entry: ; This is not a device_global variable: no "sycl-device-global-size" property attributes #0 = { "sycl-unique-id"="dg_int2" "device_image_scope" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" } -attributes #1 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #3 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #5 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } -attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } -attributes #7 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } -attributes #8 = { convergent } -attributes #9 = { convergent nounwind } +attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } +attributes #3 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } +attributes #4 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } +attributes #5 = { convergent } +attributes #6 = { convergent nounwind } !llvm.dependent-libraries = !{!0} !opencl.spir.version = !{!1} diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll index ab71595e62b80..00c2218f3d0ca 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll @@ -11,15 +11,8 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256: target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global" = type { i32 } -%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } -%"class.cl::sycl::detail::array" = type { [1 x i64] } %"class.cl::sycl::detail::accessor_common" = type { i8 } -$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any -$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any -$_ZN2cl4sycl2idILi1EEC2Ev = comdat any -$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any -$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any $_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any $_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any $_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any @@ -31,83 +24,40 @@ $dg_int2 = comdat any ; CHECK-MOD2-NOT: @dg_int2 ; Third kernel that uses no device-global variables -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #7 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #4 comdat !kernel_arg_buffer_location !6 { entry: ret void } -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 { -entry: - %this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8 - %this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)* - store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)* - call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #8 - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent mustprogress noinline norecurse optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #2 comdat align 2 { -entry: - call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #8 - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #2 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - call spir_func void @_Z14kernel1_level1v() #8 + call spir_func void @_Z14kernel1_level1v() #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define dso_local spir_func void @_Z14kernel1_level1v() #4 { +define dso_local spir_func void @_Z14kernel1_level1v() #1 { entry: - %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #9 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #1 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -120,27 +70,27 @@ entry: } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #3 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #9 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #1 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -154,15 +104,12 @@ entry: ; This device_global variable has no "device_image_scope" property attributes #0 = { "sycl-unique-id"="dg_int2" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" } -attributes #1 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #3 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #5 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } -attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } -attributes #7 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } -attributes #8 = { convergent } -attributes #9 = { convergent nounwind } +attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } +attributes #3 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } +attributes #4 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } +attributes #5 = { convergent } +attributes #6 = { convergent nounwind } !llvm.dependent-libraries = !{!0} !opencl.spir.version = !{!1} diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll index 63651e9588dbd..be15b4375947d 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll @@ -10,15 +10,8 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256: target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global" = type { i32 } -%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } -%"class.cl::sycl::detail::array" = type { [1 x i64] } %"class.cl::sycl::detail::accessor_common" = type { i8 } -$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any -$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any -$_ZN2cl4sycl2idILi1EEC2Ev = comdat any -$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any -$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any $_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any $_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any $_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any @@ -28,95 +21,53 @@ $dg_int2 = comdat any @dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 #0 ; Third kernel that uses no device-global variables -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #7 comdat !kernel_arg_buffer_location !6 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 { -entry: - %this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8 - %this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)* - store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)* - call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #8 - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent mustprogress noinline norecurse optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #2 comdat align 2 { -entry: - call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #8 - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #4 comdat !kernel_arg_buffer_location !6 { entry: ret void } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel4RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel4RN2cl4sycl5queueEEUlvE_() #3 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #2 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - call spir_func void @_Z14kernel1_level1v() #8 + call spir_func void @_Z14kernel1_level1v() #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define dso_local spir_func void @_Z14kernel1_level1v() #4 { +define dso_local spir_func void @_Z14kernel1_level1v() #1 { entry: %dg_int_ptr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 %dg_int_ptr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %dg_int_ptr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*), %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %dg_int_ptr.ascast, align 8 %0 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %dg_int_ptr.ascast, align 8 - %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %0) #10 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %0) #7 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #1 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -129,27 +80,27 @@ entry: } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #2 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #8 + call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #5 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #9 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #1 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -162,16 +113,13 @@ entry: } attributes #0 = { "sycl-unique-id"="dg_int2" "device_image_scope"="true" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" } -attributes #1 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #3 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #5 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } -attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } -attributes #7 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } -attributes #8 = { convergent } -attributes #9 = { convergent nounwind } -attributes #10 = { nobuiltin allocsize(0) "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" } +attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } +attributes #3 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } +attributes #4 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } +attributes #5 = { convergent } +attributes #6 = { convergent nounwind } +attributes #7 = { nobuiltin allocsize(0) "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" } !llvm.dependent-libraries = !{!0} !opencl.spir.version = !{!1} diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll index d2ce2ee7f9387..f12cabd6e2250 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll @@ -10,15 +10,8 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256: target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global" = type { i32 } -%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" } -%"class.cl::sycl::detail::array" = type { [1 x i64] } %"class.cl::sycl::detail::accessor_common" = type { i8 } -$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any -$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any -$_ZN2cl4sycl2idILi1EEC2Ev = comdat any -$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any -$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any $_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any $_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any $_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any @@ -31,83 +24,40 @@ $_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any ; CHECK-MOD0-NOT: @dg_int3 ; Third kernel that uses no device-global variables -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #8 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 { entry: ret void } -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #4 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #4 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 { -entry: - %this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8 - %this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)* - store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)* - call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #9 - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent mustprogress noinline norecurse optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #3 comdat align 2 { -entry: - call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #9 - ret void -} - -; Function is required to increase the function list only -; Function Attrs: convergent noinline norecurse nounwind optnone -define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #4 comdat align 2 { -entry: - ret void -} - -; Function is required to increase the function list only ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #3 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #9 + call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #5 align 2 { +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #2 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - call spir_func void @_Z14kernel1_level1v() #9 + call spir_func void @_Z14kernel1_level1v() #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define dso_local spir_func void @_Z14kernel1_level1v() #5 { +define dso_local spir_func void @_Z14kernel1_level1v() #2 { entry: - %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #10 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #7 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #5 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #2 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -120,27 +70,27 @@ entry: } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #7 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #4 comdat !kernel_arg_buffer_location !6 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* - call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #9 + call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #5 align 2 { +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #2 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int3 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #10 + %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int3 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #7 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #5 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #2 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -154,15 +104,12 @@ entry: attributes #0 = { "sycl-unique-id"="dg_int2" "device_image_scope"="true" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" } attributes #1 = { "sycl-unique-id"="dg_int3" "device_image_scope"="true" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" } -attributes #2 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #3 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #4 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #5 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } -attributes #7 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } -attributes #8 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } -attributes #9 = { convergent } -attributes #10 = { convergent nounwind } +attributes #2 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #3 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" } +attributes #4 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" } +attributes #5 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" } +attributes #6 = { convergent } +attributes #7 = { convergent nounwind } !llvm.dependent-libraries = !{!0} !opencl.spir.version = !{!1} From 29a842dfe52dcb9e91f2418a30f7d0e4bb5ef0c5 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Fri, 11 Feb 2022 14:25:56 +0300 Subject: [PATCH 13/17] [sycl-post-link] Fix Optional comparison Signed-off-by: Pavel Samolysov --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 70407369f3655..37acdf33a0b4b 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -402,7 +402,7 @@ void checkImageScopedDeviceGlobals(const Module &M, VarEntryPointModule = EntryPointModulesIt->second; return; } - if (EntryPointModulesIt->second != VarEntryPointModule) { + if (EntryPointModulesIt->second != *VarEntryPointModule) { error("device_global variable '" + Twine(GV.getName()) + "' with property \"device_image_scope\" is contained in more " "than one device image."); From 7d40f39e763dad0745f1c8b19d10ce1196d11331 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Fri, 11 Feb 2022 16:43:01 +0300 Subject: [PATCH 14/17] [NFC] Change error message according to the design document Co-authored-by: Alexey Sachkov --- .../test_global_variable_many_modules_one_var_error.ll | 2 +- llvm/tools/sycl-post-link/sycl-post-link.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll index be15b4375947d..890618e10f875 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll @@ -4,7 +4,7 @@ ; single device global variable with the 'device_image_scope' property from ; multiple device images. -; CHECK: sycl-post-link: device_global variable 'dg_int2' with property "device_image_scope" is contained in more than one device image. +; CHECK: sycl-post-link: device_global variable 'dg_int2' with property "device_image_scope" is used in more than one device image. 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" diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 37acdf33a0b4b..7fde4e1138acc 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -404,7 +404,7 @@ void checkImageScopedDeviceGlobals(const Module &M, } if (EntryPointModulesIt->second != *VarEntryPointModule) { error("device_global variable '" + Twine(GV.getName()) + - "' with property \"device_image_scope\" is contained in more " + "' with property \"device_image_scope\" is used in more " "than one device image."); } }; From 1b508b36b9c278ee082cc8ec96e2b06d36b563d2 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Fri, 11 Feb 2022 18:08:59 +0300 Subject: [PATCH 15/17] Remove redundant braces Co-authored-by: Mikhail Lychkov --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 7fde4e1138acc..db395dc91f152 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -376,9 +376,8 @@ void checkImageScopedDeviceGlobals(const Module &M, // Reverse the EntryPointGroupMap to get a map of entry point -> module's name unsigned EntryPointNumber = 0; - for (const auto &Group : GMap) { + for (const auto &Group : GMap) EntryPointNumber += static_cast(Group.second.size()); - } DenseMap EntryPointModules(EntryPointNumber); for (const auto &Group : GMap) { auto ModuleName = Group.first; From 74caa294ec452995c6f4e100d4f1ae77666c5de4 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Mon, 14 Feb 2022 10:15:37 +0300 Subject: [PATCH 16/17] Remove additional responsibilities from the ModuleSplitter class As it has been suggested by @mlychkov, the only responsibility of the ModuleSplitter class should be splitting. Therefore, the class should not group entry points by modules as well as do any checks if the splitting is allowed at all. The following refactoring has been done: all the checks and entry point grouping were removed from the class' constructor. The class constructor takes the group map by value and takes ownership through moving. If any check is failed, no instances of the ModuleSplitter class are instantiated at all. --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 24 ++++++++++++-------- 1 file changed, 14 insertions(+), 10 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index db395dc91f152..fb9cd474dd334 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -328,8 +328,9 @@ bool isEntryPoint(const Function &F) { // EntryPointsGroups which maps some key to a group of entry points. Each such // group along with IR it depends on (globals, functions from its call graph, // ...) will constitute a separate module. -void groupEntryPoints(const Module &M, EntryPointGroupMap &EntryPointsGroups, - EntryPointsGroupScope EntryScope) { +EntryPointGroupMap groupEntryPoints(const Module &M, + EntryPointsGroupScope EntryScope) { + EntryPointGroupMap EntryPointsGroups{}; // Only process module entry points: for (const auto &F : M.functions()) { if (!isEntryPoint(F)) @@ -363,6 +364,8 @@ void groupEntryPoints(const Module &M, EntryPointGroupMap &EntryPointsGroups, // No entry points met, record this. if (EntryPointsGroups.empty()) EntryPointsGroups[GLOBAL_SCOPE_NAME] = {}; + + return EntryPointsGroups; } // For device global variables with the 'device_image_scope' property, @@ -794,17 +797,15 @@ bool processCompileTimeProperties(Module &M) { // module as a split condition. class ModuleSplitter { std::unique_ptr InputModule{nullptr}; + bool IsSplit; EntryPointGroupMap GMap; EntryPointGroupMap::const_iterator GMapIt; - bool IsSplit; public: ModuleSplitter(std::unique_ptr M, bool Split, - EntryPointsGroupScope Scope) - : InputModule(std::move(M)), IsSplit(Split) { - groupEntryPoints(*InputModule, GMap, Scope); - if (DeviceGlobals) - checkImageScopedDeviceGlobals(*InputModule, GMap); + EntryPointGroupMap GroupMap) + : InputModule(std::move(M)), IsSplit(Split), + GMap(std::move(GroupMap)) { assert(!GMap.empty() && "Entry points group map is empty!"); GMapIt = GMap.cbegin(); } @@ -857,9 +858,12 @@ TableFiles processOneModule(std::unique_ptr M, bool IsEsimd, if (IsEsimd && LowerEsimd) lowerEsimdConstructs(*M); - EntryPointsGroupScope Scope = selectDeviceCodeGroupScope(*M); + EntryPointGroupMap GMap = + groupEntryPoints(*M, selectDeviceCodeGroupScope(*M)); + if (DeviceGlobals) + checkImageScopedDeviceGlobals(*M, GMap); bool DoSplit = (SplitMode.getNumOccurrences() > 0); - ModuleSplitter MSplit(std::move(M), DoSplit, Scope); + ModuleSplitter MSplit(std::move(M), DoSplit, std::move(GMap)); StringRef FileSuffix = IsEsimd ? "esimd_" : ""; From 6b9655e51a4eb9d64e78c7a17a9ae12c3996c420 Mon Sep 17 00:00:00 2001 From: Pavel Samolysov Date: Mon, 14 Feb 2022 10:22:22 +0300 Subject: [PATCH 17/17] [NFC] Fix formatting inssues --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index fb9cd474dd334..57131bc418c28 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -804,8 +804,7 @@ class ModuleSplitter { public: ModuleSplitter(std::unique_ptr M, bool Split, EntryPointGroupMap GroupMap) - : InputModule(std::move(M)), IsSplit(Split), - GMap(std::move(GroupMap)) { + : InputModule(std::move(M)), IsSplit(Split), GMap(std::move(GroupMap)) { assert(!GMap.empty() && "Entry points group map is empty!"); GMapIt = GMap.cbegin(); }