   {
     switch (threads)
     {
+
     case 512:
-      cudaFuncSetCacheConfig (k_jacobianReduceHessianGradient<512,true>, cudaFuncCachePreferShared);
-      k_jacobianReduceHessianGradient<512, true><<< dimGrid, dimBlock >>>(jacobian_input_device,
-                                                                          residual_input_device,
-                                                                          visibility_input_device,
-                                                                          gradient_output,
-                                                                          hessian_output,
-                                                                          size);
+      throw std::runtime_error(" 512 threads exceed the 48kB of available shared memory per block!");
+      //      cudaFuncSetCacheConfig (k_jacobianReduceHessianGradient<512,true>, cudaFuncCachePreferShared);
+      //      k_jacobianReduceHessianGradient<512, true><<< dimGrid, dimBlock >>>(jacobian_input_device,
+      //                                                                          residual_input_device,
+      //                                                                          visibility_input_device,
+      //                                                                          gradient_output,
+      //                                                                          hessian_output,
+      //                                                                          size);
       break;
 
     case 256:
-      cudaFuncSetCacheConfig (k_jacobianReduceHessianGradient<256,true>, cudaFuncCachePreferShared);
+      cudaFuncSetCacheConfig (k_jacobianReduceHessianGradient<256,true>, cudaFuncCachePreferL1);
       k_jacobianReduceHessianGradient<256, true><<< dimGrid, dimBlock >>>(jacobian_input_device,
                                                                           residual_input_device,
                                                                           visibility_input_device,

      break;

    case 128:
      cudaFuncSetCacheConfig (k_jacobianReduceHessianGradient<128,true>, cudaFuncCachePreferShared);
      k_jacobianReduceHessianGradient<128, true><<< dimGrid, dimBlock >>>(jacobian_input_device,
                                                                          residual_input_device,
                                                                          visibility_input_device,

      break;

    case 128:
      cudaFuncSetCacheConfig (k_jacobianReduceHessianGradient<128,false>, cudaFuncCachePreferL1);
      k_jacobianReduceHessianGradient<128, false><<< dimGrid, dimBlock >>>(jacobian_input_device,
                                                                           residual_input_device,
                                                                           visibility_input_device,
