     }
 }
 
-__global__ void kernelCheckCache_(const int * i_ptr, float * K, int * KCacheRemapIdx, int * KCacheRowIdx, int * KCacheRowPriority, 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)
+__global__ void kernelCheckCachePriority(const int * i_ptr, float * K, int * KCacheRemapIdx, int * KCacheRowIdx, int * KCacheRowPriority, int cache_rows, const float * x, const float * xT, float gamma, int num_vec, int num_vec_aligned, int dim, int dim_aligned)
 {
     extern __shared__ int2 spriority[];
+    int j = blockDim.x * blockIdx.x + threadIdx.x;
     int i = *i_ptr;
-    int j = blockDim.x * blockIdx.x + threadIdx.x;
     if (KCacheRemapIdx[i] >= 0)
     {
         if (j == 0)
             KCacheRowPriority[KCacheRemapIdx[i]] = d_cacheUpdateCnt;  // refresh priority
         return;  //item already in cache
     }
-    //int last = (KCacheRowIdx[cache_rows + lastPtrIdx] + 1) % cache_rows;
     int2 minpriority = make_int2(INT_MAX, 0);
     for (int k = 0; k < cache_rows; k += blockDim.x)
     {

        assert_cuda(cudaMemcpy(d_workingset + 1, d_reduceidx, sizeof(int), cudaMemcpyDeviceToDevice));

        //check if J is cached
        kernelCheckCache<<<dimGridCache, dimBlockCache, kernelCheckCacheSMSize>>>(d_workingset + 1, 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;
        //workaround if caching J deleted I out of cache
        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));

        //delete[] KCacheRowIdx;

        //check if I is cached
        if (usePriorityCache)
            kernelCheckCachePriority<<<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);
        else
        {
            kernelCheckCache << <dimGridCache, dimBlockCache, kernelCheckCacheSMSize >> >(d_workingset, d_K, d_KCacheRemapIdx, d_KCacheRowIdx, 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));
