Merge branch 'DrTimothyAldenDavis:dev2' into dev2

Author: VidithM

CUDA/template/GB_jit_kernel_cuda_apply_unop.cu

@@ -39,16 +39,15 @@ __global__ void GB_cuda_apply_unop_kernel
 
     #define A_iso GB_A_ISO
 
-    int tid = blockDim.x * blockIdx.x + threadIdx.x ;
-    int nthreads = blockDim.x * gridDim.x ;
-
     #if ( GB_DEPENDS_ON_Y )
         // get thunk value (of type GB_Y_TYPE)
         GB_Y_TYPE thunk_value = * ((GB_Y_TYPE *) thunk) ;
     #endif
 
     #if ( GB_A_IS_BITMAP || GB_A_IS_FULL )
         // bitmap/full case
+        int tid = blockDim.x * blockIdx.x + threadIdx.x ;
+        int nthreads = blockDim.x * gridDim.x ;
         for (int64_t p = tid ; p < anz ; p += nthreads)
         {
             if (!GBb_A (Ab, p)) { continue ; }
@@ -92,8 +91,10 @@ __global__ void GB_cuda_apply_unop_kernel
                     }
                 }
         #else
-            const int64_t avlen = A->vlen ;
             // can do normal method
+            const int64_t avlen = A->vlen ;
+            int tid = blockDim.x * blockIdx.x + threadIdx.x ;
+            int nthreads = blockDim.x * gridDim.x ;
             for (int64_t p = tid ; p < anz ; p += nthreads)
             {
                 #if ( GB_DEPENDS_ON_I )

Source/jitifyer/GB_macrofy_binop.c

@@ -113,7 +113,7 @@ void GB_macrofy_binop
             ASSERT (!is_monoid_or_build) ;
             if (is_kron)
             { 
-                fprintf (fp, " %s (&(z), &(x),ix,jx, &(y),iy,jy, theta)\n",
+                fprintf (fp, " %s (&(z), ) &(x),ix,jx, &(y),iy,jy, theta)\n",
                     op->name) ;
             }
             else

Source/jitifyer/GB_macrofy_unop.c

@@ -43,8 +43,8 @@ void GB_macrofy_unop
 
         ASSERT (op != NULL) ;
         GB_macrofy_defn (fp, 3, op->name, op->defn) ;
-        fprintf (fp, "#define %s(z,x,%s,y)  %s (&(z), &(x))\n", macro_name,
-            ij, op->name) ;
+        fprintf (fp, "#define %s(z,x,%s,y)  %s (&(z), &(x))\n",
+            macro_name, ij, op->name) ;
 
     }
     else if (ecode == 254)