         HANDLE_ERROR(cudaStatus);
 
 	int block;
-	//cudaFuncSetCacheConfig(fw_kernel_phase_three<BLOCK_SIZE, THREAD_SIZE>, cudaFuncCachePreferShared);
 	FOR(block, 0, numOfBlock) 
 	{
 		fw_kernel_phase_one<VIRTUAL_BLOCK_SIZE><<<1, dimBlockP1>>>(block, n, pitch_int, dev_d, dev_p);

	// Synchronize to make sure the all value are loaded in block
	__syncthreads();

		
	#pragma unroll
	FOR(i, 0, BLOCK_SIZE - 1)
	{

	}
	
	const int cell_current = v1 * pitch + v2;
	if (v1 < n && v2 < n)
	{
		current_d[ty][tx] = d[cell_current];
