[Comgr] Support compressed device binaries

[GZGavinZhao]

This adds support for compressed device binaries implemented in llvm@7e28234.

Previously this has not been possible on the implementation level because to decompress the device binaries, we need to know the exact size of the binary, but __hipRegisterFatBinary only gives the starting address of the binary. Thanks to llvm#88827 that was merged last week, the size info of the device binary is now included in the compressed device binary header, so this is possible now.

This allows significant space saving in GPU binaries that have grown so large that they've exceed linker limits, e.g. ROCm/composable_kernel#789 and ROCm/composable_kernel#1044. In composable_kernel's case, with this patch the total size of static libraries library shrink from 3.69 GiB to 1.3GB, and this is only with the default compression settings (i.e. without tuning compression levels to optimize further).

I have yet to notice any significant runtime penalty. Even if there are, it would only be a one-time penalty at executable startup when __hipRegisterFatBinary is called to register the device binaries, so overall I believe we should be fine. In addition, --offload-compress would probably not be used anywhere except in special cases like in composable_kernel where binary sizes are big enough to cause linking issues.

This change should be NFC for binaries that are not compiled with --offload-compress. With --offload-compress, I've verified that all tests in composable_kernel are passed with gfx1030 and gfx900.

[lamb-j]

@GZGavinZhao thanks for putting this together.

Because we're using the OfflaodBundling API in Comgr to unbundle files before linking, we should pick up the decompression changes automatically without any Comgr changes.

What contexts would you expect a user to use the new "AMD_COMGR_DATA_KIND_COMPRESSED_FATBIN" data type? We did just add a new data type "AMD_COMGR_DATA_KIND_OBJ_BUNDLE" which may cover the same use case. See these two changes, which adds more support for unbundling:

[Comgr] Add OffloadBundler APIs
[Comgr] Split get_bundle_entry_ids into two API calls

Do you think this PR is still needed on top of those changes?

[GZGavinZhao]

@lamb-j Thanks for your response! The purpose of this PR was originally to enable loading compressed device binaries in clr through ROCm/clr#75, but it seems like ROCm/clr@e53df57 just implemented that feature, so I think this PR is not needed anymore. Thanks again for your time looking at this!