         reduceMaxIdx(d_reduceval, d_reduceidx, d_reduceval2, d_reduceidx2, num_vec_aligned, reduce_block_size);
         assert_cuda(cudaMemcpy(d_workingset, d_reduceidx, sizeof(int), cudaMemcpyDeviceToDevice));
 
-        //int * KCacheVecIdx = new int[cache_rows + 2];
-        //assert_cuda(cudaMemcpy(KCacheVecIdx, d_KCacheVecIdx, (cache_rows + 2) * sizeof(int), cudaMemcpyDeviceToHost));
-        //std::cout << "KCacheVecIdx: ";
+        //int * KCacheRowIdx = new int[cache_rows + 2];
+        //assert_cuda(cudaMemcpy(KCacheRowIdx, d_KCacheRowIdx, (cache_rows + 2) * sizeof(int), cudaMemcpyDeviceToHost));
+        //std::cout << "KCacheRowIdx: ";
         //for (int k = 0; k < cache_rows + 2; k++)
-        //    std::cout << KCacheVecIdx[k] << ", ";
+        //    std::cout << KCacheRowIdx[k] << ", ";
         //std::cout << std::endl;
-        //delete[] KCacheVecIdx;
+        //delete[] KCacheRowIdx;
 
         //check if I is cached
-        kernelCheckCache<<<dimGridCache, dimBlockCache, kernelCheckCacheSMSize>>>(d_workingset, d_K, d_KCacheRemapIdx, d_KCacheVecIdx, cache_rows, d_x, d_xT, gamma, num_vec, num_vec_aligned, dim, dim_aligned, cacheLastPtrIdx);
-        cacheLastPtrIdx = (cacheLastPtrIdx + 1) % 2;
-
-        //KCacheVecIdx = new int[cache_rows + 2];
-        //assert_cuda(cudaMemcpy(KCacheVecIdx, d_KCacheVecIdx, (cache_rows + 2) * sizeof(int), cudaMemcpyDeviceToHost));
-        //std::cout << "KCacheVecIdx: ";
+        kernelCheckCache<<<dimGridCache, dimBlockCache, kernelCheckCacheSMSize>>>(d_workingset, d_K, d_KCacheRemapIdx, d_KCacheRowIdx, d_KCacheRowPriority, cache_rows, d_x, d_xT, gamma, num_vec, num_vec_aligned, dim, dim_aligned, cacheLastPtrIdx);
+        cacheLastPtrIdx = 1 - cacheLastPtrIdx;
+
+        //int * KCacheRowIdx = new int[cache_rows + 2];
+        //assert_cuda(cudaMemcpy(KCacheRowIdx, d_KCacheRowIdx, (cache_rows + 2) * sizeof(int), cudaMemcpyDeviceToHost));
+        //std::cout << "KCacheRowIdx: ";
         //for (int k = 0; k < cache_rows + 2; k++)
-        //    std::cout << KCacheVecIdx[k] << ", ";
+        //    std::cout << KCacheRowIdx[k] << ", ";
         //std::cout << std::endl;
-        //delete[] KCacheVecIdx;
+        //delete[] KCacheRowIdx;
 
         //kernelSelectJ1<<<dimGrid, dimBlock>>>(d_reduceval, d_reduceidx, d_y, d_g, d_alpha, C, num_vec);
         kernelSelectJCached<<<dimGrid, dimBlock>>>(d_reduceval, d_reduceidx, d_y, d_g, d_alpha, C, num_vec, num_vec_aligned, d_workingset, d_K, d_KDiag, d_KCacheRemapIdx);

    //export_cuda_buffer(d_K, num_vec_aligned, num_vec_aligned, sizeof(float), "K.dat");

    int cacheLastPtrIdx = 0;
    size_t kernelCheckCacheSMSize = dim * sizeof(float);
    dim3 dimBlock(reduce_block_size);
    dim3 dimGrid(getgriddim(num_vec_aligned, (size_t)dimBlock.x));
    dim3 dimBlockCache(256);
    dim3 dimGridCache(getgriddim(num_vec_aligned, (size_t)dimBlockCache.x));
    for (int iter = 0;; iter++)
    {
        kernelSelectI<<<dimGrid, dimBlock>>>(d_reduceval, d_reduceidx, d_y, d_g, d_alpha, C, num_vec);

    //kernelInitg<<<dimGrid, dimBlock>>>(d_g, num_vec_aligned);
    memsetCuda<float>(d_g, 1, num_vec_aligned);
    memsetCuda<int>(d_KCacheRemapIdx, -1, num_vec);
    memsetCuda<int>(d_KCacheRowIdx, -1, cache_rows + 2);
    memsetCuda<int>(d_KCacheRowPriority, -1, cache_rows);
    int cacheUpdateCnt = 0;
    assert_cuda(cudaMemcpyToSymbol(d_cacheUpdateCnt, &cacheUpdateCnt, sizeof(int), 0));

    //export_cuda_buffer(d_g, num_vec_aligned, 1, sizeof(float), "g.dat");

