ROCm · kiran-thumma · May 27, 2026 · May 27, 2026 · May 27, 2026 · May 27, 2026
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -49,7 +49,7 @@ install(CODE "set(CMAKE_INSTALL_LOCAL_ONLY TRUE)" ALL_COMPONENTS)
 set(PYTHON_SUPPORTED_VERSIONS "3.10" "3.11" "3.12" "3.13" "3.14")
 
 # Supported AMD GPU architectures.
-set(HIP_SUPPORTED_ARCHS "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201")
+set(HIP_SUPPORTED_ARCHS "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1250;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201")
 
 # ROCm installation prefix. Default to /opt/rocm but allow override via
 # -DROCM_PATH=/your/rocm/path when invoking cmake.
@@ -1370,6 +1370,17 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
     "csrc/rocm/torch_bindings.cpp"
     "csrc/rocm/skinny_gemms.cu"
     "csrc/rocm/attention.cu")
+  set(VLLM_ROCM_EXT_FLAGS ${VLLM_GPU_FLAGS})
+
+  # skinny_gemms.cu is built on gfx9/gfx11 ISA (MFMA, dot2/dot4, legacy
+  # s_waitcnt asm) that gfx1250 (gfx12) does not provide. Exclude it from the
+  # gfx1250 build and disable its op registrations (VLLM_SKIP_SKINNY_GEMMS);
+  # vLLM falls back to default/Triton GEMM for those ops on gfx1250.
+  if(VLLM_GPU_ARCHES MATCHES "gfx1250")
+    message(STATUS "gfx1250 detected: excluding csrc/rocm/skinny_gemms.cu from _rocm_C")
+    list(REMOVE_ITEM VLLM_ROCM_EXT_SRC "csrc/rocm/skinny_gemms.cu")
+    list(APPEND VLLM_ROCM_EXT_FLAGS "-DVLLM_SKIP_SKINNY_GEMMS")
+  endif()
 
   set(VLLM_ROCM_HAS_GFX1100 OFF)
   if(VLLM_GPU_ARCHES MATCHES "gfx1100")
@@ -1385,10 +1396,15 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
     DESTINATION vllm
     LANGUAGE ${VLLM_GPU_LANG}
     SOURCES ${VLLM_ROCM_EXT_SRC}
-    COMPILE_FLAGS ${VLLM_GPU_FLAGS}
+    COMPILE_FLAGS ${VLLM_ROCM_EXT_FLAGS}
     ARCHITECTURES ${VLLM_GPU_ARCHES}
     USE_SABI 3
     WITH_SOABI)
+
+  # Needed to pass skip skinny flag to .cpp targets
+  if(VLLM_GPU_ARCHES MATCHES "gfx1250")                               
+    target_compile_definitions(_rocm_C PRIVATE VLLM_SKIP_SKINNY_GEMMS)
+  endif()
 
   if(VLLM_ROCM_HAS_GFX1100)
     target_compile_definitions(_rocm_C PRIVATE VLLM_ROCM_GFX1100)

diff --git a/benchmarks/vllm_smoketest.sh b/benchmarks/vllm_smoketest.sh
@@ -0,0 +1,102 @@
+#!/bin/bash
+set -euo pipefail
+
+# add model args here
+declare -A MODEL_PATHS=(
+  [gpt-oss-120b-mxfp4]="/data/amd/gpt-oss-120b-w-mxfp4-a-fp8"
+  [DeepSeek-R1-0528-MXFP4]="/data/amd/DeepSeek-R1-0528-MXFP4"
+)
+declare -A MODEL_SERVE_ARGS=(
+  [gpt-oss-120b-mxfp4]="--tensor-parallel-size 1 --gpu_memory_utilization 0.7 --attention-backend TRITON_ATTN"
+  [DeepSeek-R1-0528-MXFP4]="--tensor-parallel-size 1 --gpu_memory_utilization 0.9 --dtype auto --no-enable-prefix-caching --disable-uvicorn-access-log --trust-remote-code"
+)
+declare -A MODEL_ENV=(
+  [gpt-oss-120b-mxfp4]="HSA_ENABLE_SDMA=0 USE_SVM=0 HSA_XNACK=0 VLLM_ROCM_AITER_FUSED_MOE_TRITON_GEMM_A4W4=1 VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_UNIFIED_ATTENTION=0 VLLM_ROCM_USE_SKINNY_GEMM=0 VLLM_ROCM_USE_AITER_RMSNORM=0"
+  [DeepSeek-R1-0528-MXFP4]="HSA_ENABLE_SDMA=0 USE_SVM=0 HSA_XNACK=0 VLLM_ROCM_AITER_FUSED_MOE_TRITON_GEMM_A4W4=1 VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER_UNIFIED_ATTENTION=1 VLLM_ROCM_USE_SKINNY_GEMM=0 VLLM_ROCM_USE_AITER_RMSNORM=0"
+)
+
+MODEL_NAME="gpt-oss-120b-mxfp4" # default
+MODEL_PATH_OVERRIDE=""         # overrided path
+MODEL_SET=0                    # whether --model was explicitly passed
+RUN_LM_EVAL=0
+PORT=8000
+
+usage() {
+  cat <<'EOF'
+Usage: ./vllm_smoketest.sh [--model NAME] [--model-path PATH] [--lm-eval] [--port PORT] [--list]
+  --model NAME       Registry key to run (default: gpt-oss-120b-mxfp4)
+  --model-path PATH  Override model path for this run
+  --lm-eval          Run lm_eval after curl check
+  --port PORT        Server port (default: 8000)
+  --list             List available models
+EOF
+}
+
+while [[ $# -gt 0 ]]; do
+  case "$1" in
+    --model)      MODEL_NAME="$2"; MODEL_SET=1; shift 2 ;;
+    --model-path) MODEL_PATH_OVERRIDE="$2"; shift 2 ;;
+    --lm-eval)    RUN_LM_EVAL=1; shift ;;
+    --port)       PORT="$2"; shift 2 ;;
+    --list)       for n in "${!MODEL_PATHS[@]}"; do echo "$n -> ${MODEL_PATHS[$n]}"; done; exit 0 ;;
+    -h|--help)    usage; exit 0 ;;
+    *) echo "Unknown argument: $1" >&2; usage; exit 1 ;;
+  esac
+done
+
+if [[ -n "$MODEL_PATH_OVERRIDE" && $MODEL_SET -eq 0 ]]; then
+  echo "ERROR: --model-path requires --model (which model's serve args/env to use)." >&2
+  usage; exit 1
+fi
+
+MODEL="${MODEL_PATHS[$MODEL_NAME]:-$MODEL_NAME}"
+[[ -n "$MODEL_PATH_OVERRIDE" ]] && MODEL="$MODEL_PATH_OVERRIDE"
+SERVE_ARGS="${MODEL_SERVE_ARGS[$MODEL_NAME]:-}"
+MODEL_ENV_ARGS="${MODEL_ENV[$MODEL_NAME]:-}"
+echo "Model: $MODEL | Port: $PORT | lm_eval: $RUN_LM_EVAL"
+
+# temp
+pip3 install "fastapi<0.137"
+
+LOG_DIR="$(pwd)/vllm_smoke_test_logs"
+mkdir -p "$LOG_DIR"
+echo "Logs: $LOG_DIR"
+
+server_pid=""
+cleanup() { [[ -n "$server_pid" ]] && kill "$server_pid" 2>/dev/null && wait "$server_pid" 2>/dev/null; true; }
+trap cleanup EXIT
+trap 'exit 130' INT TERM
+
+# --- start server ---
+env $MODEL_ENV_ARGS \
+vllm serve --model "$MODEL" --host localhost --port "$PORT" $SERVE_ARGS &
+server_pid=$!
+echo "Waiting for server (pid $server_pid)..."
+
+until curl -s "http://localhost:$PORT/health" &>/dev/null; do
+  ps -p "$server_pid" >/dev/null || { echo "ERROR: server died"; tail -n 50 "$LOG_DIR/server.log"; exit 1; }
+  sleep 5
+done
+echo "Server ready."
+
+# --- check if text was generated ---
+echo "=== Curl completion check ==="
+http_code=$(curl -s -o "$LOG_DIR/curl_completion.log" -w '%{http_code}' \
+  "http://localhost:$PORT/v1/completions" \
+  -H "Content-Type: application/json" \
+  -d "{\"model\":\"$MODEL\",\"prompt\":\"San Francisco is a\",\"max_tokens\":32,\"temperature\":0}")
+cat "$LOG_DIR/curl_completion.log"; echo
+[[ "$http_code" == "200" ]] || { echo "FAIL: HTTP $http_code" >&2; exit 1; }
+grep -q '"text"' "$LOG_DIR/curl_completion.log" || { echo "FAIL: no text in response" >&2; exit 1; }
+echo "PASS"
+
+# --- lm_eval ---
+if [[ $RUN_LM_EVAL -eq 1 ]]; then
+  echo "=== lm_eval ==="
+  pip install "lm-eval[api]"
+  lm_eval --model local-completions \
+    --model_args "model=$MODEL,base_url=http://localhost:$PORT/v1/completions,num_concurrent=64,max_retries=3,tokenized_requests=False" \
+    --tasks gsm8k --num_fewshot 3 --limit 100 2>&1 | tee "$LOG_DIR/lm_eval.log"
+fi
+
+echo "Done."
diff --git a/cmake/external_projects/triton_kernels.cmake b/cmake/external_projects/triton_kernels.cmake
@@ -1,6 +1,6 @@
 # Install OpenAI triton_kernels from https://github.com/triton-lang/triton/tree/main/python/triton_kernels
 
-set(DEFAULT_TRITON_KERNELS_TAG "v3.5.1")
+set(DEFAULT_TRITON_KERNELS_TAG "padroute")
 
 # Set TRITON_KERNELS_SRC_DIR for use with local development with vLLM. We expect TRITON_KERNELS_SRC_DIR to
 # be directly set to the triton_kernels python directory.
@@ -11,13 +11,14 @@ if (DEFINED ENV{TRITON_KERNELS_SRC_DIR})
           SOURCE_DIR $ENV{TRITON_KERNELS_SRC_DIR}
   )
 
+#TODO (JPVILLAM): Test new release of triton kernels and put back to normal
 else()
-  set(TRITON_GIT "https://github.com/triton-lang/triton.git")
+  set(TRITON_GIT "https://github.com/jpvillam-amd/triton.git")
   message (STATUS "[triton_kernels] Fetch from ${TRITON_GIT}:${DEFAULT_TRITON_KERNELS_TAG}")
   FetchContent_Declare(
           triton_kernels
           # TODO (varun) : Fetch just the triton_kernels directory from Triton
-          GIT_REPOSITORY https://github.com/triton-lang/triton.git
+          GIT_REPOSITORY https://github.com/jpvillam-amd/triton.git
           GIT_TAG ${DEFAULT_TRITON_KERNELS_TAG}
           GIT_PROGRESS TRUE
           SOURCE_SUBDIR python/triton_kernels/triton_kernels

diff --git a/csrc/quickreduce/base.h b/csrc/quickreduce/base.h
@@ -81,11 +81,11 @@ union BufferResource {
 
 __quickreduce_device_inline__ static int32x4_t buffer_load_dwordx4(
     int32x4_t srsrc, int32_t voffset, int32_t soffset,
-    int32_t aux) __asm("llvm.amdgcn.raw.buffer.load.v4i32");
+    int32_t aux) {}
 
 __quickreduce_device_inline__ static void buffer_store_dwordx4(
     int32x4_t data, int32x4_t srsrc, int32_t voffset, int32_t soffset,
-    int32_t aux) __asm("llvm.amdgcn.raw.buffer.store.v4i32");
+    int32_t aux) {}
 
 __quickreduce_device_inline__ static void set_fp16_ovfl(bool const value) {
 #if defined(__gfx942__)

@@ -2405,13 +2405,23 @@ template <typename T, int absz, int cbid, int blgp>
 __device__ __forceinline__ floatx8 gcn_wmma16x16x16_instr(const bit16x8& inpA,
                                                           const bit16x8& inpB,
                                                           const floatx8& inpC) {
+#if defined(__gfx1250__)
+  // gfx1250 (gfx12 family) does not provide the gfx12 WMMA variant used by
+  // gfx1200/1201 (needs wmma-128b-insts). This custom-attention WMMA path is
+  // unsupported on gfx1250; trap if ever launched (fail loud, not silent-wrong).
+  (void)inpA;
+  (void)inpB;
+  __builtin_trap();
+  return inpC;
+#else
   if constexpr (std::is_same<T, _Float16>::value) {
     return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(inpA, inpB, inpC);
   } else if constexpr (std::is_same<T, __hip_bfloat16>::value) {
     return __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(inpA, inpB, inpC);
   } else {
     static_assert(false, "unsupported 16b dtype");
   }
+#endif
 }
 
 template <typename T>

@@ -14,6 +14,10 @@
 TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, rocm_ops) {
   // vLLM custom ops for rocm
 
+// skinny_gemms.cu (LLMM1/wvSplitK/wvSplitKrc/wvSplitKQ) is excluded on gfx1250
+// (gfx9/gfx11 ISA, unsupported there); skip these registrations to avoid
+// undefined symbols. vLLM uses default/Triton GEMM for these ops on gfx1250.
+#ifndef VLLM_SKIP_SKINNY_GEMMS
   // Custom gemm op for matrix-vector multiplication
   rocm_ops.def(
       "LLMM1(Tensor in_a, Tensor in_b, int rows_per_block) -> "
@@ -38,6 +42,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, rocm_ops) {
       "Tensor scale_a, "
       "          Tensor scale_b, int CuCount) -> ()");
   rocm_ops.impl("wvSplitKQ", torch::kCUDA, &wvSplitKQ);
+#endif  // VLLM_SKIP_SKINNY_GEMMS
 
 #ifdef VLLM_ROCM_GFX1100
   // W4A16 GPTQ kernels for AMD RDNA3 (gfx1100).

@@ -15,7 +15,7 @@ ARG NIC_BACKEND=all
 # Users can specify a custom version compatible with their host drivers.
 # The default version has been tested with ioinic-dkms=25.11.1.001
 ARG AINIC_VERSION=1.117.3-hydra
-ARG UBUNTU_CODENAME=jammy
+ARG UBUNTU_CODENAME=noble
 
 # Sccache configuration. Release builds use this today; CI can opt in when a
 # shared S3-compatible cache backend is available.
@@ -28,14 +28,14 @@ ARG SCCACHE_S3_NO_CREDENTIALS=0
 
 FROM ${BASE_IMAGE} AS base
 
-ARG ARG_PYTORCH_ROCM_ARCH
+ARG ARG_PYTORCH_ROCM_ARCH=gfx942;gfx950;gfx1250
 ENV PYTORCH_ROCM_ARCH=${ARG_PYTORCH_ROCM_ARCH:-${PYTORCH_ROCM_ARCH}}
 
 # Install build dependencies and utilities
 RUN apt-get update -q -y && apt-get install -q -y \
     sqlite3 libsqlite3-dev libfmt-dev libmsgpack-dev libsuitesparse-dev \
     apt-transport-https ca-certificates wget curl \
-    libnuma-dev ccache mold
+    build-essential libnuma-dev ccache mold
 RUN --mount=type=cache,target=/root/.cache/pip \
     python3 -m pip install --upgrade pip
 # Note: mold is installed but not set as the system default linker because
@@ -61,6 +61,10 @@ ENV UV_HTTP_TIMEOUT=500
 ENV UV_INDEX_STRATEGY="unsafe-best-match"
 # Use copy mode to avoid hardlink failures with Docker cache mounts
 ENV UV_LINK_MODE=copy
+# python binary fall back for non venv builds
+ENV UV_PYTHON=${VIRTUAL_ENV:-/usr}/bin/python3
+# Expose paths from wheel installation
+ENV PKG_CONFIG_PATH=${ROCM_PATH}/lib/rocm_sysdeps/lib/pkgconfig:${PKG_CONFIG_PATH}
 # ccache directory - persisted across layer rebuilds via cache mounts.
 ENV CCACHE_DIR=/root/.cache/ccache
 ENV CCACHE_COMPILERCHECK=content
@@ -110,8 +114,8 @@ WORKDIR ${COMMON_WORKDIR}
 FROM base AS fetch_vllm_0
 ONBUILD COPY ./ vllm/
 FROM base AS fetch_vllm_1
-ARG VLLM_REPO="https://github.com/vllm-project/vllm.git"
-ARG VLLM_BRANCH="main"
+ARG VLLM_REPO="https://github.com/ROCm/vllm.git"
+ARG VLLM_BRANCH="455_wip"
 ENV VLLM_REPO=${VLLM_REPO}
 ENV VLLM_BRANCH=${VLLM_BRANCH}
 ONBUILD RUN git clone ${VLLM_REPO} \
@@ -242,7 +246,7 @@ ARG RIXL_BRANCH="39be1de8"
 ARG RIXL_REPO="https://github.com/ROCm/RIXL.git"
 ARG UCX_BRANCH="bfb51733"
 ARG UCX_REPO="https://github.com/openucx/ucx.git"
-ENV ROCM_PATH=/opt/rocm
+# ENV ROCM_PATH=/opt/rocm -> correct ROCM_PATH is set in base image
 ENV UCX_HOME=/usr/local/ucx
 ENV RIXL_HOME=/usr/local/rixl
 ENV RIXL_BENCH_HOME=/usr/local/rixl_bench
@@ -298,6 +302,7 @@ RUN --mount=type=cache,target=/root/.cache/ccache \
     git checkout ${RIXL_BRANCH} && \
     CC="ccache gcc" CXX="ccache g++" \
     meson setup build --prefix=${RIXL_HOME} \
+                     --force-fallback-for=abseil-cpp \
                      -Ducx_path=${UCX_HOME} \
                      -Drocm_path=${ROCM_PATH} && \
     cd build && \
@@ -310,6 +315,10 @@ RUN --mount=type=cache,target=/root/.cache/ccache \
 RUN cd /opt/rixl && \
     sed -i "s/--exclude 'libamdhip64\*'/--exclude 'libamdhip64*' --exclude 'libcore*' --exclude 'libpull*'/" \
         contrib/build-wheel.sh && \
+    # The wheel build re-runs meson via meson-python; force the bundled abseil
+    sed -i 's|setup = \["-Dinstall_headers=false"\]|setup = ["-Dinstall_headers=false", "--force-fallback-for=abseil-cpp"]|' \
+        pyproject.toml && \
+    grep -q 'force-fallback-for' pyproject.toml && \
     mkdir -p /app/install && \
     _ucx_install_dir=${UCX_HOME} \
     ./contrib/build-wheel.sh \
@@ -326,7 +335,7 @@ ARG ROCSHMEM_REPO="https://github.com/ROCm/rocm-systems.git"
 # DeepEP only supports gfx942 and gfx950; build ROCShmem for the same set so
 # it can be linked against DeepEP without arch mismatches.
 ARG DEEPEP_ROCM_ARCH="gfx942;gfx950"
-ENV ROCM_PATH=/opt/rocm
+# ENV ROCM_PATH=/opt/rocm -> Correct rocm_path is set in base image
 ENV ROCSHMEM_DIR=/opt/rocshmem
 
 RUN --mount=type=cache,target=/root/.cache/ccache \
@@ -386,7 +395,7 @@ RUN /bin/bash -lc 'set -euo pipefail; \
    curl -fsSL https://packages.broadcom.com/artifactory/api/security/keypair/PackagesKey/public \
      -o /etc/apt/keyrings/broadcom-nic.asc; \
    chmod a+r /etc/apt/keyrings/broadcom-nic.asc; \
-   echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/broadcom-nic.asc] https://packages.broadcom.com/artifactory/ethernet-nic-debian-public jammy main" \
+   echo "deb [arch=amd64 signed-by=/etc/apt/keyrings/broadcom-nic.asc] https://packages.broadcom.com/artifactory/ethernet-nic-debian-public ${UBUNTU_CODENAME} main" \
      > /etc/apt/sources.list.d/broadcom-nic.list; \
    apt-get update && apt-get install -y --no-install-recommends \
      bnxt-rocelib=235.2.86.0 \
@@ -575,9 +584,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
 ENV HF_XET_HIGH_PERFORMANCE=1
 ENV HF_HUB_DOWNLOAD_TIMEOUT=60
 
-# Keep torch.cuda.is_available() fork-safe (see vllm/env_override.py).
-ENV PYTORCH_NVML_BASED_CUDA_CHECK=1
-
 # Pre-install vLLM test dependencies.
 COPY requirements/test/rocm.txt /tmp/rocm-test-reqs.txt
 RUN --mount=type=cache,target=/root/.cache/uv \
@@ -634,6 +640,10 @@ COPY --from=export_vllm /vllm_v1 /usr/local/lib/python${PYTHON_VERSION}/dist-pac
 # Hide source under src/ so it won't shadow the installed package in tests.
 RUN mkdir src && mv vllm src/vllm
 
+# tblib is required by conftest.py at collection time; --no-deps prevents
+# uv from pulling CUDA torch as a transitive dependency.
+RUN uv pip install --system --no-deps tblib
+
 # -----------------------
 # Final vLLM image
 FROM mori_base AS final
@@ -698,18 +708,23 @@ ENV SAFETENSORS_FAST_GPU=1
 # Performance environment variable.
 ENV HIP_FORCE_DEV_KERNARG=1
 
-# Keep torch.cuda.is_available() fork-safe (see vllm/env_override.py).
-ENV PYTORCH_NVML_BASED_CUDA_CHECK=1
-
 # Workaround for ROCm profiler limits
 RUN echo "ROCTRACER_MAX_EVENTS=10000000" > ${COMMON_WORKDIR}/libkineto.conf
 ENV KINETO_CONFIG="${COMMON_WORKDIR}/libkineto.conf"
 RUN echo "VLLM_BASE_IMAGE=${BASE_IMAGE}" >> ${COMMON_WORKDIR}/versions.txt \
     && echo "MORI_NIC_BACKEND=${NIC_BACKEND}" >> ${COMMON_WORKDIR}/versions.txt \
     && echo "AINIC_VERSION=${AINIC_VERSION}" >> ${COMMON_WORKDIR}/versions.txt
 
+
+### Install triton from upstream for AITER Deps TODO: (JPVILLAM) If possible to get this on whls it would be better
+RUN pip3 uninstall -y triton && \
+    git clone https://github.com/triton-lang/triton.git && \
+    cd triton && \
+    git checkout c517f38c && \
+    TRITON_APPEND_CMAKE_ARGS="-DCMAKE_FIND_USE_CMAKE_ENVIRONMENT_PATH=FALSE" pip3 install .
+
 CMD ["/bin/bash"]
 
 #Set entrypoint for vllm-openai official images
 FROM final AS vllm-openai
-ENTRYPOINT ["vllm", "serve"]
+ENTRYPOINT ["vllm", "serve"]
@@ -0,0 +1,14 @@
+# FFM overlay for gfx1250 builds
+ARG BASE_IMAGE=rocm/vllm-private:juan_455_npi_test
+FROM ${BASE_IMAGE}
+
+# extract tarball into /root/x/ffm/.
+ADD ubuntu_24_04_rel.tar.gz /root/x/ffm/
+ENV HSA_MODEL_LIB=/root/x/ffm/libhsakmtmodel.so
+ENV HSA_MODEL_TOML="/root/x/ffm/ffm_config.toml"
+ENV HSA_MODEL_ARGS=ffm_enable_time_slicing
+ENV HSA_MODEL_TOPOLOGY=/root/x/ffm/topology/mi450
+ENV HSA_MODEL_NUM_THREADS=64
+
+### Remove AMD SMI
+RUN pip uninstall -y amdsmi && rm -rf ${SITE_PACKAGES}/_rocm_sdk_core/share/amd_smi/amdsmi