[cig-commits] r15463 - mc/3D/CitcomS/trunk/lib

[leif at geodynamics.org]

Author: leif
Date: 2009-07-15 17:49:42 -0700 (Wed, 15 Jul 2009)
New Revision: 15463

Modified:
   mc/3D/CitcomS/trunk/lib/multigrid_kernel.cu
Log:
Bug fixes... still doesn't work right.


Modified: mc/3D/CitcomS/trunk/lib/multigrid_kernel.cu
===================================================================
--- mc/3D/CitcomS/trunk/lib/multigrid_kernel.cu	2009-07-15 21:54:24 UTC (rev 15462)
+++ mc/3D/CitcomS/trunk/lib/multigrid_kernel.cu	2009-07-16 00:49:42 UTC (rev 15463)
@@ -355,7 +355,7 @@
     
     if (tid == 0) {
         /* Each block writes one element of Ad... */
-        Ad[eqn] = sum[0];
+        Ad[eqn] += sum[0];
         /* ..and one element of d0. */
         d0[eqn] += E->temp[eqn];
     }
@@ -425,10 +425,13 @@
     }
     
     /* wait for completion */
-    cudaThreadSynchronize();
+    if (cudaThreadSynchronize() != cudaSuccess) {
+        assert(0 && "something went wrong");
+    }
     
     /* copy output from device */
     cudaMemcpy(Ad, d_Ad, (1+neq)*sizeof(double), cudaMemcpyDeviceToHost);
+    cudaMemcpy(d0, d_d0, (1+neq)*sizeof(double), cudaMemcpyDeviceToHost);
     
     /* free device memory */
     cudaFree(d_d0);
@@ -462,12 +465,6 @@
 {
     /* Map out how to parallelize "Au[C[i]] += ..." and "Ad[C[j]] += ...". */
     
-    static int2 *cache;
-    if (cache) {
-        E->term = cache;
-        return;
-    }
-    
     const int neq = E->lmesh.NEQ;
     const int nno = E->lmesh.NNO;
     
@@ -497,8 +494,6 @@
         }
     }
     
-    cache = E->term;
-    
     return;
 }
 
@@ -537,8 +532,6 @@
     
     collect_terms(&kE);
     
-    assert(0);
-    
     do_gauss_seidel(
         &kE,
         d0[M],