 }
 
 void extractAll(int sharedMemoryOffset, int numBlocks, int numThreadsPerBlock, Aggregator* sumOverAll) {
-  Aggregator *sumOverBlock = NULL;
-  cudaMalloc((void**)&sumOverBlock, numBlocks * sizeof(Aggregator));
+  printf("eight\n");
+
+  Aggregator* sumOverBlock = NULL;
+  cudaErrorCheck(cudaMalloc((void**)&sumOverBlock, numBlocks * sizeof(Aggregator)));
+
+  printf("nine\n");
+
   extractFromBlock<<<numBlocks, 1, sharedMemoryOffset + numThreadsPerBlock * sizeof(Aggregator)>>>(sharedMemoryOffset, numThreadsPerBlock, sumOverBlock);
+  cudaErrorCheck(cudaPeekAtLastError());
+  cudaErrorCheck(cudaDeviceSynchronize());
 
-  Aggregator *sumOverBlock2 = malloc(numBlocks * sizeof(Aggregator));
-  cudaMemcpy(&sumOverBlock2, sumOverBlock, numBlocks * sizeof(Aggregator), cudaMemcpyDeviceToHost);
-  cudaFree(sumOverBlock);
+  printf("ten\n");
+
+  Aggregator* sumOverBlock2 = (Aggregator*)malloc(numBlocks * sizeof(Aggregator));
+
+  printf("eleven\n");
+
+  cudaErrorCheck(cudaMemcpy(sumOverBlock2, sumOverBlock, numBlocks * sizeof(Aggregator), cudaMemcpyDeviceToHost));
+
+  printf("twelve\n");
+
+  cudaErrorCheck(cudaFree(sumOverBlock));
+
+  printf("thirteen\n");
 
   zero(sumOverAll);
+
+  printf("fourteen %d\n", numBlocks);
+
   for (int block = 0;  block < numBlocks;  ++block)
     combine(sumOverAll, &sumOverBlock2[block]);
 
+  printf("fifteen\n");
+
   free(sumOverBlock2);
+
+  printf("sixteen\n");
 }
 
 int main(int argc, char** argv) {
+  int numBlocks = 1;
+  int numThreadsPerBlock = 1;
+  int sharedMemoryOffset = 0;
 
+  printf("one\n");
 
+  initialize<<<numBlocks, numThreadsPerBlock, numThreadsPerBlock * sizeof(Aggregator)>>>(sharedMemoryOffset);
+  cudaErrorCheck(cudaPeekAtLastError());
+  cudaErrorCheck(cudaDeviceSynchronize());
+
+  printf("two\n");
+
+  float x_cpu[10] = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f, 10.0f};
+  float y_cpu[10] = {10.0f, 9.0f, 8.0f, 7.0f, 6.0f, 5.0f, 4.0f, 3.0f, 2.0f, 1.0f};
+  float weight_cpu[10] = {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f};
+
+  float* x_gpu;
+  float* y_gpu;
+  float* weight_gpu;
+
+  printf("three\n");
+
+  cudaErrorCheck(cudaMalloc((void**)&x_gpu, 10 * sizeof(float)));
+  cudaErrorCheck(cudaMalloc((void**)&y_gpu, 10 * sizeof(float)));
+  cudaErrorCheck(cudaMalloc((void**)&weight_gpu, 10 * sizeof(float)));
+
+  printf("four\n");
+
+  cudaErrorCheck(cudaMemcpy(x_gpu, x_cpu, 10 * sizeof(float), cudaMemcpyHostToDevice));
+  cudaErrorCheck(cudaMemcpy(y_gpu, y_cpu, 10 * sizeof(float), cudaMemcpyHostToDevice));
+  cudaErrorCheck(cudaMemcpy(weight_gpu, weight_cpu, 10 * sizeof(float), cudaMemcpyHostToDevice));
+
+  printf("five\n");
+
+  fillAll<<<numBlocks, numThreadsPerBlock, numThreadsPerBlock * sizeof(Aggregator)>>>(sharedMemoryOffset, x_gpu, y_gpu, weight_gpu);
+  cudaErrorCheck(cudaPeekAtLastError());
+  cudaErrorCheck(cudaDeviceSynchronize());
+
+  printf("six\n");
+
+  cudaErrorCheck(cudaFree(x_gpu));
+  cudaErrorCheck(cudaFree(y_gpu));
+  cudaErrorCheck(cudaFree(weight_gpu));
+
+  printf("seven\n");
+
+  Aggregator sumOverAll;
+  extractAll(sharedMemoryOffset, numBlocks, numThreadsPerBlock, &sumOverAll);
+
+  printf("done %g\n", sumOverAll.test);
 }

}

extern __shared__ unsigned char sharedMemory[];

__global__ void initialize(int sharedMemoryOffset) {
  Aggregator* threadLocal = (Aggregator*)((size_t)sharedMemory + sharedMemoryOffset + threadId()*sizeof(Aggregator));


// the input fields are auto-generated
__global__ void fillAll(int sharedMemoryOffset, float* input_x, float* input_y, float* weight) {
  printf("here %g %g %g\n", input_x[0], input_y[0], weight[0]);

  int id = threadId();
  fill(sharedMemoryOffset, input_x[id], input_y[id], weight[id]);
}

__global__ void extractFromBlock(int sharedMemoryOffset, int numThreadsPerBlock, Aggregator* sumOverBlock) {
  Aggregator *blockLocal = &sumOverBlock[blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z];
  zero(sumOverBlock);
  for (int thread = 0;  thread < numThreadsPerBlock;  ++thread) {
    Aggregator* singleAggregator = (Aggregator*)((size_t)sharedMemory + sharedMemoryOffset + thread*sizeof(Aggregator));
