       reference_on_device, reference_on_host,
       num_reference_bytes, cudaMemcpyHostToDevice);
 
-    int num_threads_per_block = 512;
+    int num_threads_per_block = 2;
     int num_blocks = (num_query_points + num_threads_per_block - 1) /
                      num_threads_per_block;
 

__device__ void LoadReferencePoint_(
  int num_dimensions,
  float *reference,
  int reference_point_id) {

  extern __shared__ float reference_point_shared_mem[];

  int i;
  int source_pos = reference_point_id * num_dimensions;

    if(global_thread_id < num_query_points) {
      AccumulateTileContribution_(
        num_dimensions, bandwidth, query_point_local_mem,
        num_reference_points_in_this_tile,
        reference_point_shared_mem, & local_sum);
    }

    // Synchronize all threads within this block before loading new
