 
         //check if J is cached
         kernelCheckCache<<<dimGridCache, dimBlockCache, kernelCheckCacheSMSize>>>(d_workingset + 1, d_K, d_KCacheRemapIdx, d_KCacheVecIdx, cache_rows, d_x, d_xT, gamma, num_vec, num_vec_aligned, dim, dim_aligned, cacheLastPtrIdx);
+        kernelCheckCacheFinalize<<<1, 1>>>(d_KCacheRemapIdx);
         cacheLastPtrIdx = (cacheLastPtrIdx + 1) % 2;
         //workaround if caching J deleted I out of cache
         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);
+        kernelCheckCacheFinalize<<<1, 1>>>(d_KCacheRemapIdx);
         cacheLastPtrIdx = (cacheLastPtrIdx + 1) % 2;
 
         if (iter > 0 && iter % 1000 == 0)

    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;
    }


    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);

