[n1gh7e0us]

The lack of ROCm support has been taken up in candle for a while now

Earlier there was a PR raised #3186 which was a completely AI generated pull request and was meant to be shut down anyways.

Since then there was no follow up but in the issue #3327 so I thought it would be nice to begin with the implementation.
This is still an initial draft but would love to get reviews on the code structure and all.

The current code is AI generated (guided by me for each small thing though) and definitely needs alot of clean up and fixing, most of the structure has been kept similar to that of cuda_backend.

rocm-rs was chosen as the crate for ROCm bindings due to the recent activity and tested that it does compile the example GPU kernel on my GPU.

Fedora has 6.4.2 that's why I used that, would love to test this out for 7.1.0 as well if someone can get me access to that machine

Specs for current test machine:
OS: Fedora Linux 43 (KDE Plasma Desktop Edition) x86_64
Kernel: Linux 6.19.8-200.fc43.x86_64
CPU: AMD Ryzen 7 7700X (16) @ 5.58 GHz
GPU: AMD Radeon RX 7700 XT (Discrete, gfx1101), AMD Raphael (Integrated)
ROCM: 6.4.2

Current status:
Discovering GPU works for now and we can successfully discover the GPU present.

Next steps:

• loading tensor to GPU
• implementing rand
• mat functions
• getting example code from README.md to work
• convolutions
• hopefully train a model? someone please project manage me xD
• Update the API to make more sense, currently its AI slop coded so its a mess

[n1gh7e0us]

Tested out zeros loading to GPU memory using nvtop and the following program. Everything is pretty bare and there wasn't any GPU memory trait/function to call so I went with nvtop to view it

Program for anyone to test :)

use candle_core::{Device, Tensor};

fn main() -> Result<(), Box<dyn std::error::Error>> {
    let device = Device::new_rocm(0)?;
    println!("Device discovered: {:?}", device);

    let size = 100000000usize;
    for i in 0..1000000 {
        let a = Tensor::zeros(size, candle_core::DType::F32, &device)?;
        println!("{a}");
        drop(a);
        println!("[{i}] dropped tensor");
        // std::thread::sleep(std::time::Duration::from_secs(2));
    }
    Ok(())
}

[n1gh7e0us]

Alrighty so the example code for candle-core from README.md does work now with new_rocm device.

Code for reference that was used.

use candle_core::{Device, Tensor};

fn main() -> Result<(), Box<dyn std::error::Error>> {
    let device = Device::new_rocm(0)?;
    println!("Device discovered: {:?}", device);

    let a = Tensor::randn(0f32, 1., (2, 3), &device)?;
    let b = Tensor::randn(0f32, 1., (3, 4), &device)?;

    let c = a.matmul(&b)?;
    println!("{c}");

    Ok(())
}

This is great news :)
atleast the RX7700Xt is being used for calculation of matmul of the rand.

Next steps:
Lets see if we can get convolutions working and training a model.

[n1gh7e0us]

Current situation: Need to implement ROCm kernels for running mathematical operations like Sigmoid etc.

Cuda kernels are using JIT for which functionality is present in cudarc and nvcc but runtime compilation isn't available in hipcc but inside of hiprtc for which rocm-rs doesn't have the bindings/functionality available

Implementing AOT (Ahead of Time) compilation instead so initial compilation time would be very slow cause you will have to compile the kernels, but it would be optimized for your GPU's architecture :D

Sounds like a good idea. Open to suggestions

[n1gh7e0us]

Candle-nn now trains a primitive nn (linear layer) through using ROCm

use candle_core::{DType, Device, Result, Tensor};
use candle_nn::{Linear, Module, Optimizer, VarBuilder, VarMap, SGD};

fn main() -> Result<()> {
    let device = Device::new_rocm(0)?;
    println!("Device: {:?}", device);

    let n_samples = 1000;
    let input_dim = 10;
    let output_dim = 1;

    let x = Tensor::randn(0f32, 1., (n_samples, input_dim), &device)?;
    let true_weight = Tensor::randn(0f32, 1., (output_dim, input_dim), &device)?;
    let true_bias = Tensor::randn(0f32, 1., output_dim, &device)?;
    let noise = Tensor::randn(0f32, 0.1, (n_samples, output_dim), &device)?;

    let y = x
        .matmul(&true_weight.t()?)?
        .broadcast_add(&true_bias)?
        .add(&noise)?;

    let varmap = VarMap::new();
    let vs = VarBuilder::from_varmap(&varmap, DType::F32, &device);
    let linear = Linear::new(
        vs.get((output_dim, input_dim), "weight")?,
        Some(vs.get(output_dim, "bias")?),
    );

    let learning_rate = 0.01;
    let mut sgd = SGD::new(varmap.all_vars(), learning_rate)?;

    let n_epochs = 100;
    for epoch in 1..=n_epochs {
        let preds = linear.forward(&x)?;
        let diff = preds.sub(&y)?;
        let loss = diff.sqr()?.mean_all()?;

        sgd.backward_step(&loss)?;

        if epoch % 10 == 0 {
            let loss_val = loss.to_vec0::<f32>()?;
            println!("Epoch {}: loss = {:.6}", epoch, loss_val);
        }
    }

    println!("\nTraining completed!");
    let final_preds = linear.forward(&x)?;
    let final_loss = final_preds.sub(&y)?.sqr()?.mean_all()?.to_vec0::<f32>()?;
    println!("Final loss: {:.6}", final_loss);

    Ok(())
}

[n1gh7e0us]

another force push cause I forgot to sync things in my main branch
its inline with the main branch in huggingface/candle repo so we are gucci

[n1gh7e0us]

BERT now runs on ROCm!!

Results from: cargo run --example bert --features rocm -- --prompt "The cat sits on the mat"

Yayy

[n1gh7e0us]

alot of things are still unsafe and not implemented but slowly can be done over time. Bert works atleast which was one of the milestones.

[noot]

confirming that I got this working on a strix halo (gfx1151) on debian 13 with a few fixes:

• needed this patch add workspace to cargo template, add +nightly flag for build RustNSparks/rocm_kernel_macros#1
• needed to use TheRock 7.12.0rc1 nightly build for gfx1151 from AMD's prerelease server - ROCm 6.4 didn't support gfx1151
• added __build_in_kernels_ to the workspace exclude in Cargo.toml on this branch

built with:

THEROCK=~/therock-7.12
ROCM_PATH=$THEROCK HIP_PATH=$THEROCK LIBRARY_PATH=$THEROCK/lib \
C_INCLUDE_PATH=$THEROCK/include CPLUS_INCLUDE_PATH=$THEROCK/include \
BINDGEN_EXTRA_CLANG_ARGS="-I$THEROCK/include" \
cargo build -p candle-core --features rocm

small test program:

// main.rs
use candle_core::{Device, Tensor};

fn main() -> Result<(), Box<dyn std::error::Error>> {
    println!("attempting to create ROCm device...");
    let device = Device::new_rocm(0)?;
    println!("device: {:?}", device);
    
    let a = Tensor::randn(0f32, 1., (2, 3), &device)?;
    println!("random tensor on GPU: {}", a);
    
    let b = Tensor::randn(0f32, 1., (2, 3), &device)?;
    let c = (&a + &b)?;
    println!("a + b = {}", c);
    
    let d = Tensor::randn(0f32, 1., (2, 4), &device)?;
    let e = Tensor::randn(0f32, 1., (4, 3), &device)?;
    let f = d.matmul(&e)?;
    println!("matmul result: {}", f);
    
    println!("all tests passed!");
    Ok(())
}

# Cargo.toml
[package]
name = "candle-rocm-test"
version = "0.1.0"
edition = "2021"

[dependencies]
candle-core = { path = "~/candle/candle-core", features = ["rocm"] }

[patch.crates-io]
rocm_kernel_macros = { path = "../rocm_kernel_macros" }

run with:

THEROCK=~/therock-7.12
PATH=$THEROCK/bin:$PATH \
LD_LIBRARY_PATH=$THEROCK/lib \
ROCM_PATH=$THEROCK HIP_PATH=$THEROCK \
./target/debug/candle-rocm-test

output:

$ PATH=$THEROCK/bin:$PATH LD_LIBRARY_PATH=$THEROCK/lib ROCM_PATH=$THEROCK HIP_PATH=$THEROCK ./target/debug/candle-rocm-test 
attempting to create ROCm device...
device: Rocm(RocmDevice(DeviceId(1)))
random tensor on GPU: [[ 1.1592, -0.8498,  2.0249],
 [ 0.6686, -2.3056, -0.5295]]
Tensor[[2, 3], f32, rocm:0]
a + b = [[-0.2396, -0.5602,  0.8829],
 [-0.3299, -2.2806, -0.4325]]
Tensor[[2, 3], f32, rocm:0]
matmul result: [[ 3.6954, -0.7929, -0.9816],
 [ 0.9673, -0.2567, -0.3048]]
Tensor[[2, 3], f32, rocm:0]
all tests passed!

machine specs:

OS: Debian GNU/Linux 13 (trixie) x86_64
Host: Desktop (AMD Ryzen AI Max 300 Series) (A6)
Kernel: Linux 6.12.74+deb13+1-amd64
CPU: AMD RYZEN AI MAX+ 395 (32) @ 5.18 GHz
GPU: AMD Radeon Graphics / Radeon 8050S Graphics / Radeon 8060S Graphics [Integrated]

I'd love to see this pushed forward, happy to help in any way!

[zachelnet]

**Confirming ROCm support works on RDNA 4 (RX 9070 XT) 🎉

Hardware / OS:

• Nobara 44 (Fedora-based), AMD Radeon RX 9070 XT (gfx1201)
• ROCm 7.1.0 via Nobara repos (rocm-7.1.0-3.fc44)
• Rust 1.96.1

Build adjustments needed:

• rocm-rs requires default-features = false, features = ["miopen"] to skip the rocm_kernel_macros GPU kernel compilation (which needs nightly + amdgcn-amd-amdhsa target)
• rocm-rs 0.5.2 has a missing #[cfg(feature = "macros")] gate in rocarray/mod.rs — sorting module references hip::memory_ext unconditionally but that module requires the macros feature. Needs either a local patch or upstream fix.

Build command:

cargo build -p candle-core --features rocm

cargo run --example bert --features rocm
...
score: 0.06 'The new movie is awesome' 'The cat plays in the garden'
score: 0.04 'The cat plays in the garden' 'A woman watches TV'
score: 0.04 'The cat sits outside' 'The cat plays in the garden'
score: 0.03 'I love pasta' 'The new movie is awesome'
score: 0.02 'A man is playing guitar' 'I love pasta'

Great work — excited to see ROCm support land in Candle!

[xmiksay]

I hope I have finished your PR in #3801.
I needed to rebase it to current master so I have lost your commit authorship - how to handle it? Easiest way is to rebase your branch and continue with this start point

[n1gh7e0us]

@xmiksay thanks for taking care of it! Its fine if we lose my authorship :) I don't mind as long as candle gets AMD support