Don't remove llvm.compiler.used for NVPTX/AMDGCN

These are removed automatically during lowering.

Author: hdelan

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -788,7 +788,10 @@ processInputModule(std::unique_ptr<Module> 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.

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<int[4], TestProperties> 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<int, 1> OutBuf(&OutVal, 1);
+    Q.submit([&](handler &CGH) {
+      auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
+      CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar.get()[0]; });
+    });
+  }
+  assert(OutVal == 42 && "Read value does not match.");
+  return 0;
+}

sycl/test/check_device_code/device_global_static.cpp

@@ -0,0 +1,24 @@
+// 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" %}
+
+// Tests that the underlying pointer in a const-qualified shared device_global
+// is not optimized out during access.
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi::experimental;
+
+static device_global<int> DeviceGlobalVar;
+
+int main() {
+  sycl::queue{}.single_task([=] { volatile int ReadVal = DeviceGlobalVar; });
+}