hsaco v3 metadata kernarg size

[zgu153522]

I have a kernel compiled through HIP with: hipcc --genco --amdgpu-target=gfx906 -save-temps $< -o $@

I inspected the resulting ISA and found: .kernarg_segment_size: 72

So I assembled the hsaco with: /opt/rocm/llvm/bin/clang -target amdgcn-amd-amdhsa -mcpu=gfx906 -o $@ $<

If I open the binary I can see the plain text portion with: .kernarg_segment_sizeH<A9>

But when I load the kernel through HSA and query the kernarg size with: hsa_executable_symbol_get_info( symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &kernargSize);

I end up getting 0 for a kernargSize. From hsaco v2 (old compile with HCC) I don't have this problem.

[zgu153522]

Is there some issue with the way I'm generating the hsaco possibly, that would make ROCr not be able to read the correct kernarg size back?

[t-tye]

See https://llvm.org/docs/AMDGPUUsage.html#code-object-v3-kernel-descriptor for a definition of the kernel descriptor for the current code object version. The definition has changed across code object versions. I believe that the code object V4 released as part of ROCm 4.1 does fill in the kernarg size in the kernel descriptor, but V3 left the field 0. V2 used by HCC was likely different.

The high level languages obtain the kernarg size from the metadata. See https://llvm.org/docs/AMDGPUUsage.html#amdgpu-amdhsa-code-object-kernel-metadata-map-table-v3 . The COMgr library can be used to access code object metadata.

Hope that helps.