Fix #5116 - Skip the sematics checks of compiler-generated functions and template instances in dcompute

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;
+      }
     }
   }
 

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;
+}