 
         //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);
+        kernelCheckCacheFinalize<<<1, 1>>>(d_KCacheRemapIdx);
         cacheLastPtrIdx = (cacheLastPtrIdx + 1) % 2;
 
         //KCacheVecIdx = new int[cache_rows + 2];

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


            //float diff = xT[num_vec_aligned * d + i] - xT[num_vec_aligned * d + j];
            sum += diff * diff;
        }
        K[num_vec_aligned * last + j] = expf(-gamma * sum);
        j += gridDim.x * blockDim.x;
    }
}

__global__ void kernelCheckCacheFinalize(int * KCacheRemapIdx)
{
#pragma unroll
    for (int i = 0; i < 2; i++)
    {
        int2 c = d_KCacheRemapIdxChanges[i];
        if (c.x >= 0)
        {
            KCacheRemapIdx[c.x] = c.y;
            d_KCacheRemapIdxChanges[i].x = -1;
        }
    }
}

