[SYCL] Fix propagation of aspects from global variables (#11173)

Before this PR, the `pf-wg-atomic64.cpp` (from this PR) fails because
the `atomic64` aspect was not propagated to the calling kernel. The
reason is:

- With opaque pointers disabled, since the `this` parameter of the
constructor of `atomic_ref` would have a type of `atomic_ref*`, the call
to the constructor of `atomic_ref` would be enough to propagate the
`atomic64` aspect to the kernel. With opaque pointers enabled, this
would not be the case: the `this` parameter would just be a `ptr`.
- Additionally, for some reason I am unsure of, the code generated for
`parallel_for_work_group` (and not, say, for `parallel_for`) has the
pointer for `feature` generated as a global variable instead of
generated from an `alloca` instruction. If there were an `alloca` of an
`atomic_ref`, there would be no problem.

Now, when examining the operands of instructions in
`SYCLPropagateAspectsPass`, we now consider if the operand is a global
variable, and if so, propagate the type is points to instead of
propagating the pointer type.

Author: jzc

llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp

@@ -239,21 +239,24 @@ AspectsSetTy getAspectsUsedByInstruction(const Instruction &I,
     ReturnType = AI->getAllocatedType();
   }
   AspectsSetTy Result = getAspectsFromType(ReturnType, Types);
-  for (const auto &OperandIt : I.operands()) {
-    const AspectsSetTy &Aspects =
-        getAspectsFromType(OperandIt->getType(), Types);
+  auto AddAspectsFromType = [&](Type *Ty) {
+    const AspectsSetTy &Aspects = getAspectsFromType(Ty, Types);
     Result.insert(Aspects.begin(), Aspects.end());
+  };
+  for (const auto &OperandIt : I.operands()) {
+    if (const auto *GV =
+            dyn_cast<const GlobalValue>(OperandIt->stripPointerCasts()))
+      AddAspectsFromType(GV->getValueType());
+    else
+      AddAspectsFromType(OperandIt->getType());
   }
 
   // Opaque pointer arguments may hide types of pointer arguments until elements
   // inside the types are accessed through a GEP instruction. However, this will
   // not be caught by the operands check above, so we must extract the
   // information directly from the GEP.
-  if (auto *GEPI = dyn_cast<const GetElementPtrInst>(&I)) {
-    const AspectsSetTy &Aspects =
-        getAspectsFromType(GEPI->getSourceElementType(), Types);
-    Result.insert(Aspects.begin(), Aspects.end());
-  }
+  if (auto *GEPI = dyn_cast<const GetElementPtrInst>(&I))
+    AddAspectsFromType(GEPI->getSourceElementType());
 
   if (const MDNode *InstApsects = I.getMetadata("sycl_used_aspects")) {
     for (const MDOperand &MDOp : InstApsects->operands()) {

llvm/test/SYCLLowerIR/PropagateAspectsUsage/global-value.ll

@@ -0,0 +1,27 @@
+; This tests ensures that the aspects of a global variable are propagated.
+; RUN: opt -passes=sycl-propagate-aspects-usage < %s -S | FileCheck %s
+
+%struct.StructWithAspect = type { i32 }
+@global = internal addrspace(3) global %struct.StructWithAspect undef, align 8
+
+declare void @external(ptr addrspace(4))
+
+; CHECK: spir_kernel void @foo() !sycl_used_aspects ![[MDID:[0-9]+]]
+define spir_kernel void @foo() {
+  %res = load ptr addrspace(3), ptr addrspace(3) @global
+  ret void
+}
+
+; CHECK: spir_kernel void @bar() !sycl_used_aspects ![[MDID]]
+define spir_kernel void @bar() {
+  call void @external(ptr addrspace(4) addrspacecast(ptr addrspace(3) @global to ptr addrspace(4)))
+  ret void
+}
+
+!sycl_types_that_use_aspects = !{!1}
+!sycl_aspects = !{!2}
+
+!1 = !{!"struct.StructWithAspect", i32 6}
+!2 = !{!"fp64", i32 6}
+
+; CHECK: ![[MDID]] = !{i32 6}

sycl/test-e2e/Regression/pf-wg-atomic64.cpp

@@ -0,0 +1,27 @@
+// DISABLED: aspect-atomic64
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/sycl.hpp>
+using namespace sycl;
+
+using AtomicRefT =
+    atomic_ref<unsigned long long, memory_order::relaxed, memory_scope::device>;
+
+int main() {
+  queue q;
+  auto *p = malloc_shared<unsigned long long>(1, q);
+  try {
+    q.submit([&](sycl::handler &cgh) {
+       cgh.parallel_for_work_group(range{1}, range{1}, [=](group<1>) {
+         AtomicRefT feature(*p);
+         feature += 42;
+       });
+     }).wait();
+  } catch (sycl::exception &e) {
+    if (e.code() != sycl::errc::kernel_not_supported)
+      throw;
+    std::cout << "Caught right exception: " << e.what() << "\n";
+    return 0;
+  }
+}