 #endif
 		}
 	//	tmp_state = m_multiplyblockdiag(MatrixAp,state,2);
-		MatrixMultiplyKernel<<<dimGridA, dimBlockA>>>(dMatrixAp, dstate1, dstate2);
-		dtmp_state = dstate1;
-		dstate1 = dstate2;
-		dstate2 = dtmp_state;
+		MatrixMultiplyKernel<<<dimGridA, dimBlockA>>>(device_MatrixAp, device_state1, device_state2);
+		device_tmp_state = device_state1;
+		device_state1 = device_state2;
+		device_state2 = device_tmp_state;
 	//	state = tmp_state;
 		i = i + blocksize;
 	}

	Matrix MatrixCA;
	Matrix tmp_state;

	Matrix dMatrixAp, dblock_CA, dstate1, dstate2, dtmp_state, dblock_samples;

	cudaSetDevice(0);


	m_print(state);
#endif

	CUDA_SAFE_CALL(cudaMalloc((void**) &device_MatrixAp.elements, m_size(MatrixAp)));
	CUDA_SAFE_CALL(cudaMemcpy(device_MatrixAp.elements,MatrixAp.elements, m_size(MatrixAp), cudaMemcpyHostToDevice));

	CUDA_SAFE_CALL(cudaMalloc((void**) &device_block_CA.elements, m_size(block_CA)));
	CUDA_SAFE_CALL(cudaMemcpy(device_block_CA.elements,block_CA.elements, m_size(block_CA), cudaMemcpyHostToDevice));

	CUDA_SAFE_CALL(cudaMalloc((void**) &device_block_samples.elements, m_size(block_samples)));
	CUDA_SAFE_CALL(cudaMemcpy(device_block_samples.elements,block_samples.elements, m_size(block_samples), cudaMemcpyHostToDevice));

	CUDA_SAFE_CALL(cudaMalloc((void**) &device_state1.elements, m_size(state)));
	CUDA_SAFE_CALL(cudaMemcpy(device_state1.elements,state.elements, m_size(state), cudaMemcpyHostToDevice));

	CUDA_SAFE_CALL(cudaMalloc((void**) &device_state2.elements, m_size(state)));
	CUDA_SAFE_CALL(cudaMemcpy(device_state2.elements,state.elements, m_size(state), cudaMemcpyHostToDevice));

	dim3 dimBlockCA(1, 1);
	dim3 dimGridCA(state.cols / dimBlockCA.x, block_CA.rows / dimBlockCA.y);
