 	       	 * CUDA IMPLEMENTATION
 		 */
 
-		MatrixMultiplyKernel<<<dimGridCA, dimBlockCA>>>(device_MatrixCA, device_state1, device_output_chunk);
+		MatrixMultiplyKernel<<<dimGridCA, dimBlockCA>>>(device_MatrixCA, device_state_read, device_output_chunk);
 
 		cudaDeviceSynchronize();	
 		cudaMemcpy(output_chunk.elements, device_output_chunk.elements, m_size(output_chunk), cudaMemcpyDeviceToHost);

		for(j = 0; j < blocksize; j++) {
			output[i+j] = m_get(output_chunk,j,0)/128;
		}
		MatrixMultiplyKernel<<<dimGridA, dimBlockA>>>(device_MatrixAp, device_state1, device_state2);
		device_tmp_state = device_state1;
		device_state1 = device_state2;
		device_state2 = device_tmp_state;
#else
		/*
	       	 * CPU IMPLEMENTATION

	CUDA_SAFE_CALL(cudaMalloc((void**) &device_output_chunk.elements, m_size(output_chunk)));
	CUDA_SAFE_CALL(cudaMemcpy(device_output_chunk.elements, output_chunk.elements, m_size(output_chunk), cudaMemcpyHostToDevice));

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

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

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