 	dim3 dimBlockA(1, 1);
 	dim3 dimGridA(state.cols / dimBlockA.x, MatrixAp.rows / dimBlockA.y);
 
+	cudaEventRecord(MatrixCA_start, streams[0]);
 	MatrixMultiplyKernel<<<dimGridCA, dimBlockCA, 1, streams[0]>>>(device_MatrixCA, device_state_read, device_output_chunk_write);
+	cudaEventRecord(MatrixCA_stop, streams[0]);
+
+	cudaEventRecord(MatrixAp_start, streams[1]);
 	MatrixMultiplyKernel<<<dimGridA, dimBlockA, 1, streams[1]>>>(device_MatrixAp, device_state_read, device_state_write);
+	cudaEventRecord(MatrixAp_stop, streams[1]);
 #else
 	Matrix state_tmp;
 #endif


		cudaThreadSynchronize();

		device_state_tmp = device_state_read;
		device_state_read = device_state_write;
		device_state_write = device_state_tmp;

		output_chunk_read = output_chunk_write;
		output_chunk_write = output_chunk_tmp;

		cudaEventRecord(MatrixCA_start, streams[0]);
		MatrixMultiplyKernel<<<dimGridCA, dimBlockCA, 1, streams[0]>>>(device_MatrixCA, device_state_read, device_output_chunk_write);
		cudaEventRecord(MatrixCA_stop, streams[0]);

		cudaEventRecord(MatrixAp_start, streams[1]);
		MatrixMultiplyKernel<<<dimGridA, dimBlockA, 1, streams[1]>>>(device_MatrixAp, device_state_read, device_state_write);
		cudaEventRecord(MatrixAp_stop, streams[1]);

		cudaEventRecord(Memcpy_start, streams[2]);
		cudaMemcpyAsync(output_chunk_write.elements, device_output_chunk_read.elements, m_size(output_chunk_write), cudaMemcpyDeviceToHost, streams[2]);
		cudaEventRecord(Memcpy_stop, streams[2]);

		for(j = 0; j < blocksize; j++) {
			output[i+j] = m_get(output_chunk_read,j,0)/128;
