Install AMD ROCm on WSL

rocm-wsl.md

This assumes you already have AMD Adrenalin installed on the Windows side with up-to-date chipset drivers.

Even if you are not compiling GPU programs, you still need the shared libraries for binaries that dynamically link to them at runtime. For example, running ldd llama-server reports that it requires many libhip*.so and libroc*.so files. If these are not in your LD_LIBRARY_PATH, the program will crash.

Install amdgpu-install

Go to https://repo.radeon.com/amdgpu-install and find the latest release for your OS. I'm on Ubuntu 24 (Noble):

wget https://repo.radeon.com/amdgpu-install/7.2.4/ubuntu/noble/amdgpu-install_7.2.4.70204-1_all.deb
sudo dpkg -i ./amdgpu-install_7.2.4.70204-1_all.deb

Install rocm

Note that starting with version 7.2.1, there is no longer a wsl usecase. Also note the use of --no-dkms because the kernel driver lives on the Windows side:

sudo amdgpu-install -y --usecase=rocm --no-dkms

Install librocdxg

This must be installed after ROCm because it adds librocdxg.so to /opt/rocm/lib:

wget https://github.com/ROCm/librocdxg/releases/download/v1.2.0/rocdxg-roct_1.2.0_amd64.deb
sudo dpkg -i ./rocdxg-roct_1.2.0_amd64.deb

Set environment variables

Add these to your ~/.bashrc:

export LD_LIBRARY_PATH=/opt/rocm/lib:${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}
export HSA_ENABLE_DXG_DETECTION=1

Uninstall

If you have multiple versions installed, use --rocmrelease=<version>, otherwise use all:

sudo apt remove -y --purge --autoremove rocdxg-roct
sudo amdgpu-install --uninstall --rocmrelease=all

Verify

Paste this into rocmcheck.cpp and compile with /opt/rocm/bin/hipcc rocmcheck.cpp -o rocmcheck then run ./rocmcheck:

#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>

#define HIP_CHECK(cmd) do {                                       \
    hipError_t _e = (cmd);                                        \
    if (_e != hipSuccess) {                                       \
        fprintf(stderr, "HIP error %d (%s) at %s:%d -> %s\n",     \
                (int)_e, hipGetErrorName(_e), __FILE__, __LINE__, \
                hipGetErrorString(_e));                           \
        return 1;                                                 \
    }                                                             \
} while (0)

/* Tiny kernel that squares each element on the device. */
__global__ void square(const float *in, float *out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = in[i] * in[i];
}

int main(void) {
    int rt = 0, drv = 0;
    HIP_CHECK(hipRuntimeGetVersion(&rt));
    HIP_CHECK(hipDriverGetVersion(&drv));
    printf("HIP runtime version : %d\n", rt);
    printf("HIP driver version  : %d\n", drv);

    int ndev = 0;
    HIP_CHECK(hipGetDeviceCount(&ndev));
    if (ndev == 0) {
        fprintf(stderr, "No ROCm devices found.\n");
        return 1;
    }
    printf("\n%d device(s) detected\n", ndev);

    HIP_CHECK(hipSetDevice(0));
    hipDeviceProp_t p;
    HIP_CHECK(hipGetDeviceProperties(&p, 0));
    printf("  Name         : %s\n", p.name);
    printf("  Architecture : %s\n", p.gcnArchName);
    printf("  Visible VRAM : %.2f GiB\n", (double)p.totalGlobalMem / (1024.0 * 1024.0 * 1024.0));
    printf("  Processors   : %d\n", p.multiProcessorCount);

    /* Use plain hipMalloc, NOT hipMallocManaged. */
    const int N = 1024;
    const size_t bytes = (size_t)N * sizeof(float);

    float *h_in  = (float *)malloc(bytes);
    float *h_out = (float *)malloc(bytes);
    if (!h_in || !h_out) {
        fprintf(stderr, "host alloc failed\n"); return 1;
    }
    for (int i = 0; i < N; i++) {
        h_in[i] = (float)i;
    }

    float *d_in = NULL, *d_out = NULL;
    HIP_CHECK(hipMalloc(&d_in, bytes));
    HIP_CHECK(hipMalloc(&d_out, bytes));
    HIP_CHECK(hipMemcpy(d_in, h_in, bytes, hipMemcpyHostToDevice));

    int threads = 256, blocks = (N + threads - 1) / threads;
    square<<<blocks, threads, 0, 0>>>(d_in, d_out, N);
    HIP_CHECK(hipGetLastError());       /* catches launch-config errors */
    HIP_CHECK(hipDeviceSynchronize());  /* catches execution errors */
    HIP_CHECK(hipMemcpy(h_out, d_out, bytes, hipMemcpyDeviceToHost));

    int ok = 1;
    for (int i = 0; i < N; i++) {
        float expect = h_in[i] * h_in[i];
        if (h_out[i] != expect) {
            fprintf(stderr, "Mismatch at %d: got %f, expected %f\n", i, h_out[i], expect);
            ok = 0;
            break;
        }
    }

    /* Cast to void because hipFree is [[nodiscard]] in HIP 7.x. */
    (void)hipFree(d_in);
    (void)hipFree(d_out);
    free(h_in);
    free(h_out);
    if (!ok) {
        fprintf(stderr, "\nGPU compute verification FAILED.\n");
        return 1;
    }
    printf("\nGPU compute verified (squared %d elements on device).\n", N);
    printf("\nSUCCESS: ROCm and librocdxg are working.\n");
    return 0;
}

You should see something like this:

HIP runtime version : 70253211
HIP driver version  : 70253211

1 device(s) detected
  Name         : AMD Radeon(TM) 8060S Graphics
  Architecture : gfx1151
  Visible VRAM : 95.65 GiB
  Processors   : 20

GPU compute verified (squared 1024 elements on device).

SUCCESS: ROCm and librocdxg are working.

Note that this GPU (Strix Halo) has 40 compute units (CU), but multiProcessorCount reports workgroup processors (WGP) and each WGP has 2 CUs.