     {
         int idx = idxshift + threadIdx.x;
         if (idx < dim)
-            xi[idx] = x[num_vec_aligned * idx + i];
-            //xi[idx] = x[dim_aligned * i + idx];
+            //xi[idx] = xT[num_vec_aligned * idx + i];
+            xi[idx] = x[dim_aligned * i + idx];
         idxshift += blockDim.x;
     }
     __syncthreads();

        //delete[] KCacheVecIdx;

        //check if I is cached
        kernelCheckCache<<<dimGridCache, dimBlockCache, kernelCheckCacheSMSize>>>(d_workingset, d_K, d_KCacheRemapIdx, d_KCacheVecIdx, cache_rows, 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_workingset + 1, d_reduceidx, sizeof(int), cudaMemcpyDeviceToDevice));

        //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);
        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);
        cacheLastPtrIdx = (cacheLastPtrIdx + 1) % 2;

        if (iter > 0 && iter % 1000 == 0)
