         else
         {
             kernelCheckCache << <dimGridCache, dimBlockCache, kernelCheckCacheSMSize >> >(d_workingset + 1, d_K, d_KCacheRemapIdx, d_KCacheRowIdx, cache_rows, d_x, d_xT, gamma, num_vec, num_vec_aligned, dim, dim_aligned, cacheLastPtrIdx);
-            kernelCheckCacheFinalize<<<1, 1>>>(d_KCacheRemapIdx);
             cacheLastPtrIdx = 1 - cacheLastPtrIdx;
         }
+        kernelCheckCacheFinalize<<<1, 1>>>(d_KCacheRemapIdx);
         //workaround if caching J deleted I out of cache when not using priority cache
         if (!usePriorityCache)
         {

    {
        //KCacheRowIdx[cache_rows + (1 - lastPtrIdx)] = last;
        int del_i = KCacheRowIdx[last];
        if (del_i >= 0)
            KCacheRemapIdx[del_i] = -1;  //cache row for vector [del_i] will be overwritten, remove it from RemapIdx array
        //set correct indices
        KCacheRemapIdx[i] = last;
        KCacheRowIdx[last] = i;
        KCacheRowPriority[last] = ++d_cacheUpdateCnt;
    }

        else
        {
            kernelCheckCache << <dimGridCache, dimBlockCache, kernelCheckCacheSMSize >> >(d_workingset, d_K, d_KCacheRemapIdx, d_KCacheRowIdx, cache_rows, d_x, d_xT, gamma, num_vec, num_vec_aligned, dim, dim_aligned, cacheLastPtrIdx);
            cacheLastPtrIdx = 1 - cacheLastPtrIdx;
        }
        kernelCheckCacheFinalize<<<1, 1>>>(d_KCacheRemapIdx);

        //int * KCacheRowIdx = new int[cache_rows + 2];
        //assert_cuda(cudaMemcpy(KCacheRowIdx, d_KCacheRowIdx, (cache_rows + 2) * sizeof(int), cudaMemcpyDeviceToHost));
