[Feature]: Jax support for gfx1151 #234
Description
Problem Description
The previous version (0.7.1) caused a segmentation fault (#132 (comment))
I thought I'd check to see if there was any change in 0.8.0, and it does give a completely different error:
F1221 14:39:47.105755 12242 stream_executor_util.cc:517] Could not load RepeatBufferKernel: INTERNAL: Failed call to hipGetFuncBySymbol: hipError_t(98)
*** Check failure stack trace: ***
@ 0x7fdf818f2e74 absl::lts_20250814::log_internal::LogMessage::SendToLog()
@ 0x7fdf818f2df6 absl::lts_20250814::log_internal::LogMessage::Flush()
@ 0x7fdf7ccca963 xla::primitive_util::FloatingPointTypeSwitch<>()
@ 0x7fdf7ccc9632 xla::gpu::InitializeBuffer()
@ 0x7fdf7ccc526b stream_executor::RedzoneAllocator::CreateBuffer()
@ 0x7fdf7549dd7d xla::gpu::RedzoneBuffers::CreateInputs()
@ 0x7fdf7549d8c7 xla::gpu::RedzoneBuffers::FromProgramShape()
@ 0x7fdf7549d506 xla::gpu::RedzoneBuffers::FromInstruction()
@ 0x7fdf7547dbd0 xla::gpu::GemmFusionAutotunerImpl::MeasurePerformance()
@ 0x7fdf7547ec07 xla::gpu::GemmFusionAutotunerImpl::Profile()
@ 0x7fdf7547fd49 xla::gpu::GemmFusionAutotunerImpl::Autotune()
@ 0x7fdf754833c1 xla::gpu::GemmFusionAutotuner::Run()
@ 0x7fdf7677b80d xla::HloPassPipeline::RunHelper<>()
@ 0x7fdf767789d7 xla::HloPassPipeline::RunPassesInternal<>()
@ 0x7fdf767782be xla::HloPassPipeline::Run()
@ 0x7fdf75168cce xla::gpu::GpuCompiler::OptimizeHloPostLayoutAssignment()
@ 0x7fdf751503b6 xla::gpu::AMDGPUCompiler::OptimizeHloPostLayoutAssignment()
@ 0x7fdf75160353 xla::gpu::GpuCompiler::OptimizeHloModule()
@ 0x7fdf7516bd4e xla::gpu::GpuCompiler::RunHloPasses()
@ 0x7fdf7513ab96 xla::Service::BuildExecutable()
@ 0x7fdf75109d65 xla::LocalService::CompileExecutables()
@ 0x7fdf75104b24 xla::LocalClient::Compile()
@ 0x7fdf75061fa6 xla::PjRtStreamExecutorClient::CompileInternal()
@ 0x7fdf750633e6 xla::PjRtStreamExecutorClient::Compile()
@ 0x7fdf75063ee5 xla::PjRtStreamExecutorClient::CompileAndLoad()
@ 0x7fdf7502527d xla::StreamExecutorGpuClient::CompileAndLoad()
@ 0x7fdf75001343 pjrt::PJRT_Client_Compile()
@ 0x7fe02cd82f2f xla::InitializeArgsAndCompile()
@ 0x7fe02cd83404 xla::PjRtCApiClient::CompileAndLoad()
@ 0x7fe02968e591 xla::ifrt::PjRtLoadedExecutable::Create()
@ 0x7fe029688f27 xla::ifrt::PjRtCompiler::CompileAndLoad()
@ 0x7fe02953ea9f jax::PyClient::CompileAndLoadIfrtProgram()
@ 0x7fe0295408bd jax::PyClient::CompileAndLoad()
@ 0x7fe02954d289 nanobind::detail::func_create<>()::{lambda()#1}::__invoke()
@ 0x7fe029613788 nanobind::detail::nb_func_vectorcall_complex()
@ 0x7fe0296140dc nanobind::detail::nb_bound_method_vectorcall()
@ 0x559a1fadf03e PyObject_Vectorcall
/tmp/spyder-user/tmp8m87oh4y: line 3: 12242 Aborted (core dumped) /home/user/anaconda3/envs/jax/bin/python -Xfrozen_modules=off -m spyder_kernels.console -f /home/user/.local/share/jupyter/runtime/kernel-c7b9db4ad6dc.json
The kernel died, restarting...
Fatal Python error: Aborted
Thread 0x00007fe03ffff6c0 [Thread-3] (most recent call first):
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/zmq/utils/garbage.py", line 46 in run
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1082 in _bootstrap_inner
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1044 in _bootstrap
Thread 0x00007fe058a856c0 [Thread-1] (most recent call first):
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/spyder_kernels/console/kernelapp.py", line 52 in run
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1082 in _bootstrap_inner
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1044 in _bootstrap
Thread 0x00007fe0593866c0 [IPythonHistoryS] (most recent call first):
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 369 in wait
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 670 in wait
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/IPython/core/history.py", line 1110 in run
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/IPython/core/history.py", line 98 in only_when_enabled
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/decorator.py", line 235 in fun
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1082 in _bootstrap_inner
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1044 in _bootstrap
Thread 0x00007fe05a0a96c0 [Control] (most recent call first):
File "/home/user/anaconda3/envs/jax/lib/python3.14/selectors.py", line 452 in select
File "/home/user/anaconda3/envs/jax/lib/python3.14/asyncio/base_events.py", line 2008 in _run_once
File "/home/user/anaconda3/envs/jax/lib/python3.14/asyncio/base_events.py", line 677 in run_forever
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/tornado/platform/asyncio.py", line 211 in start
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/control.py", line 23 in run
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1082 in _bootstrap_inner
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1044 in _bootstrap
Thread 0x00007fe05b8ac6c0 [Heartbeat] (most recent call first):
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/zmq/sugar/__init__.py", line 21 in device
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/heartbeat.py", line 106 in run
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1082 in _bootstrap_inner
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1044 in _bootstrap
Thread 0x00007fe05c0ad6c0 [IOPub] (most recent call first):
File "/home/user/anaconda3/envs/jax/lib/python3.14/selectors.py", line 452 in select
File "/home/user/anaconda3/envs/jax/lib/python3.14/asyncio/base_events.py", line 2008 in _run_once
File "/home/user/anaconda3/envs/jax/lib/python3.14/asyncio/base_events.py", line 677 in run_forever
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/tornado/platform/asyncio.py", line 211 in start
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/iostream.py", line 92 in _thread_main
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1024 in run
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1082 in _bootstrap_inner
File "/home/user/anaconda3/envs/jax/lib/python3.14/threading.py", line 1044 in _bootstrap
Current thread 0x00007fe0c2539780 [python] (most recent call first):
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/compiler.py", line 375 in backend_compile_and_load
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/profiler.py", line 359 in wrapper
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/compiler.py", line 759 in _compile_and_write_cache
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/compiler.py", line 491 in compile_or_get_cached
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/interpreters/pxla.py", line 2854 in _cached_compilation
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/interpreters/pxla.py", line 3073 in from_hlo
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/interpreters/pxla.py", line 2527 in compile
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/pjit.py", line 1600 in _pjit_call_impl_python
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/pjit.py", line 146 in _python_pjit_helper
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/pjit.py", line 263 in cache_miss
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax/_src/traceback_util.py", line 180 in reraise_with_filtered_traceback
File "/home/user/Documents/Python Bits/jax_test.py", line 40 in <module>
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/spyder_kernels/customize/utils.py", line 209 in exec_encapsulate_locals
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/spyder_kernels/customize/code_runner.py", line 640 in _exec_code
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/spyder_kernels/customize/code_runner.py", line 470 in _exec_file
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/spyder_kernels/customize/code_runner.py", line 177 in runfile
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/IPython/core/interactiveshell.py", line 2504 in run_line_magic
File "/tmp/ipykernel_12242/3280664694.py", line 1 in <module>
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/IPython/core/interactiveshell.py", line 3699 in run_code
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/spyder_kernels/console/shell.py", line 461 in run_code
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/IPython/core/interactiveshell.py", line 3639 in run_ast_nodes
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/IPython/core/interactiveshell.py", line 3394 in run_cell_async
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/IPython/core/async_helpers.py", line 128 in _pseudo_sync_runner
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/IPython/core/interactiveshell.py", line 3171 in _run_cell
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/IPython/core/interactiveshell.py", line 3116 in run_cell
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/zmqshell.py", line 602 in run_cell
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/ipkernel.py", line 455 in do_execute
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/kernelbase.py", line 767 in execute_request
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/ipkernel.py", line 368 in execute_request
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/kernelbase.py", line 400 in dispatch_shell
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/kernelbase.py", line 508 in process_one
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/kernelbase.py", line 519 in dispatch_queue
File "/home/user/anaconda3/envs/jax/lib/python3.14/asyncio/events.py", line 94 in _run
File "/home/user/anaconda3/envs/jax/lib/python3.14/asyncio/base_events.py", line 2046 in _run_once
File "/home/user/anaconda3/envs/jax/lib/python3.14/asyncio/base_events.py", line 677 in run_forever
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/tornado/platform/asyncio.py", line 211 in start
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/ipykernel/kernelapp.py", line 739 in start
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/spyder_kernels/console/start.py", line 189 in main
File "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/spyder_kernels/console/__main__.py", line 24 in <module>
File "/home/user/anaconda3/envs/jax/lib/python3.14/runpy.py", line 88 in _run_code
File "/home/user/anaconda3/envs/jax/lib/python3.14/runpy.py", line 198 in _run_module_as_main
Current thread's C stack trace (most recent call first):
Binary file "/home/user/anaconda3/envs/jax/bin/python", at _Py_DumpStack+0x4a [0x559a1fa9c921]
Binary file "/home/user/anaconda3/envs/jax/bin/python", at +0x195d8f [0x559a1fa9cd8f]
Binary file "/lib64/libc.so.6", at +0x1a290 [0x7fe0c2558290]
Binary file "/lib64/libc.so.6", at +0x743cc [0x7fe0c25b23cc]
Binary file "/lib64/libc.so.6", at gsignal+0x1e [0x7fe0c255815e]
Binary file "/lib64/libc.so.6", at abort+0x26 [0x7fe0c253f6d0]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xd2f2ab6 [0x7fdf818f2ab6]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xd2f2ed4 [0x7fdf818f2ed4]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xd2f2df6 [0x7fdf818f2df6]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0x86ca963 [0x7fdf7ccca963]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0x86c9632 [0x7fdf7ccc9632]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0x86c526b [0x7fdf7ccc526b]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xe9dd7d [0x7fdf7549dd7d]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xe9d8c7 [0x7fdf7549d8c7]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xe9d506 [0x7fdf7549d506]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xe7dbd0 [0x7fdf7547dbd0]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xe7ec07 [0x7fdf7547ec07]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xe7fd49 [0x7fdf7547fd49]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xe833c1 [0x7fdf754833c1]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0x217b80d [0x7fdf7677b80d]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0x21789d7 [0x7fdf767789d7]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0x21782be [0x7fdf767782be]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xb68cce [0x7fdf75168cce]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xb503b6 [0x7fdf751503b6]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xb60353 [0x7fdf75160353]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xb6bd4e [0x7fdf7516bd4e]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xb3ab96 [0x7fdf7513ab96]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xb09d65 [0x7fdf75109d65]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xb04b24 [0x7fdf75104b24]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xa61fa6 [0x7fdf75061fa6]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xa633e6 [0x7fdf750633e6]
Binary file "/home/user/anaconda3/envs/jax/lib/python3.14/site-packages/jax_plugins/xla_rocm7/xla_rocm_plugin.so", at +0xa63ee5 [0x7fdf75063ee5]
<truncated rest of calls>
Extension modules: zmq.backend.cython._zmq, tornado.speedups, psutil._psutil_linux, numpy._core._multiarray_umath, numpy.linalg._umath_linalg, PIL._imaging, kiwisolver._cext, _pydevd_bundle.pydevd_cython, _pydevd_sys_monitoring_cython, _pydevd_sys_monitoring._pydevd_sys_monitoring_cython, jaxlib.cpu_feature_guard, numpy.random._common, numpy.random.bit_generator, numpy.random._bounded_integers, numpy.random._pcg64, numpy.random._mt19937, numpy.random._generator, numpy.random._philox, numpy.random._sfc64, numpy.random.mtrand (total: 20)
Just for a laugh I followed this PR: #37 and added gfx1151 to the targets. The performance is quite bad, but it doesn't crash.
Operating System
Fedora 43
CPU
AMD RYZEN AI MAX+ 395
GPU
Radeon 8060S
ROCm Version
ROCm 7.1.1
ROCm Component
No response
Steps to Reproduce
import jax.numpy as jnp
N = 4096
A = jnp.ones((N, N))
B = jnp.ones((N, N))
C = jnp.matmul(A, B)
(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support
ROCk module version 6.16.6 is loaded
HSA System Attributes
Runtime Version: 1.18
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 AI MAX+ 395 w/ Radeon 8060S
Uuid: CPU-XX
Marketing Name: AMD RYZEN AI MAX+ 395 w/ Radeon 8060S
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: 49152(0xc000) KB
Chip ID: 0(0x0)
ASIC Revision: 0(0x0)
Cacheline Size: 64(0x40)
Max Clock Freq. (MHz): 5187
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: 131151704(0x7d13758) 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: 131151704(0x7d13758) 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: 131151704(0x7d13758) 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: 131151704(0x7d13758) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Recommended Granule:4KB
Alloc Alignment: 4KB
Accessible by all: TRUE
ISA Info:
Agent 2
Name: gfx1151
Uuid: GPU-XX
Marketing Name: Radeon 8060S Graphics
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: 32(0x20) KB
L2: 2048(0x800) KB
L3: 32768(0x8000) KB
Chip ID: 5510(0x1586)
ASIC Revision: 0(0x0)
Cacheline Size: 128(0x80)
Max Clock Freq. (MHz): 2900
BDFID: 49664
Internal Node ID: 1
Compute Unit: 40
SIMDs per CU: 2
Shader Engines: 2
Shader Arrs. per Eng.: 2
WatchPts on Addr. Ranges:4
Coherent Host Access: FALSE
Memory Properties: APU
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: 32(0x20)
Max Work-item Per CU: 1024(0x400)
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:: 32
SDMA engine uCode:: 17
IOMMU Support:: None
Pool Info:
Pool 1
Segment: GLOBAL; FLAGS: COARSE GRAINED
Size: 65575852(0x3e89bac) 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: 65575852(0x3e89bac) 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--gfx1151
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--gfx11-generic
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
Agent 3
Name: aie2
Uuid: AIE-XX
Marketing Name: AIE-ML
Vendor Name: AMD
Feature: AGENT_DISPATCH
Profile: BASE_PROFILE
Float Round Mode: NEAR
Max Queue Number: 1(0x1)
Queue Min Size: 64(0x40)
Queue Max Size: 64(0x40)
Queue Type: SINGLE
Node: 0
Device Type: DSP
Cache Info:
L2: 2048(0x800) KB
L3: 32768(0x8000) KB
Chip ID: 0(0x0)
ASIC Revision: 0(0x0)
Cacheline Size: 0(0x0)
Max Clock Freq. (MHz): 0
BDFID: 0
Internal Node ID: 0
Compute Unit: 0
SIMDs per CU: 0
Shader Engines: 0
Shader Arrs. per Eng.: 0
WatchPts on Addr. Ranges:0
Memory Properties:
Features: AGENT_DISPATCH
Pool Info:
Pool 1
Segment: GLOBAL; FLAGS: KERNARG, COARSE GRAINED
Size: 131151704(0x7d13758) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Recommended Granule:4KB
Alloc Alignment: 4KB
Accessible by all: TRUE
Pool 2
Segment: GLOBAL; FLAGS: COARSE GRAINED
Size: 65536(0x10000) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Recommended Granule:0KB
Alloc Alignment: 4KB
Accessible by all: TRUE
Pool 3
Segment: GLOBAL; FLAGS: COARSE GRAINED
Size: 131151704(0x7d13758) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Recommended Granule:4KB
Alloc Alignment: 4KB
Accessible by all: TRUE
ISA Info:
*** Done ***
Additional Information
OS:
NAME="Fedora Linux"
VERSION="43 (Workstation Edition)"
CPU:
model name : AMD RYZEN AI MAX+ 395 w/ Radeon 8060S
GPU:
Name: AMD RYZEN AI MAX+ 395 w/ Radeon 8060S
Marketing Name: AMD RYZEN AI MAX+ 395 w/ Radeon 8060S
Name: gfx1151
Marketing Name: Radeon 8060S Graphics
Name: amdgcn-amd-amdhsa--gfx1151
Name: amdgcn-amd-amdhsa--gfx11-generic
Name: aie2
Marketing Name: AIE-ML
Metadata
Metadata
Assignees
Type
Projects
Status