Build issue with clang, lld, amdgcn-link for ROCM on linux

[dvogel]

I'm following this guide to build llama.cpp for an older AMD card. I'm confused by a link time error, along with some of the other output. Here's the

$ make -j1 GGML_NO_CCACHE=1 GGML_HIPBLAS=1 AMDGPU_TARGETS=gfx803 CC=$(which clang-19) CXX=$(which clang++-19) GGML_NO_OPENMP=1
I llama.cpp build info: 
I UNAME_S:   Linux
I UNAME_P:   unknown
I UNAME_M:   x86_64
I CFLAGS:    -Iggml/include -Iggml/src -Iinclude -Isrc -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_CPU_AARCH64 -DGGML_USE_LLAMAFILE -DGGML_USE_AMX -DGGML_USE_HIP -DGGML_USE_CUDA  -std=c11   -fPIC -O3 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wshadow -Wstrict-prototypes -Wpointer-arith -Wmissing-prototypes -Werror=implicit-int -Werror=implicit-function-declaration -pthread -march=native -mtune=native -Wunreachable-code-break -Wunreachable-code-return -Wdouble-promotion 
I CXXFLAGS:  -std=c++11 -fPIC -O3 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wunreachable-code-break -Wunreachable-code-return -Wmissing-prototypes -Wextra-semi -Iggml/include -Iggml/src -Iinclude -Isrc -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_CPU_AARCH64 -DGGML_USE_LLAMAFILE -DGGML_USE_AMX -DGGML_USE_HIP -DGGML_USE_CUDA 
I NVCCFLAGS: -std=c++11 -O3 -g 
I LDFLAGS:   -L/usr/lib -Wl,-rpath=/usr/lib -L/usr/lib64 -Wl,-rpath=/usr/lib64 -lhipblas -lamdhip64 -lrocblas 
I CC:        Debian clang version 19.1.3 (2)
I CXX:       Debian clang version 19.1.3 (2)

Skipping a bunch of successfully compiled modules, we get to the command that fails:

/usr/bin/hipcc -std=c++11 -fPIC -O3 -g -Wall -Wextra -Wpedantic -Wcast-qual -Wno-unused-function -Wmissing-declarations -Wmissing-noreturn -pthread  -march=native -mtune=native -Wunreachable-code-break -Wunreachable-code-return -Wmissing-prototypes -Wextra-semi -Iggml/include -Iggml/src -Iinclude -Isrc -Icommon -D_XOPEN_SOURCE=600 -D_GNU_SOURCE -DNDEBUG -DGGML_USE_CPU_AARCH64 -DGGML_USE_LLAMAFILE -DGGML_USE_AMX -DGGML_USE_HIP -DGGML_USE_CUDA  --offload-arch=gfx803 -x hip -c -o ggml/src/ggml-cuda/ggml-cuda.o ggml/src/ggml-cuda/ggml-cuda.cu

I think these are mostly noise with regards to my compilation issue, but please correct me if I'm missing a signal here:

In file included from ggml/src/ggml-cuda/ggml-cuda.cu:5:
In file included from ggml/src/ggml-cuda/common.cuh:20:
ggml/src/ggml-common.h:154:9: warning: anonymous types declared in an anonymous union are an extension [-Wnested-anon-types]
  154 |         struct {
      |         ^
ggml/src/ggml-common.h:175:9: warning: anonymous types declared in an anonymous union are an extension [-Wnested-anon-types]
  175 |         struct {
      |         ^
ggml/src/ggml-common.h:196:9: warning: anonymous types declared in an anonymous union are an extension [-Wnested-anon-types]
  196 |         struct {
      |         ^
ggml/src/ggml-common.h:261:9: warning: anonymous types declared in an anonymous union are an extension [-Wnested-anon-types]
  261 |         struct {
      |         ^
ggml/src/ggml-common.h:288:9: warning: anonymous types declared in an anonymous union are an extension [-Wnested-anon-types]
  288 |         struct {
      |         ^
ggml/src/ggml-common.h:305:9: warning: anonymous types declared in an anonymous union are an extension [-Wnested-anon-types]
  305 |         struct {
      |         ^
In file included from ggml/src/ggml-cuda/ggml-cuda.cu:22:
In file included from ggml/src/ggml-cuda/mmq.cuh:5:
ggml/src/ggml-cuda/mma.cuh:172:105: warning: function 'mma_K4' could be declared with attribute 'noreturn' [-Wmissing-noreturn]
  172 |     __device__ __forceinline__ void mma_K4(const mma_int_A_I16K4 & mma_A, const mma_int_B_J8K4 & mma_B) {
      |                                                                                                         ^
ggml/src/ggml-cuda/mma.cuh:194:105: warning: function 'mma_K8' could be declared with attribute 'noreturn' [-Wmissing-noreturn]
  194 |     __device__ __forceinline__ void mma_K8(const mma_int_A_I16K8 & mma_A, const mma_int_B_J8K8 & mma_B) {
      |                                                                                                         ^
In file included from ggml/src/ggml-cuda/ggml-cuda.cu:22:
ggml/src/ggml-cuda/mmq.cuh:866:99: warning: unused parameter 'k00' [-Wunused-parameter]
  866 |     const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int & k00) {
      |                                                                                                   ^
ggml/src/ggml-cuda/mmq.cuh:1053:99: warning: unused parameter 'k00' [-Wunused-parameter]
 1053 |     const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int & k00) {
      |                                                                                                   ^
ggml/src/ggml-cuda/mmq.cuh:1704:99: warning: unused parameter 'k00' [-Wunused-parameter]
 1704 |     const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int & k00) {
      |                                                                                                   ^
ggml/src/ggml-cuda/mmq.cuh:2497:17: warning: unused parameter 'ne00' [-Wunused-parameter]
 2497 |     const int & ne00, const int & ne01, const int & stride01, const int & ne10, const int & ne11, const int & stride11, const int & ne0,
      |                 ^
ggml/src/ggml-cuda/mmq.cuh:2497:75: warning: unused parameter 'ne10' [-Wunused-parameter]
 2497 |     const int & ne00, const int & ne01, const int & stride01, const int & ne10, const int & ne11, const int & stride11, const int & ne0,
      |                                                                           ^
ggml/src/ggml-cuda/mmq.cuh:2821:15: warning: unused variable 'nsm' [-Wunused-variable]
 2821 |     const int nsm   = ggml_cuda_info().devices[id].nsm;
      |               ^~~
In file included from ggml/src/ggml-cuda/ggml-cuda.cu:1:
In file included from ggml/include/ggml-cuda.h:3:
ggml/include/ggml.h:584:12: warning: 'backend' is deprecated: use the buffer type to find the storage location of the tensor [-Wdeprecated-declarations]
  584 |     struct ggml_tensor {
      |            ^
ggml/src/ggml-cuda/ggml-cuda.cu:1855:28: note: in implicit copy constructor for 'ggml_tensor' first required here
 1855 |     ggml_tensor src0_row = *src0;
      |                            ^
ggml/include/ggml.h:587:9: note: 'backend' has been explicitly marked deprecated here
  587 |         GGML_DEPRECATED(enum ggml_backend_type backend, "use the buffer type to find the storage location of the tensor");
      |         ^
ggml/include/ggml.h:192:61: note: expanded from macro 'GGML_DEPRECATED'
  192 | #    define GGML_DEPRECATED(func, hint) func __attribute__((deprecated(hint)))
      |                                                             ^
ggml/src/ggml-cuda/ggml-cuda.cu:2670:52: warning: unused parameter 'buffer' [-Wunused-parameter]
 2670 | bool ggml_backend_cuda_register_host_buffer(void * buffer, size_t size) {
      |                                                    ^
ggml/src/ggml-cuda/ggml-cuda.cu:2670:67: warning: unused parameter 'size' [-Wunused-parameter]
 2670 | bool ggml_backend_cuda_register_host_buffer(void * buffer, size_t size) {
      |                                                                   ^
ggml/src/ggml-cuda/ggml-cuda.cu:2966:13: warning: 'break' will never be executed [-Wunreachable-code-break]
 2966 |             break;
      |             ^~~~~
ggml/src/ggml-cuda/ggml-cuda.cu:2962:15: warning: 'break' will never be executed [-Wunreachable-code-break]
 2962 |             } break;
      |               ^~~~~
ggml/src/ggml-cuda/ggml-cuda.cu:2953:15: warning: 'break' will never be executed [-Wunreachable-code-break]
 2953 |             } break;
      |               ^~~~~
ggml/src/ggml-cuda/ggml-cuda.cu:2946:15: warning: 'break' will never be executed [-Wunreachable-code-break]
 2946 |             } break;
      |               ^~~~~
ggml/src/ggml-cuda/ggml-cuda.cu:2941:15: warning: 'break' will never be executed [-Wunreachable-code-break]
 2941 |             } break;
      |               ^~~~~
ggml/src/ggml-cuda/ggml-cuda.cu:2936:15: warning: 'break' will never be executed [-Wunreachable-code-break]
 2936 |             } break;
      |               ^~~~~
ggml/src/ggml-cuda/ggml-cuda.cu:2931:15: warning: 'break' will never be executed [-Wunreachable-code-break]
 2931 |             } break;
      |               ^~~~~
ggml/src/ggml-cuda/ggml-cuda.cu:2889:15: warning: 'break' will never be executed [-Wunreachable-code-break]
 2889 |             } break;
      |               ^~~~~
ggml/src/ggml-cuda/ggml-cuda.cu:2872:15: warning: 'break' will never be executed [-Wunreachable-code-break]
 2872 |             } break;
      |               ^~~~~
ggml/src/ggml-cuda/ggml-cuda.cu:2811:13: warning: 'break' will never be executed [-Wunreachable-code-break]
 2811 |             break;
      |             ^~~~~
27 warnings generated when compiling for gfx803.

Finally, the link error:

lld: error: unknown argument '--lto-CGO3'
clang++-17: error: amdgcn-link command failed with exit code 1 (use -v to see invocation)
make: *** [Makefile:800: ggml/src/ggml-cuda/ggml-cuda.o] Error 1

This error puzzles me for a few reasons:

• I pointed the built at clang v19 and it confirmed it would use it but this error is about lld v17.
• It seems the --lto-CGO# option shipped with lld v17

My best theory is that the older hipcc toolchain shipped with Debian (v5.7 vs current 6.3) is built on clang-17 and perhaps built with options that disabled this specific LTO feature. I'm not sure where this amdgcn-link command is coming from, so I'm guessing it is a pseudo name. There is a directory full of bitcode files /usr/lib/llvm-17/lib/clang/17/amdgcn from the rocm-device-libs-17 debian package. Perhaps hipcc 5.7 is sufficient but it requires a newer rocm-device-libs pacakge? I'd appreciate guidance from anyone more familiar with building for older AMD hardware. Thanks!

[oscarbg]

where is hipcc 6.3?