   // const int data_total = data.rows * data.cols;
 
   // for (int chunk = data_start;
-  // 	   chunk < data_start + blockDim.x * data.stride;
-  // 	   chunk += num_threads)
+  //       chunk < data_start + blockDim.x * data.stride;
+  //       chunk += num_threads)
   // {
-  // 	idx = chunk + tid;
-  // 	if (idx < data_total)
-  // 	  sh_data[idx - data_start] = data.buf[idx];
+  //    idx = chunk + tid;
+  //    if (idx < data_total)
+  //      sh_data[idx - data_start] = data.buf[idx];
   // }
 
   // const int params_start = blockDim.y * blockIdx.y * params.stride;
   // const int params_total = params.rows * params.stride;
   // for (int chunk = params_start;
-  // 	   chunk < params_start + blockDim.y * params.stride;
-  // 	   chunk += num_threads)
+  //       chunk < params_start + blockDim.y * params.stride;
+  //       chunk += num_threads)
   // {
-  // 	idx = chunk + tid;
-  // 	if (idx < params_total)
-  // 	  sh_params[idx - params_start] = params.buf[idx];
+  //    idx = chunk + tid;
+  //    if (idx < params_total)
+  //      sh_params[idx - params_start] = params.buf[idx];
   // }
 
   __syncthreads();
 
   int sh_idx = thidy * blockDim.x + thidx;
   if (obs_num < data.rows & param_index < params.rows) {
-  	float d = compute_pdf(sh_data + thidx * data.cols,
-  						  sh_params + thidy * params.stride,
-  						  data.cols);
-  	sh_result[sh_idx] = d;
+    float d = compute_pdf(sh_data + thidx * data.cols,
+                          sh_params + thidy * params.stride,
+                          data.cols);
+    sh_result[sh_idx] = d;
   }
   __syncthreads();
 
   int result_idx = params.rows * obs_num + param_index;
   if (obs_num < data.rows & param_index < params.rows) {
-  	output[result_idx] = sh_result[sh_idx];
+    output[result_idx] = sh_result[sh_idx];
   }
 
   // // write out in other order to coalesce

  dim3 gridPDF(grid_x, grid_y);

  dim3 blockPDF(tune_info.data_per_block,
				tune_info.params_per_block);

  int sharedMemSize = compute_shmem(&data, &params,
									tune_info.params_per_block,
									tune_info.data_per_block);

  printf("number params: %d, number data points: %d\n",
		 tune_info.params_per_block, tune_info.data_per_block);
  printf("sharedMemSize: %d\n", sharedMemSize);
  printf("block: %d x %d, grid: %d x %d\n", blockPDF.x, blockPDF.y,
		 gridPDF.x, gridPDF.y);
  printf("nparams: %d\n", params.rows);

  mvnpdf_k<<<gridPDF,blockPDF,sharedMemSize>>>(data, params, d_pdf);

            float discrim;
            float* tData = x + obs * DATA_PADDED_DIM;
            float* tDensityInfo = density + component * PACK_DIM;
            float* tMean = tDensityInfo;            //do we need to unallocate shared/register variables?
            float* tSigma = tDensityInfo + D;
            float  tP = tDensityInfo[LOGDET_OFFSET];
            float  tLogDet = tDensityInfo[LOGDET_OFFSET+1];
