             //float diff = xT[num_vec_aligned * d + i] - xT[num_vec_aligned * d + j];
             sum += diff * diff;
         }
-        K[num_vec_aligned * last + j] = exp(-gamma * sum);
+        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;
+        }
     }
 }
 

    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 = 200;
    size_t cache_rows = cache_size_mb * 1024 * 1024 / (num_vec_aligned * sizeof(float));
    cache_rows = std::min(cache_rows, num_vec);



        //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];
