       break;
 
     case  1:
-      cudaFuncSetCacheConfig (k_jacobianReduceHessianGradient<1,true>, cudaFuncCachePreferShared);
+      cudaFuncSetCacheConfig (k_jacobianReduceHessianGradient<1,true>, cudaFuncCachePreferL1);
       k_jacobianReduceHessianGradient<1, true><<< dimGrid, dimBlock >>>(jacobian_input_device,
                                                                         residual_input_device,
                                                                         visibility_input_device,

      break;

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

      break;

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