 }
 
 __device__ void copy_data(const PMatrix* data, float* sh_data,
-						  int thidx, int thidy, int obs_num)
+                          int thidx, int thidy, int obs_num)
 {
   if (obs_num >= data->rows)
-	return;
+    return;
 
   for (int chunk = 0; chunk < data->cols; chunk += blockDim.y)
   {

  PMatrix pdata, pparams;
  CATCH_ERR(cudaMalloc((void**) &d_pdf, total_obs * nparams * sizeof(float)));
  CATCH_ERR(cudaMalloc((void**) &d_data,
					   data_stride * total_obs * sizeof(float)));
  CATCH_ERR(cudaMalloc((void**) &d_params,
					   param_stride * nparams * sizeof(float)));

  h_to_d(h_data, d_data, total_obs * data_stride);
  h_to_d(h_params, d_params, nparams * param_stride);

  // allocated enough shared memory so that this will not walk out of bounds
  // no matter what
  sh_result[sh_idx] = compute_pdf(sh_data + thidx * data.cols,
                                  sh_params + thidy * params.stride,
                                  data.cols);
  __syncthreads();

  // does this coalesce?
