 	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

		output_chunk_read = output_chunk_write;
		output_chunk_write = output_chunk_tmp;

		MatrixMultiplyKernel<<<dimGridCA, dimBlockCA, 1, streams[0]>>>(device_MatrixCA, device_state_read, device_output_chunk_write);
		MatrixMultiplyKernel<<<dimGridA, dimBlockA, 1, streams[1]>>>(device_MatrixAp, device_state_read, device_state_write);

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

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


		cudaThreadSynchronize();

		cudaEventElapsedTime(&MatrixCA_time, MatrixCA_start, MatrixCA_stop);
		cudaEventElapsedTime(&MatrixAp_time, MatrixAp_start, MatrixAp_stop);
		cudaEventElapsedTime(&Memcpy_time, Memcpy_start, Memcpy_stop);

		if(i == 5*blocksize) {
			printf("MatrixCA: %d\n", MatrixCA_time);
			printf("MatrixAp: %d\n", MatrixAp_time);
			printf("  Memcpy: %d\n", Memcpy_time);
		}

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