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/5] 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/5] 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;
+}

From 65c59b3e7a25f94657597ee76d3141e80644b469 Mon Sep 17 00:00:00 2001
From: gulugulubing <413153391@qq.com>
Date: Sun, 24 May 2026 11:04:37 -0600
Subject: [PATCH 3/5] clean the logic

---
 gen/semantic-dcompute.cpp    | 32 +++++++-------------------------
 tests/compilable/issue5116.d | 22 +++++++++++++++-------
 2 files changed, 22 insertions(+), 32 deletions(-)

diff --git a/gen/semantic-dcompute.cpp b/gen/semantic-dcompute.cpp
index 556a699ad2..92d1e7d43c 100644
--- a/gen/semantic-dcompute.cpp
+++ b/gen/semantic-dcompute.cpp
@@ -236,7 +236,13 @@ struct DComputeSemanticAnalyser : public StoppableVisitor {
     }
       
     Module *m = e->f->getModule();
-    if ((m == nullptr || (hasComputeAttr(m) == DComputeCompileFor::hostOnly)) &&
+    // Template-instantiated functions are cross-module by nature: the template
+    // declaration and the instantiated function live in different modules.
+    // getModule() returns the *declaration* module, which says nothing about
+    // whether the generated code can run on GPU. Skip the module check for them.
+    const bool isTemplateFunc = e->f->isInstantiated() != nullptr;
+    if (!isTemplateFunc &&
+        (m == nullptr || (hasComputeAttr(m) == DComputeCompileFor::hostOnly)) &&
         !isNonComputeCallExpVaild(e)) {
       error(e->loc, "can only call functions from other `@compute` modules in "
                     "`@compute` code");
@@ -251,16 +257,6 @@ 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;
   }
@@ -275,20 +271,6 @@ 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;
-      }
     }
   }
 
diff --git a/tests/compilable/issue5116.d b/tests/compilable/issue5116.d
index 6709ea7f2b..a2964b52ec 100644
--- a/tests/compilable/issue5116.d
+++ b/tests/compilable/issue5116.d
@@ -1,10 +1,8 @@
-// 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).
+// Issue #5116: defining a struct with a static array field in a @compute
+// module caused DMD to generate __xopEquals, dragging in template
+// instantiations from core.internal.array.equality. The semantic walker
+// then either reported spurious errors on the nested __equals calls or
+// crashed on a null function pointer inside the instantiated body.
 
 // REQUIRES: target_NVPTX
 // RUN: %ldc -mdcompute-targets=cuda-350 %s
@@ -17,3 +15,13 @@ private enum N = 16u;
 struct S {
     float[N] data;
 }
+
+@kernel void testEqualExp() {
+    float[N] a, b;
+    bool c = (a == b);
+}
+
+@kernel void testExplicitEquals() {
+    float[N] a, b;
+    bool c = __equals(a, b);
+}

From c63d7691844adb60486528e1b73a4b88a393f26b Mon Sep 17 00:00:00 2001
From: gulugulubing <413153391@qq.com>
Date: Sun, 24 May 2026 11:16:00 -0600
Subject: [PATCH 4/5] remove trailing blank

---
 gen/semantic-dcompute.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/gen/semantic-dcompute.cpp b/gen/semantic-dcompute.cpp
index 92d1e7d43c..acc67af1c3 100644
--- a/gen/semantic-dcompute.cpp
+++ b/gen/semantic-dcompute.cpp
@@ -234,7 +234,7 @@ struct DComputeSemanticAnalyser : public StoppableVisitor {
       stop = true;
       return;
     }
-      
+
     Module *m = e->f->getModule();
     // Template-instantiated functions are cross-module by nature: the template
     // declaration and the instantiated function live in different modules.

From 670b1a41348b4804ed79b5870a6b17350b70b4fa Mon Sep 17 00:00:00 2001
From: gulugulubing <413153391@qq.com>
Date: Sun, 24 May 2026 11:24:31 -0600
Subject: [PATCH 5/5] tweak comment in test

---
 tests/compilable/issue5116.d | 9 ++++-----
 1 file changed, 4 insertions(+), 5 deletions(-)

diff --git a/tests/compilable/issue5116.d b/tests/compilable/issue5116.d
index a2964b52ec..f59b57a245 100644
--- a/tests/compilable/issue5116.d
+++ b/tests/compilable/issue5116.d
@@ -1,8 +1,7 @@
-// Issue #5116: defining a struct with a static array field in a @compute
-// module caused DMD to generate __xopEquals, dragging in template
-// instantiations from core.internal.array.equality. The semantic walker
-// then either reported spurious errors on the nested __equals calls or
-// crashed on a null function pointer inside the instantiated body.
+// Without this patch, template instantiations from non-@compute modules (e.g.
+// __equals from core.internal.array.equality) were walked, producing spurious errors.
+// And indirect calls via function pointers inside those bodies additionally caused a
+// null-pointer dereference crash.
 
 // REQUIRES: target_NVPTX
 // RUN: %ldc -mdcompute-targets=cuda-350 %s