Error with sub-8B kernel arguments

[dhjoo98]

Possibly related to this comment in intercept.cc,

it seems like Omniprobe returns the following when each kernel argument is smaller that 8 Bytes.

>>>>>>>> HSA intercept registered.
ADDRESS_MESSAGE,timestamp,kernel,src_line,dispatch,exec_mask,xcc_id,se_id,cu_id,kind,address
host_device_combined_omniprobe: /app/omniprobe/src/interceptor.cc:649: void hsaInterceptor::fixupKernArgs(void *, void *, void *, arg_descriptor_t): Assertion `desc.clone_hidden_args_length <= desc.kernarg_length - desc.explicit_args_length' failed.

This was encountered with the following kernel argument - when 32b integer arguments are used,

__global__ void hgemm_kernel(const __half* __restrict__ A,
                             const __half* __restrict__ B,
                             __half* __restrict__ C,
                             int M, int N, int K)

but not encountered (proceeds to instrumentation) with the following kernel argument - when replaced with 64b integers,

__global__ void hgemm_kernel(const __half* __restrict__ A,
                             const __half* __restrict__ B,
                             __half* __restrict__ C,
                             int64_t M, int64_t N, int64_t K)

[rwvo]

This may have been fixed already; I looked into kernel argument alignment before. Should be testable by writing a kernel that has similar arguments as the one above with the three int arguments, launch it with a single block of 64 threads, let thread 0 printf the values of the integer arguments. Then instrument it using instrumentation scoping such that nothing is actually instrumented, but a cloned kernel with an extra void* argument is created. Then run with and without instrumentation under omniprobe. Both runs should print the same values.

[rwvo]

@Copilot: I'm unconvinced that your PR isn't an attempt to fix something that isn't broken. Write a test that shows that the implementation of kernarg size calculations as on the main branch fails, while the implementation in your PR succeeds. Use the idea for a test kernel in my previous comment. If you think that idea doesn't distinguish between broken and correct behavior, then address that. You may have to write a kernel that forces the presence of hidden kernel arguments to do that, so that you can compare the hidden arguments of the original kernel with the modified kernel. Report back with your findings.