Too many arguments for injection function will cause `IllegalAddress`

[Ivecia]

If we add some arguments for injection function of Accel-Sim, IllegalAddress will be reported.

NVBit version: 1.7.6

Injection function:

https://github.com/accel-sim/accel-sim-framework/blob/3c96d324180d0023bfbddbd693d8d3cbc1170967/util/tracer_nvbit/tracer_tool/inject_funcs.cu#L22C1-L28C57

After modification (which causes IllegalAddress):

extern "C" __device__ __noinline__ void
instrument_inst(int pred, int opcode_id, int32_t vpc, bool is_mem,
                uint64_t addr, int32_t width, int32_t desReg, int32_t srcReg1,
                int32_t srcReg2, int32_t srcReg3, int32_t srcReg4,
                int32_t srcReg5, int32_t srcNum,
                // newly added arguments
                int32_t dstType, int32_t srcType1, int32_t srcType2, int32_t srcType3,
                int32_t srcType4, int32_t srcType5, int32_t predReg,
                // end here
                uint64_t immediate,
                uint64_t pchannel_dev, uint64_t ptotal_dynamic_instr_counter,
                uint64_t preported_dynamic_instr_counter, uint64_t pstop_report,
                uint32_t line_num, uint32_t instr_idx)

Probably some arguments are stored in local memory? And interestingly, if we remove one of arguments, it will not report illegal address.

After modification (which does not cause IllegalAddress):

extern "C" __device__ __noinline__ void
instrument_inst(int pred, int opcode_id, int32_t vpc, bool is_mem,
                uint64_t addr, int32_t width, int32_t desReg, int32_t srcReg1,
                int32_t srcReg2, int32_t srcReg3, int32_t srcReg4,
                int32_t srcReg5,                                    // remove srcNum (since it is not used eventually)
                // newly added arguments
                int32_t dstType, int32_t srcType1, int32_t srcType2, int32_t srcType3,
                int32_t srcType4, int32_t srcType5, int32_t predReg,
                // end here
                uint64_t immediate,
                uint64_t pchannel_dev, uint64_t ptotal_dynamic_instr_counter,
                uint64_t preported_dynamic_instr_counter, uint64_t pstop_report,
                uint32_t line_num, uint32_t instr_idx)