[device-libs]: Compile failure on amd/device-libs/ockl/src/cluster.cl with LLVM 21 on Gentoo

[mattmcadoo]

Problem Description

OS:
NAME='Gentoo'
VERSION='2.18'
CPU:
model name : AMD Ryzen Threadripper PRO 3955WX 16-Cores
GPU:
Name: AMD Ryzen Threadripper PRO 3955WX 16-Cores
Marketing Name: AMD Ryzen Threadripper PRO 3955WX 16-Cores
Name: gfx1012
Marketing Name: AMD Radeon Pro W5500
Name: amdgcn-amd-amdhsa--gfx1012:xnack-
Name: amdgcn-amd-amdhsa--gfx10-1-generic:xnack-

FAILED: [code=1] ockl/cluster.bc /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs_build/ockl/cluster.bc
cd /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs_build/ockl && /usr/lib/llvm/21/bin/clang-21 -I/var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/../irif/inc -I/var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/../oclc/inc -I/var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/inc -fcolor-diagnostics -Werror -Wno-error=atomic-alignment -x cl -Xclang -cl-std=CL2.0 -target amdgcn-amd-amdhsa -fvisibility=hidden -fomit-frame-pointer -Xclang -finclude-default-header -Xclang -fexperimental-strict-floating-point -Xclang -fdenormal-fp-math=dynamic -nogpulib -cl-no-stdinc -Xclang -mcode-object-version=none -emit-llvm -c /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/src/cluster.cl -o /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs_build/ockl/cluster.bc -MD -MF /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs_build/ockl/cluster.cl.d && /usr/bin/cmake -E cmake_transform_depfile Ninja gccdepfile /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs_build /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs_build/ockl /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs_build/ockl/cluster.cl.d /var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs_build/CMakeFiles/d/9cb277e447a0840f639afb4d76af9594233b317d09f80da5e7f68083027135ce.d
/var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/src/cluster.cl:16:16: error: use of undeclared identifier '__builtin_amdgcn_cluster_workgroup_max_id_x'
   16 |         return __builtin_amdgcn_cluster_workgroup_max_id_x() + 1;
      |                ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/src/cluster.cl:18:16: error: use of undeclared identifier '__builtin_amdgcn_cluster_workgroup_max_id_y'
   18 |         return __builtin_amdgcn_cluster_workgroup_max_id_y() + 1;
      |                ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/src/cluster.cl:20:16: error: use of undeclared identifier '__builtin_amdgcn_cluster_workgroup_max_id_z'
   20 |         return __builtin_amdgcn_cluster_workgroup_max_id_z() + 1;
      |                ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/src/cluster.cl:31:16: error: use of undeclared identifier '__builtin_amdgcn_cluster_workgroup_id_x'
   31 |         return __builtin_amdgcn_cluster_workgroup_id_x();
      |                ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/src/cluster.cl:33:16: error: use of undeclared identifier '__builtin_amdgcn_cluster_workgroup_id_y'
   33 |         return __builtin_amdgcn_cluster_workgroup_id_y();
      |                ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/src/cluster.cl:35:16: error: use of undeclared identifier '__builtin_amdgcn_cluster_workgroup_id_z'
   35 |         return __builtin_amdgcn_cluster_workgroup_id_z();
      |                ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/var/tmp/portage/dev-libs/rocm-device-libs-7.2.0/work/llvm-project-rocm-7.2.0/amd/device-libs/ockl/src/cluster.cl:44:12: error: use of undeclared identifier '__builtin_amdgcn_cluster_workgroup_max_flat_id'
   44 |     return __builtin_amdgcn_cluster_workgroup_max_flat_id() + 1;
      |            ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
7 errors generated.

Operating System

NAME='Gentoo' VERSION='2.18'

CPU

model name : AMD Ryzen Threadripper PRO 3955WX 16-Cores

GPU

Name: gfx1012 Marketing Name: AMD Radeon Pro W5500 Name: amdgcn-amd-amdhsa--gfx1012:xnack- Name: amdgcn-amd-amdhsa--gfx10-1-generic:xnack-

ROCm Version

7.2.0

ROCm Component

No response

Steps to Reproduce

No response

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

# rocminfo --support
=====================
HSA System Attributes
=====================
Runtime Version:         1.1
Runtime Ext Version:     1.14
System Timestamp Freq.:  1000.000000MHz
Sig. Max Wait Duration:  18446744073709551615 (0xFFFFFFFFFFFFFFFF) (timestamp count)
Machine Model:           LARGE
System Endianness:       LITTLE
Mwaitx:                  DISABLED
XNACK enabled:           NO
DMAbuf Support:          YES
VMM Support:             YES

==========
HSA Agents
==========
*******
Agent 1
*******
  Name:                    AMD Ryzen Threadripper PRO 3955WX 16-Cores
  Uuid:                    CPU-XX
  Marketing Name:          AMD Ryzen Threadripper PRO 3955WX 16-Cores
  Vendor Name:             CPU
  Feature:                 None specified
  Profile:                 FULL_PROFILE
  Float Round Mode:        NEAR
  Max Queue Number:        0(0x0)
  Queue Min Size:          0(0x0)
  Queue Max Size:          0(0x0)
  Queue Type:              MULTI
  Node:                    0
  Device Type:             CPU
  Cache Info:
    L1:                      32768(0x8000) KB
  Chip ID:                 0(0x0)
  ASIC Revision:           0(0x0)
  Cacheline Size:          64(0x40)
  Max Clock Freq. (MHz):   4405
  BDFID:                   0
  Internal Node ID:        0
  Compute Unit:            32
  SIMDs per CU:            0
  Shader Engines:          0
  Shader Arrs. per Eng.:   0
  WatchPts on Addr. Ranges:1
  Memory Properties:
  Features:                None
  Pool Info:
    Pool 1
      Segment:                 GLOBAL; FLAGS: FINE GRAINED
      Size:                    395800364(0x17976f2c) KB
      Allocatable:             TRUE
      Alloc Granule:           4KB
      Alloc Recommended Granule:4KB
      Alloc Alignment:         4KB
      Accessible by all:       TRUE
    Pool 2
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    395800364(0x17976f2c) KB
      Allocatable:             TRUE
      Alloc Granule:           4KB
      Alloc Recommended Granule:4KB
      Alloc Alignment:         4KB
      Accessible by all:       TRUE
    Pool 3
      Segment:                 GLOBAL; FLAGS: KERNARG, FINE GRAINED
      Size:                    395800364(0x17976f2c) KB
      Allocatable:             TRUE
      Alloc Granule:           4KB
      Alloc Recommended Granule:4KB
      Alloc Alignment:         4KB
      Accessible by all:       TRUE
    Pool 4
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED
      Size:                    395800364(0x17976f2c) KB
      Allocatable:             TRUE
      Alloc Granule:           4KB
      Alloc Recommended Granule:4KB
      Alloc Alignment:         4KB
      Accessible by all:       TRUE
  ISA Info:
*******
Agent 2
*******
  Name:                    gfx1012
  Uuid:                    GPU-XX
  Marketing Name:          AMD Radeon Pro W5500
  Vendor Name:             AMD
  Feature:                 KERNEL_DISPATCH
  Profile:                 BASE_PROFILE
  Float Round Mode:        NEAR
  Max Queue Number:        128(0x80)
  Queue Min Size:          64(0x40)
  Queue Max Size:          131072(0x20000)
  Queue Type:              MULTI
  Node:                    1
  Device Type:             GPU
  Cache Info:
    L1:                      16(0x10) KB
    L2:                      2048(0x800) KB
  Chip ID:                 29505(0x7341)
  ASIC Revision:           1(0x1)
  Cacheline Size:          128(0x80)
  Max Clock Freq. (MHz):   1900
  BDFID:                   25344
  Internal Node ID:        1
  Compute Unit:            22
  SIMDs per CU:            2
  Shader Engines:          1
  Shader Arrs. per Eng.:   2
  WatchPts on Addr. Ranges:4
  Coherent Host Access:    FALSE
  Memory Properties:
  Features:                KERNEL_DISPATCH
  Fast F16 Operation:      TRUE
  Wavefront Size:          32(0x20)
  Workgroup Max Size:      1024(0x400)
  Workgroup Max Size per Dimension:
    x                        1024(0x400)
    y                        1024(0x400)
    z                        1024(0x400)
  Max Waves Per CU:        40(0x28)
  Max Work-item Per CU:    1280(0x500)
  Grid Max Size:           4294967295(0xffffffff)
  Grid Max Size per Dimension:
    x                        2147483647(0x7fffffff)
    y                        65535(0xffff)
    z                        65535(0xffff)
  Max fbarriers/Workgrp:   32
  Packet Processor uCode:: 123
  SDMA engine uCode::      41
  IOMMU Support::          None
  Pool Info:
    Pool 1
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED
      Size:                    8372224(0x7fc000) KB
      Allocatable:             TRUE
      Alloc Granule:           4KB
      Alloc Recommended Granule:2048KB
      Alloc Alignment:         4KB
      Accessible by all:       FALSE
    Pool 2
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    8372224(0x7fc000) KB
      Allocatable:             TRUE
      Alloc Granule:           4KB
      Alloc Recommended Granule:2048KB
      Alloc Alignment:         4KB
      Accessible by all:       FALSE
    Pool 3
      Segment:                 GROUP
      Size:                    64(0x40) KB
      Allocatable:             FALSE
      Alloc Granule:           0KB
      Alloc Recommended Granule:0KB
      Alloc Alignment:         0KB
      Accessible by all:       FALSE
  ISA Info:
    ISA 1
      Name:                    amdgcn-amd-amdhsa--gfx1012:xnack-
      Machine Models:          HSA_MACHINE_MODEL_LARGE
      Profiles:                HSA_PROFILE_BASE
      Default Rounding Mode:   NEAR
      Default Rounding Mode:   NEAR
      Fast f16:                TRUE
      Workgroup Max Size:      1024(0x400)
      Workgroup Max Size per Dimension:
        x                        1024(0x400)
        y                        1024(0x400)
        z                        1024(0x400)
      Grid Max Size:           4294967295(0xffffffff)
      Grid Max Size per Dimension:
        x                        2147483647(0x7fffffff)
        y                        65535(0xffff)
        z                        65535(0xffff)
      FBarrier Max Size:       32
    ISA 2
      Name:                    amdgcn-amd-amdhsa--gfx10-1-generic:xnack-
      Machine Models:          HSA_MACHINE_MODEL_LARGE
      Profiles:                HSA_PROFILE_BASE
      Default Rounding Mode:   NEAR
      Default Rounding Mode:   NEAR
      Fast f16:                TRUE
      Workgroup Max Size:      1024(0x400)
      Workgroup Max Size per Dimension:
        x                        1024(0x400)
        y                        1024(0x400)
        z                        1024(0x400)
      Grid Max Size:           4294967295(0xffffffff)
      Grid Max Size per Dimension:
        x                        2147483647(0x7fffffff)
        y                        65535(0xffff)
        z                        65535(0xffff)
      FBarrier Max Size:       32
*** Done ***

Additional Information

This is the commit that added that file: fcc50fb

[zichguan-amd]

Hi @mattmcadoo, can you provide the exact build steps you used?

[mattmcadoo]

I use Gentoo, using LLVM 21 and this is the ebuild I'm using:

cat /var/db/repos/local-grotto/dev-libs/rocm-device-libs/rocm-device-libs-7.2.0.ebuild
# Copyright 1999-2025 Gentoo Authors
# Distributed under the terms of the GNU General Public License v2

EAPI=8

LLVM_COMPAT=( 21 )
inherit cmake flag-o-matic llvm-r1

MY_P=llvm-project-rocm-${PV}
components=( "amd/device-libs" )

if [[ ${PV} == *9999 ]] ; then
	EGIT_REPO_URI="https://github.com/ROCm/llvm-project"
	inherit git-r3
	S="${WORKDIR}/${P}/${components[0]}"
else
	SRC_URI="https://github.com/ROCm/llvm-project/archive/rocm-${PV}.tar.gz -> ${MY_P}.tar.gz"
	S="${WORKDIR}/${MY_P}/${components[0]}"
	KEYWORDS="~amd64"
fi

DESCRIPTION="Radeon Open Compute Device Libraries"
HOMEPAGE="https://github.com/ROCm/llvm-project/tree/amd-staging/amd/device-libs"

LICENSE="MIT"
SLOT="0/$(ver_cut 1-2)"
IUSE="test"
RESTRICT="!test? ( test )"

BDEPEND="
	dev-build/rocm-cmake
	$(llvm_gen_dep "
		llvm-core/clang:\${LLVM_SLOT}
		llvm-core/lld:\${LLVM_SLOT}
	")
"

CMAKE_BUILD_TYPE=Release

PATCHES=(
	"${FILESDIR}/${PN}-6.2.0-test-bitcode-dir.patch"
)

src_unpack() {
	if [[ ${PV} == *9999 ]] ; then
		git-r3_fetch
		git-r3_checkout '' . '' "${components[@]}"
	else
		archive="${MY_P}.tar.gz"
		ebegin "Unpacking from ${archive}"
		tar -x -z -o \
			-f "${DISTDIR}/${archive}" \
			"${components[@]/#/${MY_P}/}" || die
		eend ${?}
	fi
}

src_prepare() {
	sed -e "s:amdgcn/bitcode:lib/amdgcn/bitcode:" \
		-i "${S}/cmake/OCL.cmake" \
		-i "${S}/cmake/Packages.cmake" || die
	cmake_src_prepare
}

src_configure() {
	# Do not trust CMake with autoselecting Clang, as it autoselects the latest one
	# producing too modern LLVM bitcode and causing linker errors in other packages.
	llvm_prepend_path "${LLVM_SLOT}"
	local -x CC=${CHOST}-clang
	local -x CXX=${CHOST}-clang++
	# Clean up unsupported flags for the switched compiler, see #936099
	strip-unsupported-flags

	cmake_src_configure
}

src_install() {
	cmake_src_install
	# install symlink, so that clang won't ask for "--rocm-device-lib-path" flag anymore
	local bitcodedir="$(clang -print-resource-dir)/$(get_libdir)/amdgcn/bitcode"
	dosym -r "/usr/lib/amdgcn/bitcode" "${bitcodedir#${EPREFIX}}"
}

src_test() {
	# https://github.com/ROCm/llvm-project/issues/76
	# "Failing tests are on gfx that are not supported"
	local CMAKE_SKIP_TESTS=(
		compile_frexp__gfx600
		compile_fract__gfx600
		compile_fract__gfx700
		compile_atomic_work_item_fence__*
	)

	cmake_src_test
}

The patch this is being applied is this:

cat /var/db/repos/local-grotto/dev-libs/rocm-device-libs/files/rocm-device-libs-6.2.0-test-bitcode-dir.patch
`--rocm-device-lib-path` is needed to execute the tests because they are not
installed to system yet
--- a/test/compile/CMakeLists.txt
+++ b/test/compile/CMakeLists.txt
@@ -40,6 +40,7 @@ function(add_compile_test test_name func_name)
      -DFILECHECK_BIN=${FILECHECK_BIN}
      -DOUTPUT_FILE=output.${test_name}.${test_cpu}.s
      -DINPUT_FILE=${file_name}
+     -DAMDGCN_BITCODES=${PROJECT_BINARY_DIR}/lib/amdgcn/bitcode
      -DTEST_CPU=${test_cpu}
      -DCOMPILE_FLAGS=${COMPILE_TEST_COMPILE_FLAGS}
      -DEXTRA_CHECK_PREFIX=${COMPILE_TEST_EXTRA_CHECK_PREFIXES}
--- a/test/compile/RunCompileTest.cmake
+++ b/test/compile/RunCompileTest.cmake
@@ -16,6 +16,7 @@ execute_process(COMMAND
   -target amdgcn-amd-amdhsa -mcpu=${TEST_CPU}
   -Xclang -finclude-default-header
   --rocm-path=${BINARY_DIR}
+  --rocm-device-lib-path=${AMDGCN_BITCODES}
   -mllvm -amdgpu-simplify-libcall=0
   ${COMPILE_FLAGS}
   -o ${OUTPUT_FILE} ${INPUT_FILE}
--- a/test/compile/RunConstantFoldTest.cmake
+++ b/test/compile/RunConstantFoldTest.cmake
@@ -16,6 +16,7 @@ execute_process(COMMAND
   -target amdgcn-amd-amdhsa -mcpu=${TEST_CPU}
   -Xclang -finclude-default-header
   --rocm-path=${BINARY_DIR}
+  --rocm-device-lib-path=${AMDGCN_BITCODES}
   -mllvm -amdgpu-simplify-libcall=0
   -o ${OUTPUT_FILE} ${INPUT_FILE}
   RESULT_VARIABLE CLANG_RESULT

[zichguan-amd]

I see, this seems to be gentoo specific as I cannot repro it on Ubuntu. I'll try a gentoo docker.

[zichguan-amd]

Can you use amdclang and amdclang++ from ROCm to build? or build ROCm llvm with device libs. I was able to repro the issue with upstream clang 21. I think the upstream build doesn't have these target specific symbols.

[zichguan-amd]

The builtins were added by llvm@110ab5a and only available since llvm 22. I tested with upstream clang 22 and the build works as expected on Ubuntu. Can you please give that a try and let me know if it solves the issue?

[mattmcadoo]

Using LLVM 22 I've gotten different errors than before. Build log attached.

-- The C compiler identification is Clang 22.1.0
-- The CXX compiler identification is Clang 22.1.0

rocm-device-libs-7.2.0-build-with-llvm22.log.gz

[zichguan-amd]

There's a bit of desync with upstream LLVM 22 and ROCm fork at 7.2.0. llvm@52a58a4 got into upstream LLVM 22, but the associated device-libs changes didn't make it to 7.2.0. So if you want to build device-libs at 7.2.0 you need to use the ROCm fork-ed LLVM. Or you can build with upstream LLVM at device-libs amd-staging branch tot. This is the associated device-libs change bc15782