Skip to content

[Issue]: printf in a kernel doesn't work and casues the kernel to fail silently. #3794

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Closed
evstigneevnm opened this issue May 21, 2025 · 8 comments
Closed

[Issue]: printf in a kernel doesn't work and casues the kernel to fail silently. #3794

evstigneevnm opened this issue May 21, 2025 · 8 comments
Labels
Under Investigation

Comments

@evstigneevnm
Copy link

Problem Description

running the following test to use printf from the kernel, assuming ROCm/hip is installed. The source file test_printf.hip is:

#define HIP_ENABLE_PRINTF

#include "test_common.h"
#include <stdexcept>
#include <string>
#include <hip/hip_runtime.h>

#define __STR_HELPER(x) #x
#define __STR(x) __STR_HELPER(x)

#define HIP_SAFE_CALL(X)                                                                                                                                                                       \
    do {                                                                                                                                                                                        \
        hipError_t hip_res = (X);                                                                                                                                                             \
        if (hip_res != hipSuccess) throw std::runtime_error(std::string("HIP_SAFE_CALL " __FILE__ " " __STR(__LINE__) " : " #X " failed: ") + std::string(hipGetErrorString(hip_res)));    \
    } while (0)



__global__ void run_printf(int i, int* a) 
{ 
	a[0] = 4+i;
	printf("Hello World\n"); 
//	assert(a[0] == 4);
}

int main() {
    int device_count = 0;
    HIP_SAFE_CALL( hipGetDeviceCount(&device_count) );
    printf("device_count: %i\n", device_count);
    int a_h = 0;
    int *a_d;
    for (int i = 0; i < device_count; ++i) {
        hipDeviceProp_t prop;
	HIP_SAFE_CALL(hipGetDeviceProperties(&prop, i)	);
	printf("%i: %s, %lu, PCI: %i:%i:%i; %i, clock:%i, busw:%i, l2:%i, maxth:%i \n", i, prop.name, prop.totalGlobalMem, prop.pciBusID, prop.pciDeviceID, prop.pciDomainID, prop.tccDriver, prop.memoryClockRate, prop.memoryBusWidth, prop.l2CacheSize, prop.maxThreadsPerMultiProcessor );
	HIP_SAFE_CALL(hipSetDevice(i) );
	HIP_SAFE_CALL( hipMalloc((void**) &a_d, sizeof(int))  );
	HIP_SAFE_CALL( hipMemcpy(a_d, &a_h, sizeof(int), hipMemcpyHostToDevice) );
	hipLaunchKernelGGL(HIP_KERNEL_NAME(run_printf), dim3(1), dim3(1), 0, 0, i, a_d);
        HIP_SAFE_CALL(hipDeviceSynchronize() );
	HIP_SAFE_CALL(hipMemcpy(&a_h, a_d, sizeof(int), hipMemcpyDeviceToHost));
        HIP_SAFE_CALL(hipFree(a_d));
        printf("%i: %i\n", i, a_h);
    }
    passed();
}

compiled with hipcc.
execued with:

export HCC_ENABLE_PRINTF=1
./a.out

PROBLEM: uncommenting printf causes no output and the kernel is not executed, as can be checked by the output. The same goes with the assert
Output with commented lines:

device_count: 2
0: AMD Radeon™ RX 7600 XT, 17163091968, PCI: 3:0:0; 0, clock:1124000, busw:128, l2:2097152, maxth:2048 
0: 4
1: AMD Radeon™ RX 7600 XT, 17163091968, PCI: 6:0:0; 0, clock:1124000, busw:128, l2:2097152, maxth:2048 
1: 5
PASSED!

Output with printf or assert uncommented:

device_count: 2
0: AMD Radeon™ RX 7600 XT, 17163091968, PCI: 3:0:0; 0, clock:1124000, busw:128, l2:2097152, maxth:2048 
0: 0
1: AMD Radeon™ RX 7600 XT, 17163091968, PCI: 6:0:0; 0, clock:1124000, busw:128, l2:2097152, maxth:2048 
1: 0
PASSED!

Notce the output lines for the a_h variable.

What can be wrong?
Thank you for the support.

Operating System

Ubuntu 24.04.2 LTS x86_64, 6.11.0-25-generic

CPU

Intel(R) Xeon(R) CPU E5-2670 0 @ 2.60GHz

GPU

AMD Radeon™ RX 7600 XT amdgcn-amd-amdhsa--gfx1102

ROCm Version

rocm-6.3.4

ROCm Component

No response

Steps to Reproduce

No response

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

ROCk module version 6.12.12 is loaded

HSA System Attributes

Runtime Version: 1.15
Runtime Ext Version: 1.7
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: Intel(R) Xeon(R) CPU E5-2670 0 @ 2.60GHz
Uuid: CPU-XX
Marketing Name: Intel(R) Xeon(R) CPU E5-2670 0 @ 2.60GHz
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): 3300
BDFID: 0
Internal Node ID: 0
Compute Unit: 16
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: 65780988(0x3ebbcfc) 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: 65780988(0x3ebbcfc) 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: 65780988(0x3ebbcfc) 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: 65780988(0x3ebbcfc) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Recommended Granule:4KB
Alloc Alignment: 4KB
Accessible by all: TRUE
ISA Info:

Agent 2

Name: gfx1102
Uuid: GPU-XX
Marketing Name: AMD Radeon™ RX 7600 XT
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
Chip ID: 29824(0x7480)
ASIC Revision: 0(0x0)
Cacheline Size: 128(0x80)
Max Clock Freq. (MHz): 2539
BDFID: 768
Internal Node ID: 1
Compute Unit: 32
SIMDs per CU: 2
Shader Engines: 2
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: 32(0x20)
Max Work-item Per CU: 1024(0x400)
Grid Max Size: 4294967295(0xffffffff)
Grid Max Size per Dimension:
x 4294967295(0xffffffff)
y 4294967295(0xffffffff)
z 4294967295(0xffffffff)
Max fbarriers/Workgrp: 32
Packet Processor uCode:: 542
SDMA engine uCode:: 21
IOMMU Support:: None
Pool Info:
Pool 1
Segment: GLOBAL; FLAGS: COARSE GRAINED
Size: 16760832(0xffc000) 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: 16760832(0xffc000) 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--gfx1102
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 4294967295(0xffffffff)
y 4294967295(0xffffffff)
z 4294967295(0xffffffff)
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 4294967295(0xffffffff)
y 4294967295(0xffffffff)
z 4294967295(0xffffffff)
FBarrier Max Size: 32

Agent 3

Name: gfx1102
Uuid: GPU-XX
Marketing Name: AMD Radeon™ RX 7600 XT
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: 2
Device Type: GPU
Cache Info:
L1: 32(0x20) KB
L2: 2048(0x800) KB
Chip ID: 29824(0x7480)
ASIC Revision: 0(0x0)
Cacheline Size: 128(0x80)
Max Clock Freq. (MHz): 2493
BDFID: 1536
Internal Node ID: 2
Compute Unit: 32
SIMDs per CU: 2
Shader Engines: 2
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: 32(0x20)
Max Work-item Per CU: 1024(0x400)
Grid Max Size: 4294967295(0xffffffff)
Grid Max Size per Dimension:
x 4294967295(0xffffffff)
y 4294967295(0xffffffff)
z 4294967295(0xffffffff)
Max fbarriers/Workgrp: 32
Packet Processor uCode:: 542
SDMA engine uCode:: 21
IOMMU Support:: None
Pool Info:
Pool 1
Segment: GLOBAL; FLAGS: COARSE GRAINED
Size: 16760832(0xffc000) 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: 16760832(0xffc000) 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--gfx1102
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 4294967295(0xffffffff)
y 4294967295(0xffffffff)
z 4294967295(0xffffffff)
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 4294967295(0xffffffff)
y 4294967295(0xffffffff)
z 4294967295(0xffffffff)
FBarrier Max Size: 32
*** Done ***

Additional Information

No response
@ppanchad-amd
Copy link

Hi @evstigneevnm. Internal ticket has been created to investigate this issue. Thanks!

@lucbruni-amd
Copy link

lucbruni-amd commented May 28, 2025 •
edited
Loading

Hi @evstigneevnm! Thanks for opening an issue with us.

Before I share my results with the kernel, I have one question for my own clarification:

PROBLEM: uncommenting printf causes no output and the kernel is not executed

Which lines specifically are you commenting, other than the assert line, for the kernel to seemingly not execute? I initially thought you were referring to the printf("%i: %i\n", i, a_h); line, but your "output with commented lines" still displays output from that line (although the value of a_h appears to be 0).

I just tested your kernel on the latest ROCm version, 6.4.1, and was unable to spot any obvious issues. Here's my output along with some system information:

$ hipcc test_printf.hip -o test_printf
$ export HCC_ENABLE_PRINTF=1
$ ./test_printf

device_count: 1
0: Radeon RX 7900 XT, 21458059264, PCI: 3:0:0; 0, clock:1249000, busw:320, l2:6291456, maxth:2048 
Hello World
0: 4

$ sudo apt show rocm -a
Package: rocm
Version: 6.4.1.60401-83~24.04
...

Is the output I got more of what you were expecting? My first suggestion would be to try upgrading your ROCm version to the latest. If possible, please try again on the latest version and report back if you are still receiving the issue. Thanks!

@evstigneevnm
Copy link
Author

Hello @lucbruni-amd ! Thank you for replying!
I have updated the driver, but still got the same result.
For your reference, i post the kernel code, the results and the driver version here:

$ sudo apt show rocm -a
Package: rocm
Version: 6.4.1.60401-83~24.04
Priority: optional
Section: devel
Maintainer: ROCm dev support <[email protected]>
Installed-Size: 13.3 kB
Depends: rocm-utils (= 6.4.1.60401-83~24.04), rocm-developer-tools (= 6.4.1.60401-83~24.04), rocm-openmp-sdk (= 6.4.1.60401-83~24.04), rocm-opencl-sdk (= 6.4.1.60401-83~24.04), rocm-ml-sdk (= 6.4.1.60401-83~24.04), mivisionx (= 3.2.0.60401-83~24.04), migraphx (= 2.12.0.60401-83~24.04), rpp (= 1.9.10.60401-83~24.04), rocm-core (= 6.4.1.60401-83~24.04), migraphx-dev (= 2.12.0.60401-83~24.04), mivisionx-dev (= 3.2.0.60401-83~24.04), rpp-dev (= 1.9.10.60401-83~24.04)
Homepage: https://github.com/RadeonOpenCompute/ROCm
Download-Size: 2140 B
APT-Manual-Installed: yes
APT-Sources: https://repo.radeon.com/rocm/apt/6.4.1 noble/main amd64 Packages
Description: Radeon Open Compute (ROCm) software stack meta package

Package: rocm
Version: 6.1.5.60105-134~22.04
Priority: optional
Section: devel
Maintainer: ROCm dev support <[email protected]>
Installed-Size: 13.3 kB
Depends: rocm-utils (= 6.1.5.60105-134~22.04), rocm-developer-tools (= 6.1.5.60105-134~22.04), rocm-openmp-sdk (= 6.1.5.60105-134~22.04), rocm-opencl-sdk (= 6.1.5.60105-134~22.04), rocm-ml-sdk (= 6.1.5.60105-134~22.04), mivisionx (= 2.5.0.60105-134~22.04), migraphx (= 2.9.0.60105-134~22.04), rpp (= 1.5.0.60105-134~22.04), rocm-core (= 6.1.5.60105-134~22.04), migraphx-dev (= 2.9.0.60105-134~22.04)
Homepage: https://github.com/RadeonOpenCompute/ROCm
Download-Size: 2052 B
APT-Sources: https://repo.radeon.com/rocm/apt/6.1.5 jammy/main amd64 Packages
Description: Radeon Open Compute (ROCm) software stack meta package

$ which hipcc
/opt/rocm/bin/hipcc

$ hipcc --version
HIP version: 6.4.43483-a187df25c
AMD clang version 19.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.4.1 25184 c87081df219c42dc27c5b6d86c0525bc7d01f727)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-6.4.1/lib/llvm/bin
Configuration file: /opt/rocm-6.4.1/lib/llvm/bin/clang++.cfg

$ cat ./test_printf.hip | grep __global__ -A5
__global__ void run_printf(int i, int* a) 
{ 
	a[0] = 4+i;
	printf("Hello World\n"); 
//	assert(a[0] == 4);
}

$ hipcc test_printf.hip -o test_printf
$ export HCC_ENABLE_PRINTF=1
$ ./test_printf
device_count: 2
0: AMD Radeon™ RX 7600 XT, 17163091968, PCI: 3:0:0; 0, clock:1124000, busw:128, l2:2097152, maxth:2048 
0: 0
1: AMD Radeon™ RX 7600 XT, 17163091968, PCI: 6:0:0; 0, clock:1124000, busw:128, l2:2097152, maxth:2048 
1: 0
PASSED!

I bought a new SSD, will reinstall a new linux and test all possible driver combinations, retaining other hardware intact. Then will report back here.

@evstigneevnm
Copy link
Author

Hello dear @lucbruni-amd !
I did some research and found that the support of the PCI atomics might be a problem.
My current configuration is:

$ lspci -t
...
 \-[0000:00]-+-00.0
             +-01.0-[07]--
             +-02.0-[01-03]----00.0-[02-03]----00.0-[03]--+-00.0
             |                                            \-00.1
             +-03.0-[04-06]----00.0-[05-06]----00.0-[06]--+-00.0
             |                                            \-00.1
...

and their atomic bits:

$ sudo lspci -kkvvv -s 01:00.0 | grep -i atomic 
			 AtomicOpsCap: Routing+ 32bit- 64bit- 128bitCAS-
			 AtomicOpsCtl: EgressBlck-

$ sudo lspci -kkvvv -s 2:00.0 | grep -i atomic
			 AtomicOpsCap: Routing+
			 AtomicOpsCtl: EgressBlck-

$ sudo lspci -kkvvv -s 03:00.0 | grep -i atomic 
			 AtomicOpsCap: 32bit+ 64bit+ 128bitCAS-
			 AtomicOpsCtl: ReqEn+

$ sudo lspci -kkvvv -s 04:00.0 | grep -i atomic 
			 AtomicOpsCap: Routing+ 32bit- 64bit- 128bitCAS-
			 AtomicOpsCtl: EgressBlck-

$ sudo lspci -kkvvv -s 5:00.0 | grep -i atomic
			 AtomicOpsCap: Routing+
			 AtomicOpsCtl: EgressBlck-

$ sudo lspci -kkvvv -s 06:00.0 | grep -i atomic 
			 AtomicOpsCap: 32bit+ 64bit+ 128bitCAS-
			 AtomicOpsCtl: ReqEn+

I also downloaded rocm-examples and tried to run reduciton tests. I got the following error:

$ ./rocm-examples/build/bin/Tutorials/reduction/reduction_test_v6
[==========] Running 85 tests from 1 test suite.
[----------] Global test environment set-up.
[----------] 85 tests from Reduction
[ RUN      ] Reduction.10_64_2
hipKernelLaunchGGL(the operation cannot be performed in the present state)

Can this be related to PCI-express atomics?

@lucbruni-amd
Copy link

Hi @evstigneevnm,

Thanks for your research thus far!

This could indeed be related to PCIe atomics as printf/assert/malloc/free etc. host calls will not work if they are not supported.
What motherboard are you running at the moment?

Also, could you share the full output of the following command:

sudo lspci -vvv.

Thanks!

@evstigneevnm
Copy link
Author

Dear @lucbruni-amd ,
Sorry for the delay. The PC is currently disassembled and i will test a new hardware soon. For the reference i will assemble it back on the old hardware and test

lspci -vvv

for you. Be back tomorrow.

@evstigneevnm
Copy link
Author

Dear @lucbruni-amd !
I have installed a newer hardware and everything is working as expected!
For my example with the kernel

__global__ void run_printf(int i, int* a) 
{ 
	a[0] = 4+i;
	printf("Hello World\n"); 
}

I' getting the expected output:

$ ./test_printf.bin 
device_count: 2
0: AMD Radeon™ RX 7600 XT, 17163091968, PCI: 14:0:0; 0, clock:1124000, busw:128, l2:2097152, maxth:2048 
Hello World
0: 4
1: AMD Radeon™ RX 7600 XT, 17163091968, PCI: 8:0:0; 0, clock:1124000, busw:128, l2:2097152, maxth:2048 
Hello World
1: 5

The problem was with the outdated motherboard. Sorry for this misleading issue and thank you for the support.
I hope this helps to someone else. For the reference i post here old and new lspci -vvv outputs (old one with the other older GPU installed).
My only fancy is that you might want to add an explicit error message for these kind of mishappens.
Best regards, Nick.

lspcivvv_new.txt
lspcivvv_old.txt

@lucbruni-amd
Copy link

Hi, Nick!

Thanks for getting back to me, and I'm glad to hear you were able to resolve the issue with a new motherboard!
This will certainly help others who may stumble across these issues.

I agree with your suggestion for a more explicit message or warning for this. I'll forward this to the team for review.

Closing this issue as resolved by upgrading hardware to support PCIe atomics resulting in expected kernel output. Feel free to open new issues if you encounter problems with ROCm - we're here to help. Thanks!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
Under Investigation
Projects
None yet
Milestone
No milestone
Development

No branches or pull requests
3 participants
@evstigneevnm @ppanchad-amd @lucbruni-amd