     alpha[j] -= l * y[j];
 }
 
+//contains changes to KCacheRemapIdx buffer, which should be written after kernelCheckCache ends
+//each change to buffer is contained in int2 variable (x,y) such that
+//KCacheRemapIdx[x] = y
+__device__ int2 d_KCacheRemapIdxChanges[2];
+
 __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 j = blockDim.x * blockIdx.x + threadIdx.x;
     int i = *i_ptr;
     if (KCacheRemapIdx[i] >= 0)
+    {
+        if (j == 0)
+            KCacheVecIdx[cache_rows + (1 - lastPtrIdx)] = KCacheVecIdx[cache_rows + lastPtrIdx];
         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;
+        KCacheVecIdx[cache_rows + (1 - lastPtrIdx)] = 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
+            d_KCacheRemapIdxChanges[1] = make_int2(del_i, 1);  //cache row for vector [del_i] will be overwritten, remove it from RemapIdx array
         //set correct indices
-        KCacheRemapIdx[i] = last;
+        d_KCacheRemapIdxChanges[0] = make_int2(i, last);
         KCacheVecIdx[last] = i;
     }
 


        //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(d_x, x, num_vec_aligned * dim_aligned * sizeof(float), cudaMemcpyHostToDevice));
    assert_cuda(cudaMemcpy(d_y, y, num_vec_aligned * sizeof(float), cudaMemcpyHostToDevice));

    int KCacheRemapIdxChanges[4] = { -1, -1, -1, -1 };
    assert_cuda(cudaMemcpyToSymbol(d_KCacheRemapIdxChanges, KCacheRemapIdxChanges, sizeof(KCacheRemapIdxChanges), 0));

    float a = 1, b = 0;
    assert_cublas(cublasSgeam(cublas, CUBLAS_OP_T, CUBLAS_OP_N, num_vec, dim, &a, d_x, dim_aligned, &b, d_x, num_vec_aligned, d_xT, num_vec_aligned));

