<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/86332>86332</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[AMDGPU] With Clang>17, -amdgpu-early-inline-all=true consumes 8x more memory
</td>
</tr>
<tr>
<th>Labels</th>
<td>
clang
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
AngryLoki
</td>
</tr>
</table>
<pre>
There is some kind of regression in `-amdgpu-early-inline-all=true` option, which is set for every HIP application in hipcc.
While this option makes no significant performance/memory impact in Clang 17, attempt to migrate to Clang 18.1.0 or nightly Clang 19 build consumes 8x more memory, which makes Clang unusable for HIP (i. e. when multiple compile units consume 10GB each in parallel, there is just not enough RAM eventually, even when compiling for single target GPU arch).
Environment:
```
/usr/lib/llvm/17/bin/clang-17 --version | grep version
clang version 17.0.6
/usr/lib/llvm/18/bin/clang-18 --version | grep version
clang version 18.1.0
/usr/lib/llvm/19/bin/clang-19 --version | grep version
clang version 19.0.0git6d3cec01
```
Common flags (verbose output of [composable-kernel-6.0.2](https://github.com/ROCm/composable_kernel/tree/rocm-6.0.2)):
```
export FLAGS="-cc1 -triple amdgcn-amd-amdhsa -aux-triple x86_64-pc-linux-gnu -emit-obj -disable-free -clear-ast-before-backend -disable-llvm-verifier -discard-value-names -main-file-name device_batchnorm_forward_f32_instance.cpp -mrelocation-model pic -pic-level 2 -fhalf-no-semantic-interposition -mframe-pointer=none -fno-rounding-math -mconstructor-aliases -aux-target-cpu x86-64 -fcuda-is-device -mllvm -amdgpu-internalize-symbols -fcuda-allow-variadic-functions -fvisibility=hidden -fapply-global-visibility-to-externs -mlink-builtin-bitcode /usr/lib/amdgcn/bitcode/hip.bc -mlink-builtin-bitcode /usr/lib/amdgcn/bitcode/ocml.bc -mlink-builtin-bitcode /usr/lib/amdgcn/bitcode/ockl.bc -mlink-builtin-bitcode /usr/lib/amdgcn/bitcode/oclc_daz_opt_off.bc -mlink-builtin-bitcode /usr/lib/amdgcn/bitcode/oclc_unsafe_math_off.bc -mlink-builtin-bitcode /usr/lib/amdgcn/bitcode/oclc_finite_only_off.bc -mlink-builtin-bitcode /usr/lib/amdgcn/bitcode/oclc_correctly_rounded_sqrt_on.bc -mlink-builtin-bitcode /usr/lib/amdgcn/bitcode/oclc_wavefrontsize64_off.bc -mlink-builtin-bitcode /usr/lib/amdgcn/bitcode/oclc_isa_version_1030.bc -mlink-builtin-bitcode /usr/lib/amdgcn/bitcode/oclc_abi_version_400.bc -target-cpu gfx1030 -debugger-tuning=gdb -fdebug-compilation-dir=/var/tmp/portage/sci-libs/composable-kernel-6.0.2/work/composable_kernel-rocm-6.0.2_build -resource-dir /usr/lib/clang/17 -dependency-file library/src/tensor_operation_instance/gpu/batchnorm/CMakeFiles/device_batchnorm_instance.dir/device_batchnorm_forward_f32_instance.cpp.o.d -MT library/src/tensor_operation_instance/gpu/batchnorm/CMakeFiles/device_batchnorm_instance.dir/device_batchnorm_forward_f32_instance.cpp.o -sys-header-deps -internal-isystem /usr/lib/clang/17/include/cuda_wrappers -idirafter /usr/local/include -include __clang_hip_runtime_wrapper.h -include /usr/include/gentoo/fortify.h -include /usr/include/gentoo/maybe-stddefs.h -D CK_ENABLE_BF16 -D CK_ENABLE_BF8 -D CK_ENABLE_FP16 -D CK_ENABLE_FP32 -D CK_ENABLE_FP64 -D CK_ENABLE_FP8 -D CK_ENABLE_INT8 -D USE_PROF_API=1 -D __HIP_PLATFORM_AMD__=1 -D __HIP_PLATFORM_HCC__=1 -I /var/tmp/portage/sci-libs/composable-kernel-6.0.2/work/composable_kernel-rocm-6.0.2/library/include -I /var/tmp/portage/sci-libs/composable-kernel-6.0.2/work/composable_kernel-rocm-6.0.2/include -I /var/tmp/portage/sci-libs/composable-kernel-6.0.2/work/composable_kernel-rocm-6.0.2_build/include -internal-isystem /usr/lib/gcc/x86_64-pc-linux-gnu/13/include/g++-v13 -internal-isystem /usr/lib/gcc/x86_64-pc-linux-gnu/13/include/g++-v13/x86_64-pc-linux-gnu -internal-isystem /usr/lib/gcc/x86_64-pc-linux-gnu/13/include/g++-v13/backward -internal-isystem /usr/lib/gcc/x86_64-pc-linux-gnu/13/include/g++-v13 -internal-isystem /usr/lib/gcc/x86_64-pc-linux-gnu/13/include/g++-v13/x86_64-pc-linux-gnu -internal-isystem /usr/lib/gcc/x86_64-pc-linux-gnu/13/include/g++-v13/backward -internal-isystem /usr/lib/clang/17/include -internal-isystem /usr/local/include -internal-isystem /usr/x86_64-pc-linux-gnu/include -internal-externc-isystem /include -internal-externc-isystem /usr/include -internal-isystem /usr/lib/clang/17/include -internal-isystem /usr/local/include -internal-isystem /usr/x86_64-pc-linux-gnu/include -internal-externc-isystem /include -internal-externc-isystem /usr/include -O3 -std=c++17 -fdeprecated-macro -fno-autolink -ferror-limit 19 -fmessage-length=173 -fhip-new-launch-api -fgnuc-version=4.2.1 -fcxx-exceptions -fexceptions -fcolor-diagnostics -vectorize-loops -vectorize-slp -mllvm -amdgpu-function-calls=false -cuid=aa0b75146f478e4b -fcuda-allow-variadic-functions -faddrsig -D__GCC_HAVE_DWARF2_CFI_ASM=1 -o /tmp/device_batchnorm_forward_f32_instance-gfx1030-437c24.o -x hip /var/tmp/portage/sci-libs/composable-kernel-6.0.2/work/composable_kernel-rocm-6.0.2/library/src/tensor_operation_instance/gpu/batchnorm/device_batchnorm_forward_f32_instance.cpp"
```
Without `-amdgpu-early-inline-all=true` everything is fine:
```
/usr/bin/time -f 'Memory: %M KB, Time: %E' /usr/lib/llvm/17/bin/clang-17 $FLAGS
Memory: 818272 KB, Time: 0:20.62
/usr/bin/time -f 'Memory: %M KB, Time: %E' /usr/lib/llvm/18/bin/clang-18 $FLAGS
Memory: 830300 KB, Time: 0:18.28
/usr/bin/time -f 'Memory: %M KB, Time: %E' /usr/lib/llvm/19/bin/clang-19 $FLAGS
Memory: 861772 KB, Time: 0:22.69
```
With `-amdgpu-early-inline-all=true` Clang 18 and 19 are hungry and slow:
```
/usr/bin/time -f 'Memory: %M KB, Time: %E' /usr/lib/llvm/17/bin/clang-17 $FLAGS -mllvm -amdgpu-early-inline-all=true
Memory: 818240 KB, Time: 0:20.80
/usr/bin/time -f 'Memory: %M KB, Time: %E' /usr/lib/llvm/18/bin/clang-18 $FLAGS -mllvm -amdgpu-early-inline-all=true
Memory: 6402824 KB, Time: 1:02.50
/usr/bin/time -f 'Memory: %M KB, Time: %E' /usr/lib/llvm/19/bin/clang-19 $FLAGS -mllvm -amdgpu-early-inline-all=true
Memory: 6343976 KB, Time: 1:12.43
```
I don't provide preprocessed version of [device_batchnorm_forward_f32_instance.cpp](https://github.com/ROCm/composable_kernel/blob/rocm-6.0.2/library/src/tensor_operation_instance/gpu/batchnorm/device_batchnorm_forward_f32_instance.cpp), because for some reason I can't rebuild it after preprocessing (complaints about constexprs). However if you need it or some other dumps, please ask and I will attach.
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzsWVtv47by_zTMy4CCRPn6kAfHjnaDbrrBdvvvo0BRI4kNRepPUk7cT39AyU42rpNN9xIUOAdwHGtEzo3D3wyH3DlZa8RzMr0g080Z731j7PlK13b3wdzKs8KUu_PPDVoE6cCZFuFW6hJMBRZri85Jo0FqILOY8rasu54it2pHpVZSI-VKkXTjbY9kFoPpvDSasDXcNVI0A0_0UBkLuEW7g_dXN8C7TknB_Z5zIzshIhJvSLwav_9opELwjXR7htDyW3SgDQRzZCUF1x46tJWxLdcCCctabI3dgWw7Lnzgu1Zc15DMgzbce2w7D95AK2vLPYaf-xGLKIliMBa0rBuvdgf6EopeqhKE0a5v0cHiHlpjEUZRj1aO2o2zet07XigcbA7WEraQEWAEdw1qaHvlZacQhGm7YGWvpXcHEZDE7y4AefCcho5brhSqIMgflujP3nnQxgNq09cNfFpdB9dq33OlBp3C0yhslCF1PSjjpK6DW7mt0cO7m9-BW9EQtozgS-df6q20RreoPUn3NDKL95_xkWW9s4RlShbhW21bwrLg6ayQmrBMBFfQZA6UbtEOIUTma6gtdrAnjJyGgQcSJPMojmZfKvOMpMWxpMU_lTSs-SskLY8lLf-ppGUUR3Et_axMBYo4Oe3R4Xtt2tZoqBSvXYibLdrCOATT-673YU-S6UVYVDOEGL1Fq1HRWRRHjEw3hC0a7zsXlo1lhGW19E1fRMIEUz59XId_j9PzcTphmbcYdpA1ot0zY8vweWb58b4z1kP2YfXuN5JuCGNUiASot0NkB5QQOoBF-GscB8r7-8Pb-8Usn01oJ6iSur-nte6BYis9NcWfQEs5mlZZRKBCIbeUO08LrIxFWnBxi7p8HBfWKayHrCTagSy4LemWqx6p5mHX0pZLTSupRgKUuJUC84J70Whj27wy9o7bMq9SlkvtfACUSHQd0NaiMiNS0daUqKCTAmgnBVW4RQUMaNVwVVFtqMOWay8Fldqj7YyTA3TRtrK8RdqZgU7SjTYagVbaUGt6XUpd05b7BmgbYMDbXnhjKVeSu6D-4Lxh01LR9cGBdDYBWom-5FQ6OtoDtA2-gANID8I0V_IvpG7XFka5wxyulLmjW24lL6WgVa9F0DS830onC6mk35F008iyRA20CoC9o7UyBVf0cQj1huJ9EBOcrKS-pQEwvdS0kF6YEuFoV42hMWyp4T1hWSO7qBDfPt-IVn0ng9vvZaBEXvK_ctP53FTV9zPrteMV5iEofgzDSmrpMTda7X4MQ2GsReHVLh9CGMvc_b_1udHfz_qOb7GyRnsn_8LZ5MfoKx3P94icJ3Eafz9HXsgHjpN4ZPjFLq2r-yAHaIlFX9doqe-11DVJN3VZAK0GOh0z9AgwpbQDmmZbHsT7tiMsC0jL6yDVCUmVLNwTDH-aAlh2Z-ztKZCnj9iej0UNtehMbwUGucdGDylsSOnBgA51iVrsBgwFJQvLQ_WTOSuCnqidsbnp0A52PEBoSEBdHzx3QFrCsvU1v8VMKgx2_A2JH9A3-OLE--eQOjJRCfT6879UOaBu52iDvEQb_OngAZ6pdDvnsX1-BQjLpBaqH0Iv4Hd-Z3nXoQ1cSml55fHLBTSCq8c5QdL4I88Hnnkju9z22ssWD5yi5nHYA6NHoTVqbwxhWWWsl9XutcNbviuQOl-WWLkwaQPrX_LLX1cXHy7ziyyZHVMWTwnZzfGQ7CZlx5SQC59Sjthc_fp5oPz-22V-8-ljlq9urki6SQItz99f3eQ3H1afs4-frvPV9SbPn3n3fr0-vLuCn75Lx0DYx_LDWr6N4LcUN6LR03h9eWfUImzsE2Vk2C3p00gk7IKwC7pN0p_F9_SUnygtVMABZv7nqB_mqFNY--Ksv0Psc0NP6_33iWMZLb5k8KpBT7D3v8nSjymEzELSjRhXPZQqVYmdRcE9lrTlwprxlMV7b0KpB7RCa42lSrbSQzjKVy06x2ukCnXtmwDv8zSc6WRHNd5RxXstGso7CbSqdS8OZ3-SbiYRi5Jwprq_p3gvsDscop48CKOMpaXktTbOS-GAbjEc8MLBTBnTPSE41R0f5A7HMyq4Uo6km4orF87GvQzmcx4X82kymVWT-QInxSsOebwsrZM10E2ev1uv8_er_7vMN3-sPmUsX2dX-eq36zHPGXhA_ldVO3Rf9dJJOhdsEgqfe2hk98bp8ltKv1dXc4SxFzo4f0jfmN6_sls6tER9I3UN0kElNX614zb2okLxBrQCwubXYy8yXQFh02v45YKwNXyWLe5Jl4TNj2HghYYdYZOxrzOIfWS-SBZszo7YxyRdsTiasZONtB-q6omO37OqpnEaxydUTRYRW_x8VU-0DJ9VdZbMT3uVRbPlV-LslUF2aHID12XAPG4Rml7XdjdQnDJ3_5qoO8a-58w6EZyTUyvO4mhxusv7VsH5jRbNJjFbsMmRGglJVzGLpm9g0gtB_K0mpZN0OZ-dMClh0SR9IdivoDSasLmHzpqtLBE6i501Ap3D8qHTPrbHXw_k39EwL5Qpjhrmb5d_lsF_BQreu_GWabi1s8id0XAFgo--sjg2eqSHsUXw6LSQcAhbBMsUl9o74EXIWkP7Ge8764ZboffmLmQokBXsTA8aceB2kGh8gxbKvu1c0KhTyB0Cd7cDrlzBnVQKuPdcNNFZeZ6Wy3TJz_A8mSdJks6TOTtrzqs4xjjmOC3TebGcsfl0upgtUjadVXzGJ8WZPGcxm8QpY8lyyqYswmSa8mnF05iliwo5mcTYcqmiEJSRsfWZdK7H88UsTdmZ4gUqN1x_MrYvfxmZbs7s-XB1UPS1I5NYSefdIwcvvRruTFfXm3c3v5PpBgbEHbCUpJfjpeJXNsBz94ZnvVXnL8TdfgsO6nXW_InCh8o3GBXqpMGu_wQAAP__0wrAgQ">