[SYCL][ESIMD] Don't error on functions with local names (#7996)

In the ESIMD Verifier, we have a known good list of functions and want
to check that any ESIMD code only calls those functions. We do that by
demangling the names of all ESIMD functions and comparing them against
the known good functions. In the case where there's a function that is
from a local name such as a lambda, functor, or local class, the
demangled name may include the return type of the parent function if
it's a template function. If this happens, we end up using the return
type in the function name check which makes us check the wrong thing.

We don't care about local names, so just skip them.

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -118,6 +118,14 @@ class ESIMDVerifierImpl {
           if (!NameNode) // Can it be null?
             continue;
 
+          // Skip local names, which are functions whose type
+          // is not exposed outside of the current function,
+          // such as lambdas or local classes. Note we will
+          // still analyze functions called by these constructs,
+          // assuming they are marked as ESIMD functions.
+          if (NameNode->getKind() == id::Node::KLocalName)
+            continue;
+
           id::OutputBuffer NameBuf;
           NameNode->print(NameBuf);
           StringRef Name(NameBuf.getBuffer(), NameBuf.getCurrentPosition());

sycl/test/esimd/esimd_verify_local_names.cpp

@@ -0,0 +1,45 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -S %s -o /dev/null
+// Test that the ESIMD Verifier doesn't error on locally defined types
+#include <sycl.hpp>
+#include <sycl/ext/intel/esimd.hpp>
+
+using namespace sycl::ext::intel::esimd;
+
+constexpr size_t n = 64;
+
+constexpr size_t size = 4;
+
+class functor {
+public:
+  template <typename T> sycl::event operator()(sycl::queue &q, T *x) {
+    return q.submit([&](sycl::handler &cgh) {
+      cgh.parallel_for(n / size, [=](auto item) SYCL_ESIMD_KERNEL {
+        size_t offset = item.get_id(0);
+        simd<T, size> vec(offset);
+        block_store(x + offset * size, vec);
+      });
+    });
+  }
+};
+
+template <typename T> sycl::event func(sycl::queue &q, T *x) {
+  return q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for(n / size, [=](auto item) SYCL_ESIMD_KERNEL {
+      size_t offset = item.get_id(0);
+      simd<T, size> vec(offset);
+      block_store(x + offset * size, vec);
+    });
+  });
+}
+
+int main() {
+  sycl::queue q;
+  float *x = sycl::malloc_shared<float>(n, q);
+  auto event = func(q, x);
+  event.wait();
+  functor f;
+  auto event2 = f(q, x);
+  event2.wait();
+  sycl::free(x, q);
+  return 0;
+}