From 72a603288fc94c39f2dd0d2923a6d960677da888 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Thu, 29 Aug 2024 12:04:15 +0100 Subject: [PATCH 1/7] Don't remove llvm.compiler.used for NVPTX/AMDGCN These are removed automatically during lowering. --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 5 ++- .../DeviceGlobal/device_global_static.cpp | 31 +++++++++++++++++++ .../device_global_static.cpp | 26 ++++++++++++++++ 3 files changed, 61 insertions(+), 1 deletion(-) create mode 100644 sycl/test-e2e/DeviceGlobal/device_global_static.cpp create mode 100644 sycl/test/check_device_code/device_global_static.cpp diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 734e71e91183a..6da3eaf1ed0a5 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -788,7 +788,10 @@ processInputModule(std::unique_ptr M) { // its purpose, these device_global variables can be removed. If they are not // used inside the device code after they have been removed from // "llvm.compiler.used" they can be erased safely. - Modified |= removeDeviceGlobalFromCompilerUsed(*M.get()); + if (auto Triple = M->getTargetTriple(); + Triple.find("nvptx") != std::string::npos && + Triple.find("amdgcn") != std::string::npos) + Modified |= removeDeviceGlobalFromCompilerUsed(*M.get()); // Instrument each image scope device globals if the module has been // instrumented by sanitizer pass. diff --git a/sycl/test-e2e/DeviceGlobal/device_global_static.cpp b/sycl/test-e2e/DeviceGlobal/device_global_static.cpp new file mode 100644 index 0000000000000..68c0fc1ede7cb --- /dev/null +++ b/sycl/test-e2e/DeviceGlobal/device_global_static.cpp @@ -0,0 +1,31 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// +// The OpenCL GPU backends do not currently support device_global backend +// calls. +// UNSUPPORTED: opencl && gpu +// +// Tests static device_global access through device kernels. + +#include "common.hpp" + +static device_global DeviceGlobalVar; + +int main() { + queue Q; + + Q.single_task([=]() { DeviceGlobalVar.get()[0] = 42; }); + // Make sure that the write happens before subsequent read + Q.wait(); + + int OutVal = 0; + { + buffer OutBuf(&OutVal, 1); + Q.submit([&](handler &CGH) { + auto OutAcc = OutBuf.get_access(CGH); + CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar.get()[0]; }); + }); + } + assert(OutVal == 42 && "Read value does not match."); + return 0; +} diff --git a/sycl/test/check_device_code/device_global_static.cpp b/sycl/test/check_device_code/device_global_static.cpp new file mode 100644 index 0000000000000..f97550a182704 --- /dev/null +++ b/sycl/test/check_device_code/device_global_static.cpp @@ -0,0 +1,26 @@ +// Tests that the llvm.compiler.used symbol, which is used to implement static +// device globals, is removed at some point in compilation. For SPIR-V this +// symbol is removed at sycl-post-link and for NVPTX/AMDGCN it is removed at +// lowering. + +// UNSUPPORTED: windows + +// RUN: %clangxx -fsycl %s -o %t +// RUN: strings %t | not grep "llvm.compiler.used" + +// RUN: %if cuda %{ %clangxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda %s -o %t %} +// RUN: %if cuda %{ strings %t | not grep "llvm.compiler.used" %} + +// RUN: %if hip_amd %{ %clangxx -fsycl -fsycl-targets=amd_gpu_gfx906 %s -o %t %} +// RUN: %if hip_amd %{ strings %t | not grep "llvm.compiler.used" %} + +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental; + +static device_global DeviceGlobalVar; + +int main() { + sycl::queue{}.single_task([=] { volatile int ReadVal = DeviceGlobalVar; }); +} From 74fa7ba37259bdaf3a29a8813a901a486e99f10c Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Tue, 3 Sep 2024 16:15:49 +0100 Subject: [PATCH 2/7] Don't remove llvm.compiler.used at sycl post link This is removed at lowering for NVPTX/AMDGCN. TODO: check that this is working for SPIR-V. --- .../CodeGen/AMDGPU/remove-compiler-used.ll | 19 +++++ .../CodeGen/NVPTX/remove-compiler-used.ll | 17 +++++ llvm/tools/sycl-post-link/sycl-post-link.cpp | 74 ------------------- 3 files changed, 36 insertions(+), 74 deletions(-) create mode 100644 llvm/test/CodeGen/AMDGPU/remove-compiler-used.ll create mode 100644 llvm/test/CodeGen/NVPTX/remove-compiler-used.ll diff --git a/llvm/test/CodeGen/AMDGPU/remove-compiler-used.ll b/llvm/test/CodeGen/AMDGPU/remove-compiler-used.ll new file mode 100644 index 0000000000000..d33c41a36dddb --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/remove-compiler-used.ll @@ -0,0 +1,19 @@ +;; This test verifies llc on AMDGCN will delete the llvm.compiler.used symbol +;; while keeping the symbol in the outputted ASM. + +; RUN: llc < %s -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 | FileCheck %s +; RUN: llc < %s -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 | FileCheck %s +; RUN: llc < %s -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a | FileCheck %s + +@keep_this = internal global i32 2, align 4 +@llvm.compiler.used = appending global [1 x ptr] [ptr @keep_this], section "llvm.metadata" + +; CHECK-NOT: llvm.metadata +; CHECK-NOT: llvm{{.*}}used +; CHECK-NOT: llvm{{.*}}compiler{{.*}}used + +; CHECK: .type keep_this,@object ; + +; CHECK-NOT: llvm.metadata +; CHECK-NOT: llvm{{.*}}used +; CHECK-NOT: llvm{{.*}}compiler{{.*}}used diff --git a/llvm/test/CodeGen/NVPTX/remove-compiler-used.ll b/llvm/test/CodeGen/NVPTX/remove-compiler-used.ll new file mode 100644 index 0000000000000..ba9030818debd --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/remove-compiler-used.ll @@ -0,0 +1,17 @@ +;; This test verifies llc on NVPTX will delete the llvm.compiler.used symbol +;; while keeping the symbol in the outputted ASM. + +; RUN: llc < %s -march=nvptx64 | FileCheck %s + +@keep_this = internal global i32 2, align 4 +@llvm.compiler.used = appending global [1 x ptr] [ptr @keep_this], section "llvm.metadata" + +; CHECK-NOT: llvm.metadata +; CHECK-NOT: llvm{{.*}}used +; CHECK-NOT: llvm{{.*}}compiler{{.*}}used + +; CHECK: .global .align 4 .u32 keep_this + +; CHECK-NOT: llvm.metadata +; CHECK-NOT: llvm{{.*}}used +; CHECK-NOT: llvm{{.*}}compiler{{.*}}used diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 6da3eaf1ed0a5..1a4d20a16e13f 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -572,70 +572,6 @@ static bool removeSYCLKernelsConstRefArray(Module &M) { return true; } -// Removes all device_global variables from the llvm.compiler.used global -// variable. A device_global with internal linkage will be in llvm.compiler.used -// to avoid the compiler wrongfully removing it during optimizations. However, -// as an effect the device_global variables will also be distributed across -// binaries, even if llvm.compiler.used has served its purpose. To avoid -// polluting other binaries with unused device_global variables, we remove them -// from llvm.compiler.used and erase them if they have no further uses. -static bool removeDeviceGlobalFromCompilerUsed(Module &M) { - GlobalVariable *GV = M.getGlobalVariable("llvm.compiler.used"); - if (!GV) - return false; - - // Erase the old llvm.compiler.used. A new one will be created at the end if - // there are other values in it (other than device_global). - assert(GV->user_empty() && "Unexpected llvm.compiler.used users"); - Constant *Initializer = GV->getInitializer(); - const auto *VAT = cast(GV->getValueType()); - GV->setInitializer(nullptr); - GV->eraseFromParent(); - - // Destroy the initializer. Keep the operands so we keep the ones we need. - SmallVector IOperands; - for (auto It = Initializer->op_begin(); It != Initializer->op_end(); It++) - IOperands.push_back(cast(*It)); - assert(llvm::isSafeToDestroyConstant(Initializer) && - "Cannot remove initializer of llvm.compiler.used global"); - Initializer->destroyConstant(); - - // Iterate through all operands. If they are device_global then we drop them - // and erase them if they have no uses afterwards. All other values are kept. - SmallVector NewOperands; - for (auto It = IOperands.begin(); It != IOperands.end(); It++) { - Constant *Op = *It; - auto *DG = dyn_cast(Op->stripPointerCasts()); - - // If it is not a device_global we keep it. - if (!DG || !isDeviceGlobalVariable(*DG)) { - NewOperands.push_back(Op); - continue; - } - - // Destroy the device_global operand. - if (llvm::isSafeToDestroyConstant(Op)) - Op->destroyConstant(); - - // Remove device_global if it no longer has any uses. - if (!DG->isConstantUsed()) - DG->eraseFromParent(); - } - - // If we have any operands left from the original llvm.compiler.used we create - // a new one with the new size. - if (!NewOperands.empty()) { - ArrayType *ATy = ArrayType::get(VAT->getElementType(), NewOperands.size()); - GlobalVariable *NGV = - new GlobalVariable(M, ATy, false, GlobalValue::AppendingLinkage, - ConstantArray::get(ATy, NewOperands), ""); - NGV->setName("llvm.compiler.used"); - NGV->setSection("llvm.metadata"); - } - - return true; -} - SmallVector handleESIMD(module_split::ModuleDesc &&MDesc, bool &Modified, bool &SplitOccurred) { @@ -783,16 +719,6 @@ processInputModule(std::unique_ptr M) { // actions. Modified |= removeSYCLKernelsConstRefArray(*M.get()); - // There may be device_global variables kept alive in "llvm.compiler.used" - // to keep the optimizer from wrongfully removing them. Since it has served - // its purpose, these device_global variables can be removed. If they are not - // used inside the device code after they have been removed from - // "llvm.compiler.used" they can be erased safely. - if (auto Triple = M->getTargetTriple(); - Triple.find("nvptx") != std::string::npos && - Triple.find("amdgcn") != std::string::npos) - Modified |= removeDeviceGlobalFromCompilerUsed(*M.get()); - // Instrument each image scope device globals if the module has been // instrumented by sanitizer pass. if (isModuleUsingAsan(*M)) From 3f9ae01b9494d6b8f9923e69579ef451cce9358e Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Tue, 3 Sep 2024 17:29:13 +0100 Subject: [PATCH 3/7] Update sycl post link tests sycl post link no longer removes llvm.compiler.used. --- .../device-globals/test_global_variable_drop_used.ll | 1 - .../device-globals/test_global_variable_drop_used_opaque_ptr.ll | 1 - .../device-globals/test_global_variable_trim_used.ll | 2 +- .../device-globals/test_global_variable_trim_used_opaque_ptr.ll | 2 +- 4 files changed, 2 insertions(+), 4 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll index 2e2c6127083e1..f38e1ef16a5a9 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll @@ -13,7 +13,6 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 } %class.anon.0 = type { i8 } -; CHECK-IR-NOT: @llvm.compiler.used = @llvm.compiler.used = appending global [4 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool4 to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7no_dg_int1 to i8 addrspace(4)*)] @_ZL7dg_int1 = weak_odr addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll index 80905a39b9518..817ac2a0a3aab 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll @@ -13,7 +13,6 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 } %class.anon.0 = type { i8 } -; CHECK-IR-NOT: @llvm.compiler.used = @llvm.compiler.used = appending global [4 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL8dg_bool4 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7no_dg_int1 to ptr addrspace(4))] @_ZL7dg_int1 = weak_odr addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll index b4d3bf3f557e1..c606a1dc9f196 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll @@ -11,7 +11,7 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global.0" = type { ptr addrspace(4) } %class.anon.0 = type { i8 } -; CHECK-IR: @llvm.compiler.used = appending global [1 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4))] +; CHECK-IR: @llvm.compiler.used = = appending global [3 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4))] @llvm.compiler.used = appending global [3 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4)), ptr addrspace(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to ptr addrspace(4))] @_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll index 771f8ab1bada8..79a64eb5bb90e 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll @@ -11,7 +11,7 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global.0" = type { ptr addrspace(4) } %class.anon.0 = type { i8 } -; CHECK-IR: @llvm.compiler.used = appending global [1 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4))] +; CHECK-IR: @llvm.compiler.used = appending global [3 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4))] @llvm.compiler.used = appending global [3 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4))] @_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0 From 1b8c225be6929aa7a93103187452d0d1da99fe7f Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Tue, 3 Sep 2024 17:46:41 +0100 Subject: [PATCH 4/7] Reinstate llvm.compiler.used removal at SPL For spirv remove the llvm.compiler.used symbol at sycl-post-link. This is because we use llvm-spirv for lowering, not the SPIR-V backend. --- llvm/tools/sycl-post-link/sycl-post-link.cpp | 71 ++++++++++++++++++++ 1 file changed, 71 insertions(+) diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 1a4d20a16e13f..b59da64120a6c 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -572,6 +572,70 @@ static bool removeSYCLKernelsConstRefArray(Module &M) { return true; } +// Removes all device_global variables from the llvm.compiler.used global +// variable. A device_global with internal linkage will be in llvm.compiler.used +// to avoid the compiler wrongfully removing it during optimizations. However, +// as an effect the device_global variables will also be distributed across +// binaries, even if llvm.compiler.used has served its purpose. To avoid +// polluting other binaries with unused device_global variables, we remove them +// from llvm.compiler.used and erase them if they have no further uses. +static bool removeDeviceGlobalFromCompilerUsed(Module &M) { + GlobalVariable *GV = M.getGlobalVariable("llvm.compiler.used"); + if (!GV) + return false; + + // Erase the old llvm.compiler.used. A new one will be created at the end if + // there are other values in it (other than device_global). + assert(GV->user_empty() && "Unexpected llvm.compiler.used users"); + Constant *Initializer = GV->getInitializer(); + const auto *VAT = cast(GV->getValueType()); + GV->setInitializer(nullptr); + GV->eraseFromParent(); + + // Destroy the initializer. Keep the operands so we keep the ones we need. + SmallVector IOperands; + for (auto It = Initializer->op_begin(); It != Initializer->op_end(); It++) + IOperands.push_back(cast(*It)); + assert(llvm::isSafeToDestroyConstant(Initializer) && + "Cannot remove initializer of llvm.compiler.used global"); + Initializer->destroyConstant(); + + // Iterate through all operands. If they are device_global then we drop them + // and erase them if they have no uses afterwards. All other values are kept. + SmallVector NewOperands; + for (auto It = IOperands.begin(); It != IOperands.end(); It++) { + Constant *Op = *It; + auto *DG = dyn_cast(Op->stripPointerCasts()); + + // If it is not a device_global we keep it. + if (!DG || !isDeviceGlobalVariable(*DG)) { + NewOperands.push_back(Op); + continue; + } + + // Destroy the device_global operand. + if (llvm::isSafeToDestroyConstant(Op)) + Op->destroyConstant(); + + // Remove device_global if it no longer has any uses. + if (!DG->isConstantUsed()) + DG->eraseFromParent(); + } + + // If we have any operands left from the original llvm.compiler.used we create + // a new one with the new size. + if (!NewOperands.empty()) { + ArrayType *ATy = ArrayType::get(VAT->getElementType(), NewOperands.size()); + GlobalVariable *NGV = + new GlobalVariable(M, ATy, false, GlobalValue::AppendingLinkage, + ConstantArray::get(ATy, NewOperands), ""); + NGV->setName("llvm.compiler.used"); + NGV->setSection("llvm.metadata"); + } + + return true; +} + SmallVector handleESIMD(module_split::ModuleDesc &&MDesc, bool &Modified, bool &SplitOccurred) { @@ -719,6 +783,13 @@ processInputModule(std::unique_ptr M) { // actions. Modified |= removeSYCLKernelsConstRefArray(*M.get()); + // There may be device_global variables kept alive in "llvm.compiler.used" + // to keep the optimizer from wrongfully removing them. llvm.compiler.used + // symbols are usually removed at backend lowering, but this is handled here + // for SPIR-V since SYCL compilation uses llvm-spirv, not the SPIR-V backend. + if (auto Triple = M->getTargetTriple().find("spir") != std::string::npos) + Modified |= removeDeviceGlobalFromCompilerUsed(*M.get()); + // Instrument each image scope device globals if the module has been // instrumented by sanitizer pass. if (isModuleUsingAsan(*M)) From 4b06c768f546f9bcaef25dd781a3360516adb605 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Tue, 3 Sep 2024 17:50:56 +0100 Subject: [PATCH 5/7] Undo sycl-post-link tests --- .../device-globals/test_global_variable_drop_used.ll | 1 + .../device-globals/test_global_variable_drop_used_opaque_ptr.ll | 1 + .../device-globals/test_global_variable_trim_used.ll | 2 +- .../device-globals/test_global_variable_trim_used_opaque_ptr.ll | 2 +- 4 files changed, 4 insertions(+), 2 deletions(-) diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll index f38e1ef16a5a9..2e2c6127083e1 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll @@ -13,6 +13,7 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 } %class.anon.0 = type { i8 } +; CHECK-IR-NOT: @llvm.compiler.used = @llvm.compiler.used = appending global [4 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool4 to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7no_dg_int1 to i8 addrspace(4)*)] @_ZL7dg_int1 = weak_odr addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll index 817ac2a0a3aab..80905a39b9518 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll @@ -13,6 +13,7 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 } %class.anon.0 = type { i8 } +; CHECK-IR-NOT: @llvm.compiler.used = @llvm.compiler.used = appending global [4 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL8dg_bool4 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7no_dg_int1 to ptr addrspace(4))] @_ZL7dg_int1 = weak_odr addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll index c606a1dc9f196..b4d3bf3f557e1 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll @@ -11,7 +11,7 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global.0" = type { ptr addrspace(4) } %class.anon.0 = type { i8 } -; CHECK-IR: @llvm.compiler.used = = appending global [3 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4))] +; CHECK-IR: @llvm.compiler.used = appending global [1 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4))] @llvm.compiler.used = appending global [3 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4)), ptr addrspace(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to ptr addrspace(4))] @_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0 diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll index 79a64eb5bb90e..771f8ab1bada8 100644 --- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll +++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll @@ -11,7 +11,7 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global.0" = type { ptr addrspace(4) } %class.anon.0 = type { i8 } -; CHECK-IR: @llvm.compiler.used = appending global [3 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4))] +; CHECK-IR: @llvm.compiler.used = appending global [1 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4))] @llvm.compiler.used = appending global [3 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL16NotADeviceGlobal to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4))] @_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0 From d3c9529acf8ac3a9943cae73cdf544d857dedb90 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Wed, 4 Sep 2024 16:47:35 +0100 Subject: [PATCH 6/7] Check that `strings` can find the symbol Make sure that `strings` can successfully detect the symbol in an object file. --- sycl/test/check_device_code/device_global_static.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/sycl/test/check_device_code/device_global_static.cpp b/sycl/test/check_device_code/device_global_static.cpp index f97550a182704..ff7281452f3c7 100644 --- a/sycl/test/check_device_code/device_global_static.cpp +++ b/sycl/test/check_device_code/device_global_static.cpp @@ -2,15 +2,25 @@ // device globals, is removed at some point in compilation. For SPIR-V this // symbol is removed at sycl-post-link and for NVPTX/AMDGCN it is removed at // lowering. +// +// It also checks that the symbol can be found in an object file for a given +// triple, thus validating that `strings` can successfully be used to check for +// the presence of the symbol. // UNSUPPORTED: windows +// RUN: %clangxx -fsycl -fsycl-device-only %s -o %t +// RUN: strings %t | grep "llvm.compiler.used" // RUN: %clangxx -fsycl %s -o %t // RUN: strings %t | not grep "llvm.compiler.used" +// RUN: %if cuda %{ %clangxx -fsycl -fsycl-device-only -fsycl-targets=nvptx64-nvidia-cuda %s -o %t %} +// RUN: %if cuda %{ strings %t | grep "llvm.compiler.used" %} // RUN: %if cuda %{ %clangxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda %s -o %t %} // RUN: %if cuda %{ strings %t | not grep "llvm.compiler.used" %} +// RUN: %if hip_amd %{ %clangxx -fsycl -fsycl-device-only -fsycl-targets=amd_gpu_gfx906 %s -o %t %} +// RUN: %if hip_amd %{ strings %t | grep "llvm.compiler.used" %} // RUN: %if hip_amd %{ %clangxx -fsycl -fsycl-targets=amd_gpu_gfx906 %s -o %t %} // RUN: %if hip_amd %{ strings %t | not grep "llvm.compiler.used" %} From e852cbd45c288c5285eb1cd723dfec924066ab13 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Thu, 5 Sep 2024 17:31:22 +0100 Subject: [PATCH 7/7] Change strings usage to llvm-strings --- .../check_device_code/device_global_static.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/sycl/test/check_device_code/device_global_static.cpp b/sycl/test/check_device_code/device_global_static.cpp index ff7281452f3c7..a2d7768067de6 100644 --- a/sycl/test/check_device_code/device_global_static.cpp +++ b/sycl/test/check_device_code/device_global_static.cpp @@ -4,25 +4,25 @@ // lowering. // // It also checks that the symbol can be found in an object file for a given -// triple, thus validating that `strings` can successfully be used to check for -// the presence of the symbol. +// triple, thus validating that `llvm-strings` can successfully be used to +// check for the presence of the symbol. // UNSUPPORTED: windows // RUN: %clangxx -fsycl -fsycl-device-only %s -o %t -// RUN: strings %t | grep "llvm.compiler.used" +// RUN: llvm-strings %t | grep "llvm.compiler.used" // RUN: %clangxx -fsycl %s -o %t -// RUN: strings %t | not grep "llvm.compiler.used" +// RUN: llvm-strings %t | not grep "llvm.compiler.used" // RUN: %if cuda %{ %clangxx -fsycl -fsycl-device-only -fsycl-targets=nvptx64-nvidia-cuda %s -o %t %} -// RUN: %if cuda %{ strings %t | grep "llvm.compiler.used" %} +// RUN: %if cuda %{ llvm-strings %t | grep "llvm.compiler.used" %} // RUN: %if cuda %{ %clangxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda %s -o %t %} -// RUN: %if cuda %{ strings %t | not grep "llvm.compiler.used" %} +// RUN: %if cuda %{ llvm-strings %t | not grep "llvm.compiler.used" %} // RUN: %if hip_amd %{ %clangxx -fsycl -fsycl-device-only -fsycl-targets=amd_gpu_gfx906 %s -o %t %} -// RUN: %if hip_amd %{ strings %t | grep "llvm.compiler.used" %} +// RUN: %if hip_amd %{ llvm-strings %t | grep "llvm.compiler.used" %} // RUN: %if hip_amd %{ %clangxx -fsycl -fsycl-targets=amd_gpu_gfx906 %s -o %t %} -// RUN: %if hip_amd %{ strings %t | not grep "llvm.compiler.used" %} +// RUN: %if hip_amd %{ llvm-strings %t | not grep "llvm.compiler.used" %} #include