Bug fixes for 1) over-counting instructions 2) broken functional sim

[SerinaTan]

Apologies for accidentally merging two separate commits into a single PR. (I didn't know pushing a new commit to the same branch will update the existing PR...). Anyhow, here are the descriptions for the two bug fixes:

Fix No.1:
In case of a "vector load" instruction which has multiple register destinations (e.g. ld.global.v4.u32 {%r1903, %r1902, %r1905, %r1904}, [%rd384]), ldst_unit::L1_latency_queue_cycle() would call warp_inst_complete() multiple times and hence over-count the number of completed instructions. This behavior is inconsistent with other ldst_unit functions such as ldst_unit::writeback().

Fix: move the warp_inst_complete() call out of the for loop iterating output registers.

Fix No.2:
In gpgpu_cuda_ptx_sim_main_func, kernel.increment_cta_id() is NOT called when checkpoint option is disabled during functional simulation. This causes functionalCoreSim::initializeCTA to be called twice in a row with the same cta_id and eventually a seg fault.

Fix: move kernel.increment_cta_id() out of the else block so that it's always executed. I am assuming the intended behaviour of the if-block is to only allow 1) checkpoint off or 2) prior to perf sim resume to execute the cta instructions. However, whether this if-condition is true/false, we should always increment the kernel's cta id or else the while loop won't break.

[tgrogers]

Hey, Serina - do you know what this if is testing for?
I know this is not your code - but your change effect code flows when this thing is true and I cannot tell from quick inspection what it is for.

Also, thanks for the fixes! :)

[SerinaTan]

I think the if-statement guards whether to perform functional simulation for each CTA. It is evaluated to true when any of the following is true: 1) not doing checkpoint cp_op==0 2) checkpoint is on and we have reached the checkpoint kernel boundary, but we have not reached the CTA boundary cp_op==1 && cta_launched<cp_cta_resume && kernel.get_uid()==cp_kernel 3) we have not reached the checkpoint kernel boundary kernel.get_uid() < cp_kernel. We need to call kernel.increment_cta_id() so that the functional simulation can progress with an updated cta id (returned by kernel.get_next_cta_id_single()).

If this if-statement is off, we have reached the checkpoint boundary and we should halt functional simulation and transition into performance simulation. However, we still need to break the while loop by incrementing the kernel cta id (kernel.increment_cta_id()).

That being said, I am actually not sure how this code EVER worked with checkpointing...