From 91994af37676de916c1f6221eeb27409b25b77ff Mon Sep 17 00:00:00 2001
From: gulugulubing <413153391@qq.com>
Date: Mon, 18 May 2026 16:28:25 -0600
Subject: [PATCH 1/2] skip compiler-generated functions and template instances

---
 gen/semantic-dcompute.cpp | 29 +++++++++++++++++++++++++++++
 1 file changed, 29 insertions(+)

diff --git a/gen/semantic-dcompute.cpp b/gen/semantic-dcompute.cpp
index 1cf3fb4843..556a699ad2 100644
--- a/gen/semantic-dcompute.cpp
+++ b/gen/semantic-dcompute.cpp
@@ -212,6 +212,11 @@ struct DComputeSemanticAnalyser : public StoppableVisitor {
     }
   }
   void visit(CallExp *e) override {
+    // Indirect calls via function pointers / delegates have no associated
+    // FuncDeclaration, so there is no module to check.
+    if (!e->f)
+      return;
+
     // SynchronizedStatement is lowered to
     //    Critsec __critsec105; // 105 == line number
     //    _d_criticalenter(& __critsec105); <--
@@ -246,6 +251,16 @@ struct DComputeSemanticAnalyser : public StoppableVisitor {
       return;
     }
 
+    // Skip compiler-generated struct support functions (e.g. __xopEquals,
+    // __xopCmp, postblit, destructor). Their bodies may reference non-@compute
+    // templates (such as __equals for static array comparison) that are outside
+    // of user control. They are only codegenerated if actually referenced by
+    // user code, at which point the codegen layer will report any issues.
+    if (fd->isGenerated()) {
+      stop = true;
+      return;
+    }
+
     IF_LOG Logger::println("current function = %s", fd->toChars());
     currentFunction = fd;
   }
@@ -260,6 +275,20 @@ struct DComputeSemanticAnalyser : public StoppableVisitor {
     // as they contain unsupported global variables.
     if (ti->tempdecl == Type::rtinfo || ti->tempdecl == Type::rtinfoImpl) {
       stop = true;
+      return;
+    }
+
+    // Template instantiations for templates declared in non-@compute modules
+    // (e.g. __equals and isEqual from core.internal.array.equality) are
+    // created as a side effect of compiler-generated support functions. They
+    // contain calls back into their declaring (non-@compute) module, which
+    // would produce spurious errors. Skip them.
+    if (ti->tempdecl) {
+      Module *m = ti->tempdecl->getModule();
+      if (m && hasComputeAttr(m) == DComputeCompileFor::hostOnly) {
+        stop = true;
+        return;
+      }
     }
   }
 

From 4c478b19b474ea1205939c4182fbdfe243aebb1a Mon Sep 17 00:00:00 2001
From: gulugulubing <413153391@qq.com>
Date: Mon, 18 May 2026 16:46:23 -0600
Subject: [PATCH 2/2] add a test

---
 tests/compilable/issue5116.d | 19 +++++++++++++++++++
 1 file changed, 19 insertions(+)
 create mode 100644 tests/compilable/issue5116.d

diff --git a/tests/compilable/issue5116.d b/tests/compilable/issue5116.d
new file mode 100644
index 0000000000..6709ea7f2b
--- /dev/null
+++ b/tests/compilable/issue5116.d
@@ -0,0 +1,19 @@
+// Regression test for issue #5116: defining a struct with a static array
+// field in a @compute module previously caused a spurious semantic error
+// ("can only call functions from other `@compute` modules") followed by a
+// null-pointer dereference crash in DComputeSemanticAnalyser::visit(CallExp*).
+// The crash happened because compiler-generated support functions (__xopEquals)
+// triggered instantiation of __equals templates from core.internal.array.equality,
+// whose body contains an indirect call through a function pointer (e->f == null).
+
+// REQUIRES: target_NVPTX
+// RUN: %ldc -mdcompute-targets=cuda-350 %s
+
+@compute(CompileFor.deviceOnly) module tests.compilable.issue5116;
+import ldc.dcompute;
+
+private enum N = 16u;
+
+struct S {
+    float[N] data;
+}