Dcompute fixes (#2195)

* Disable debug info for now Until I get it working.

* Don't recurse into uninstansiated templates

* Remove device only modules from objfiles.

* Remove typeid generation for dcompute.

* Dont generate typeinfo member functions

* Whitelist `__dcompute_reflect`

driver/main.cpp

@@ -1106,9 +1106,19 @@ void codegenModules(Modules &modules) {
       {
         cg.emit(m);
       }
-      if (atCompute != DComputeCompileFor::hostOnly)
+      if (atCompute != DComputeCompileFor::hostOnly) {
         computeModules.push_back(m);
-
+        if (atCompute == DComputeCompileFor::deviceOnly) {
+          // Remove m's object file from list of object files
+          auto s = m->objfile->name->str;
+          for (size_t j = 0; j < global.params.objfiles->dim; j++) {
+            if (s == (*global.params.objfiles)[j]) {
+              global.params.objfiles->remove(j);
+              break;
+            }
+          }
+        }
+      }
       if (global.errors)
         fatal();
     }
@@ -1119,6 +1129,10 @@ void codegenModules(Modules &modules) {
 
       dccg.writeModules();
     }
+    // We may have removed all object files, if so don't link.
+    if (global.params.objfiles->dim == 0)
+      global.params.link = false;
+
   }
 
   cache::pruneCache();

gen/dibuilder.cpp

@@ -78,11 +78,19 @@ llvm::StringRef uniqueIdent(Type* t) {
 
 bool ldc::DIBuilder::mustEmitFullDebugInfo() {
   // only for -g and -gc 
+  // TODO: but not dcompute (yet)
+
+  if (IR->dcomputetarget) return false;
+
   return global.params.symdebug == 1 || global.params.symdebug == 2;
 }
 
 bool ldc::DIBuilder::mustEmitLocationsDebugInfo() {
   // for -g -gc and -gline-tables-only 
+  // TODO:but not dcompute (yet)
+    
+  if (IR->dcomputetarget) return false;
+
   return (global.params.symdebug > 0) || global.params.outputSourceLocations;
 }
 

gen/functions.cpp

@@ -838,6 +838,17 @@ void DtoDefineFunction(FuncDeclaration *fd, bool linkageAvailableExternally) {
     return;
   }
 
+  if (gIR->dcomputetarget) {
+    auto id = fd->ident;
+    if (id == Id::xopEquals || id == Id::xopCmp || id == Id::xtoHash) {
+      IF_LOG Logger::println(
+          "No code generation for typeinfo member %s in @compute code",
+          fd->toChars());
+      fd->ir->setDefined();
+      return;
+    }
+  }
+
   if (!linkageAvailableExternally && !alreadyOrWillBeDefined(*fd)) {
     IF_LOG Logger::println("Skipping '%s'.", fd->toPrettyChars());
     fd->ir->setDefined();

gen/semantic-dcompute.cpp

@@ -38,13 +38,15 @@ struct DComputeSemanticAnalyser : public StoppableVisitor {
   // template declaration, it's module of origin is the module at the point of
   // instansiation so we need to check for that.
   bool isNonComputeCallExpVaild(CallExp *ce) {
+    FuncDeclaration *f = ce->f;
+    if (f->ident == Id::dcReflect)
+      return true;
     if (currentFunction == nullptr)
       return false;
     TemplateInstance *inst = currentFunction->isInstantiated();
     if (!inst)
       return false;
 
-    FuncDeclaration *f = ce->f;
     Objects *tiargs = inst->tiargs;
     size_t i = 0, len = tiargs->dim;
     IF_LOG Logger::println("checking against: %s (%p) (dyncast=%d)",
@@ -77,8 +79,12 @@ struct DComputeSemanticAnalyser : public StoppableVisitor {
 
   void visit(VarDeclaration *decl) override {
     // Don't print multiple errors for 'synchronized'. see visit(CallExp*)
-    if (decl->isDataseg() && strncmp(decl->toChars(), "__critsec", 9)) {
-      decl->error("global variables not allowed in @compute code");
+    if (decl->isDataseg()) {
+      if (strncmp(decl->toChars(), "__critsec", 9) &&
+        strncmp(decl->toChars(), "typeid", 6)) {
+        decl->error("global variables not allowed in @compute code variable=%s",decl->toChars());
+      }
+      // Ignore typeid: it is ignored by codegen.
       stop = true;
       return;
     }
@@ -213,7 +219,7 @@ struct DComputeSemanticAnalyser : public StoppableVisitor {
 
   void visit(FuncDeclaration *fd) override {
     if (hasKernelAttr(fd) && fd->vthis) {
-      fd->error("@kernel functions msut not require 'this'");
+      fd->error("@kernel functions must not require 'this'");
       stop = true;
       return;
     }
@@ -221,6 +227,11 @@ struct DComputeSemanticAnalyser : public StoppableVisitor {
     IF_LOG Logger::println("current function = %s", fd->toChars());
     currentFunction = fd;
   }
+    
+  void visit(TemplateDeclaration*) override {
+    // Don't try to analyse uninstansiated templates.
+    stop = true;
+  }
   // Override the default assert(0) behavior of Visitor:
   void visit(Statement *) override {}   // do nothing
   void visit(Expression *) override {}  // do nothing

runtime/druntime

@@ -1 +1 @@
-Subproject commit 1dd6a96b34f4a2222db636270b0e4269ce3a8f9b
+Subproject commit cc9335131218cf157af7e562309b64e5b0c51bb3

tests/compilable/dcompute.d

@@ -0,0 +1,43 @@
+// Tests that
+//  - we dont try to link with one file on the commandline that is @compute
+//  - truning on debugging doesn't ICE
+//  - dont analyse uninstansiated templates
+//  - typeid generated for hashing of struct (typeid(const(T))) is ignored and does not error
+
+// REQUIRES: target_NVPTX
+// RUN: %ldc -mdcompute-targets=cuda-350 -g
+
+@compute(Compilefor.deviceOnly) module dcompute;
+import ldc.dcompute;
+
+@kernel void foo()
+{
+
+}
+
+struct AutoIndexed(T)
+{
+    T p = void;
+    alias U = typeof(*T);
+
+    @property U index()
+    {
+        return p[0];
+    }
+
+    @property void index(U t)
+    {
+        p[0] = t;
+    }
+    @disable this();
+    alias index this;
+}
+alias aagf = AutoIndexed!(GlobalPointer!(float));
+
+@kernel void auto_index_test(aagf a,
+                             aagf b,
+                             aagf c)
+{
+    a = b + c;
+}
+