   cudaMalloc((void **)&d_trace, d_trace_size);
   cudaMemcpy(d_trace, trace, d_trace_size, cudaMemcpyHostToDevice);
 
+  // also record the final state of A
+  int final[8];
+  for (int i=0; i<8; i++) {
+    final[i] = 99;
+  }
+  int *d_final;
+  size_t d_final_size = sizeof(int)*8;
+  cudaMalloc((void **)&d_final, d_final_size);
+  cudaMemcpy(d_final, final, d_final_size, cudaMemcpyHostToDevice);
+
   // run kernel
   printf("Set x and y through xyvals[%d,%d]...", xyvals[0], xyvals[1]);
   ASSERT_NO_CUDA_ERROR();
-  k2<<</*gridDim=*/1, GROUPSIZE>>>(d_xyvals, d_trace);
+  k2<<</*gridDim=*/1, GROUPSIZE>>>(d_xyvals, d_trace, d_final);
   ASSERT_NO_CUDA_ERROR();
   printf("[done]\n");
 

    printf("---\n");
  }

  cudaFree(d_xyvals);
  cudaFree(d_trace);
  delete[] trace;
  return 0;


      j++;
    }
    i++;
  }

  __syncthreads();
  if (lid == 0) {
    final[0] = A[0][0]; final[1] = A[0][1];
    final[2] = A[0][2]; final[3] = A[0][3];
    final[4] = A[1][0]; final[5] = A[1][1];
    final[6] = A[1][2]; final[7] = A[1][3];
  }
}

