     //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_KCacheVecIdx, -1, cache_rows + 2);
+    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");
 

    alpha[j] -= l * y[j];
}

__global__ void kernelCheckCache(const int * i_ptr, float * K, int * KCacheRemapIdx, int * KCacheVecIdx, int cache_rows, const float * x, const float * xT, float gamma, int num_vec, int num_vec_aligned, int dim, int dim_aligned, int lastPtrIdx)
{
    int i = *i_ptr;
    if (KCacheRemapIdx[i] >= 0)
        return;  //item already in cache
    int j = blockDim.x * blockIdx.x + threadIdx.x;
    int last = (KCacheVecIdx[cache_rows + lastPtrIdx] + 1) % cache_rows;
    if (j == 0)
    {
        KCacheVecIdx[cache_rows + (lastPtrIdx + 1) % 2] = last;
        int del_i = KCacheVecIdx[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;
        KCacheVecIdx[last] = i;
    }

    //calculate cache matrix row [last], original index is [i]
    extern __shared__ float xi[];
    int idxshift = 0;
    while (idxshift < dim)
    {
        int idx = idxshift + threadIdx.x;
        if (idx < dim)
            //xi[idx] = xT[num_vec_aligned * idx + i];
            xi[idx] = x[dim_aligned * i + idx];
        idxshift += blockDim.x;
    }
    __syncthreads();
    while (j < num_vec)

    float *d_lambda = nullptr;
    int *d_workingset = nullptr,
        *d_KCacheRemapIdx = nullptr,
        *d_KCacheRowIdx = nullptr,  // items at index [cache_rows] and [cache_rows+1] are indices of last inserted item
        *d_KCacheRowPriority = nullptr;  // the higher the priority is, the later was the item added

    size_t reduce_block_size = 256;
    size_t reduce_buff_size = rounduptomult(num_vec, reduce_block_size);
    size_t ones_size = std::max(num_vec_aligned, dim_aligned);
    size_t cache_size_mb = 2000;
    size_t cache_rows = cache_size_mb * 1024 * 1024 / (num_vec_aligned * sizeof(float));
    cache_rows = std::min(cache_rows, num_vec);

