From 074479544d25393c2b2d35e35564a8e3f83f2d42 Mon Sep 17 00:00:00 2001 From: Ameil Kumar Date: Mon, 17 Aug 2026 10:03:34 -0700 Subject: [PATCH 1/7] Add PyTorch v2.13.0 ROCm build script for UBI 10 (ppc64le) --- p/pytorch/build_info.json | 7 + p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh | 236 ++++++++++++++++++ ...ch_v2.13.0_rocm_cuda_guard_narrowing.patch | 89 +++++++ .../pytorch_v2.13.0_rocm_fastgeluasm.patch | 176 +++++++++++++ 4 files changed, 508 insertions(+) create mode 100644 p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh create mode 100644 p/pytorch/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch create mode 100644 p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch diff --git a/p/pytorch/build_info.json b/p/pytorch/build_info.json index 8db2ebbd0b..b465ddaf59 100644 --- a/p/pytorch/build_info.json +++ b/p/pytorch/build_info.json @@ -21,6 +21,13 @@ "v2.4.0": { "build_script": "pytorch_ubi_9.3.sh" }, + "v2.13.0-rocm": { + "build_script": "pytorch_2.13.0_rocm_ubi_10.sh", + "patches": { + "PATCH1": "https://raw.githubusercontent.com/ppc64le/build-scripts/refs/heads/master/p/pytorch/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch", + "PATCH2": "https://raw.githubusercontent.com/ppc64le/build-scripts/refs/heads/master/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch" + } + }, "*": { "build_script":"pytorch_2.7.1_ubi_9.3.sh" }, diff --git a/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh b/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh new file mode 100644 index 0000000000..02e23656c4 --- /dev/null +++ b/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh @@ -0,0 +1,236 @@ +#!/bin/bash -e +# ----------------------------------------------------------------------------- +# +# Package : pytorch +# Version : v2.13.0 +# Source repo : https://github.com/pytorch/pytorch.git +# Tested on : UBI:10 (ppc64le) +# Language : Python +# Ci-Check : True +# Script License: Apache License, Version 2 or later +# Maintainer : Ameil Kumar +# +# Disclaimer: This script has been tested in root mode on given +# ========== platform using the mentioned version of the package. +# It may not work as expected with newer versions of the +# package and/or distribution. In such case, please +# contact "Maintainer" of this script. +# +# ---------------------------------------------------------------------------- +# +# ROCm installation mode (--rocm-install-mode): +# rpms (default) - Install ROCm RPMs from a provided repo URL +# path - Assume ROCm is already present; use ROCM_PATH as-is +# +# Usage: +# ./pytorch_2.13.0_rocm_ubi_10.sh [--rocm-install-mode rpms|path] +# [--rocm-arch "gfx90a;gfx950"] +# [--version v2.13.0] +# +# Environment variables honoured (can be set before running): +# ROCM_PATH - Path to the ROCm installation (default: /opt/rocm) +# PYTORCH_ROCM_ARCH - Semicolon-separated GPU targets +# (default: "gfx90a;gfx950") +# ROCM_REPO_URL - RPM repo baseurl +# PACKAGE_VERSION - PyTorch tag to build (default: v2.13.0) +# +# --------------------------------------------------------------------------- + +set -e + +# Variables +PACKAGE_NAME=pytorch +PACKAGE_URL=https://github.com/pytorch/pytorch.git +PACKAGE_VERSION=${PACKAGE_VERSION:-v2.13.0} +SCRIPT_DIR=$(pwd) +OS_NAME=$(grep ^PRETTY_NAME /etc/os-release | cut -d= -f2) + +ROCM_INSTALL_MODE="rpms" # rpms | path +ROCM_REPO_URL=${ROCM_REPO_URL:-"https://public.dhe.ibm.com/software/server/POWER/Linux/AMD/ROCm/RHEL/10/ppc64le"} +ROCM_PATH=${ROCM_PATH:-/opt/rocm} + +# GPU architecture targets — override via env or --rocm-arch flag +PYTORCH_ROCM_ARCH=${PYTORCH_ROCM_ARCH:-"gfx90a;gfx950"} + +# --------------------------------------------------------------------------- +# Argument parsing +# --------------------------------------------------------------------------- +while [[ $# -gt 0 ]]; do + case "$1" in + --rocm-install-mode) + ROCM_INSTALL_MODE="$2" + shift 2 + ;; + --rocm-arch) + PYTORCH_ROCM_ARCH="$2" + shift 2 + ;; + --version) + PACKAGE_VERSION="$2" + shift 2 + ;; + *) + echo "Unknown argument: $1" + echo "Usage: $0 [--rocm-install-mode rpms|path] [--rocm-arch \"gfx90a;gfx950\"] [--version v2.13.0]" + exit 1 + ;; + esac +done + +if [[ "$ROCM_INSTALL_MODE" != "rpms" && "$ROCM_INSTALL_MODE" != "path" ]]; then + echo "ERROR: --rocm-install-mode must be one of: rpms, path" + exit 1 +fi + +echo "=== PyTorch ROCm Build ===" +echo " PACKAGE_VERSION : $PACKAGE_VERSION" +echo " ROCM_INSTALL_MODE : $ROCM_INSTALL_MODE" +echo " ROCM_PATH : $ROCM_PATH" +echo " PYTORCH_ROCM_ARCH : $PYTORCH_ROCM_ARCH" +echo "==========================" + +# --------------------------------------------------------------------------- +# Install system build dependencies +# --------------------------------------------------------------------------- +EPEL_URL="https://dl.fedoraproject.org/pub/epel/epel-release-latest-10.noarch.rpm" +if ! rpm -q epel-release &>/dev/null; then + echo "Installing EPEL" + dnf install -y "$EPEL_URL" +fi + +echo "Installing system dependencies" +dnf install -y git make wget patch cmake ninja-build \ + python3.13 python3.13-devel python3.13-pip \ + gcc gcc-c++ \ + openblas openblas-devel +echo "Installed required deps from RH" + +# Use Python 3.13 for the build so the produced wheel is cp313 +PYTHON=python3.13 + +# --------------------------------------------------------------------------- +# MODE: rpms — install ROCm from a provided RPM repository +# --------------------------------------------------------------------------- +if [[ "$ROCM_INSTALL_MODE" == "rpms" ]]; then + if [[ ! "$ROCM_REPO_URL" =~ ^(https?|file):// ]]; then + echo "ERROR: ROCM_REPO_URL does not look like a valid URL (got: ${ROCM_REPO_URL})" + exit 1 + fi + echo "Installing ROCm from ${ROCM_REPO_URL}" + + cat > /etc/yum.repos.d/rocm.repo </dev/null; then + echo "ERROR: hipcc not found under ROCM_PATH=${ROCM_PATH}. Check your ROCm installation." + exit 1 +fi +echo "ROCm hipcc: $(hipcc --version | head -1)" + +# Set ENV vars +# Use OpenBLAS instead of MKL (MKL does not support Power). +# Disable CUDA/Intel tooling so cmake does not search for them. +export PYTORCH_ROCM_ARCH +export BLAS=OpenBLAS +export USE_CUDA=0 +export USE_XPU=0 +export USE_ROCM=1 + +export CMAKE_PREFIX_PATH="${ROCM_PATH}:${CMAKE_PREFIX_PATH:-}" + +# Clone PyTorch +echo "Cloning PyTorch ${PACKAGE_VERSION}" +if [ -d "${SCRIPT_DIR}/pytorch" ]; then + echo "pytorch directory already exists, reusing." + cd "${SCRIPT_DIR}/pytorch" + git checkout "$PACKAGE_VERSION" +else + if ! git clone --recursive --branch "$PACKAGE_VERSION" "$PACKAGE_URL" "${SCRIPT_DIR}/pytorch"; then + echo "------------------$PACKAGE_NAME:clone_fails---------------------------------------" + echo "$PACKAGE_URL $PACKAGE_NAME" + echo "$PACKAGE_NAME | $PACKAGE_URL | $PACKAGE_VERSION | $OS_NAME | GitHub | Fail | Clone_Fails" + exit 1 + fi + cd "${SCRIPT_DIR}/pytorch" +fi + +git submodule sync +git submodule update --init --recursive + +# Install build dependencies +$PYTHON -m pip install --upgrade pip +$PYTHON -m pip install --group dev || $PYTHON -m pip install -r requirements.txt + +# Run ROCm source transformation +echo "Running ROCm hipify transformation" +$PYTHON tools/amd_build/build_amd.py + +# Apply patches +echo "Applying patches" +PATCH_BASE_URL=${PATCH_BASE_URL:-"https://raw.githubusercontent.com/ppc64le/build-scripts/refs/heads/master/p/pytorch"} + +# Fix CUDAGuard narrowing conversion errors under GCC 14 in ROCm HIP flash-attn files +wget -q --spider "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" && \ + wget -q -O "$SCRIPT_DIR/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" && \ + git apply "$SCRIPT_DIR/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" || echo "pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch missing or already applied, skipped" + +# Fix FastGeluAsm explicit specializations rejected by AMD clang 23.0 in composable_kernel +wget -q --spider "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_fastgeluasm.patch" && \ + wget -q -O "$SCRIPT_DIR/pytorch_v2.13.0_rocm_fastgeluasm.patch" "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_fastgeluasm.patch" && \ + git apply --directory=third_party/composable_kernel "$SCRIPT_DIR/pytorch_v2.13.0_rocm_fastgeluasm.patch" || echo "pytorch_v2.13.0_rocm_fastgeluasm.patch missing or already applied, skipped" + +# Build +echo "Building PyTorch (this will take a while)" +export PYTORCH_BUILD_VERSION=${PACKAGE_VERSION#v}+rocm +export PYTORCH_BUILD_NUMBER=1 + +if ! MAX_JOBS=$(nproc) $PYTHON -m pip install --no-build-isolation -v -e .; then + echo "------------------$PACKAGE_NAME:install_fails---------------------------------------" + echo "$PACKAGE_URL $PACKAGE_NAME" + echo "$PACKAGE_NAME | $PACKAGE_URL | $PACKAGE_VERSION | $OS_NAME | GitHub | Fail | Install_Fails" + exit 1 +fi + +# Build Wheels +echo "Building distribution wheel" +$PYTHON -m pip wheel --no-build-isolation -v -w dist . + +# Basic import test +echo "Running basic import test" +cd "${SCRIPT_DIR}" + +if ! $PYTHON -c "import torch; print('torch version :', torch.__version__); print('ROCm available:', torch.cuda.is_available())"; then + echo "------------------$PACKAGE_NAME:install_success_but_test_fails---------------------" + echo "$PACKAGE_URL $PACKAGE_NAME" + echo "$PACKAGE_NAME | $PACKAGE_URL | $PACKAGE_VERSION | $OS_NAME | GitHub | Fail | Install_success_but_Import_Fails" + exit 2 +fi + +echo "------------------$PACKAGE_NAME:install_&_test_both_success-------------------------" +echo "$PACKAGE_URL $PACKAGE_NAME" +echo "$PACKAGE_NAME | $PACKAGE_URL | $PACKAGE_VERSION | $OS_NAME | GitHub | Pass | Both_Install_and_Import_Success" +exit 0 diff --git a/p/pytorch/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch b/p/pytorch/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch new file mode 100644 index 0000000000..550b32fb98 --- /dev/null +++ b/p/pytorch/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch @@ -0,0 +1,89 @@ +diff --git a/aten/src/ATen/native/transformers/hip/flash_attn/aot/mha_all_aot.hip b/aten/src/ATen/native/transformers/hip/flash_attn/aot/mha_all_aot.hip +index c16f7d1aad233..49d7917b12852 100644 +--- a/aten/src/ATen/native/transformers/hip/flash_attn/aot/mha_all_aot.hip ++++ b/aten/src/ATen/native/transformers/hip/flash_attn/aot/mha_all_aot.hip +@@ -332,7 +332,7 @@ mha_varlen_fwd_aot(const at::Tensor &q, // total_q x num_heads x head_size, tot + TORCH_CHECK(!paged_KV, "[ROCm] mha_varlen_fwd: block_table_ must be nullopt"); + TORCH_CHECK(!alibi_slopes_.has_value(), "[ROCm] mha_varlen_fwd: alibi_slopes_ must be nullopt"); + +- at::cuda::CUDAGuard device_guard{(char)q.get_device()}; ++ at::cuda::CUDAGuard device_guard{q.get_device()}; + auto stream = at::cuda::getCurrentCUDAStream().stream(); + check_gpu_arch(stream); + +@@ -535,8 +535,7 @@ mha_bwd_aot(const at::Tensor &dout, // batch_size x seqlen_q x num_heads, x hea + const at::Tensor& philox_seed, + const at::Tensor& philox_offset) { + // Otherwise the kernel will be launched from cuda:0 device +- // Cast to char to avoid compiler warning about narrowing +- at::cuda::CUDAGuard device_guard{(char)q.get_device()}; ++ at::cuda::CUDAGuard device_guard{q.get_device()}; + auto stream = at::cuda::getCurrentCUDAStream().stream(); + check_gpu_arch(stream); + +@@ -717,8 +716,7 @@ mha_varlen_bwd_aot(const at::Tensor &dout, // total_q x num_heads, x head_size + TORCH_CHECK(!alibi_slopes_.has_value(), "[ROCm] mha_varlen_fwd: alibi_slopes_ must be nullopt"); + + // Otherwise the kernel will be launched from cuda:0 device +- // Cast to char to avoid compiler warning about narrowing +- at::cuda::CUDAGuard device_guard{(char)q.get_device()}; ++ at::cuda::CUDAGuard device_guard{q.get_device()}; + auto stream = at::cuda::getCurrentCUDAStream().stream(); + check_gpu_arch(stream); + +diff --git a/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_bwd_ck.hip b/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_bwd_ck.hip +index 86a580406c20d..5a08b26eb9c08 100644 +--- a/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_bwd_ck.hip ++++ b/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_bwd_ck.hip +@@ -455,8 +455,7 @@ mha_bwd_ck(const at::Tensor &dout, // batch_size x seqlen_q x + dout_padded = dout; + } + +- // Cast to char to avoid compiler warning about narrowing +- at::cuda::CUDAGuard device_guard{(char)q.get_device()}; ++ at::cuda::CUDAGuard device_guard{q.get_device()}; + + auto opts = q.options(); + auto softmax_d = at::empty({batch_size, num_heads, seqlen_q}, opts.dtype(at::kFloat)); +diff --git a/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_fwd_ck.hip b/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_fwd_ck.hip +index 1810031d92d52..9db5f7e350cfe 100644 +--- a/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_fwd_ck.hip ++++ b/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_fwd_ck.hip +@@ -362,8 +362,7 @@ mha_fwd_ck(const at::Tensor &q, // batch_size x seqlen_q x + const int head_size_8x = round_multiple(head_size, 8); + + // Otherwise the kernel will be launched from cuda:0 device +- // Cast to char to avoid compiler warning about narrowing +- at::cuda::CUDAGuard device_guard{(char)q.get_device()}; ++ at::cuda::CUDAGuard device_guard{q.get_device()}; + + auto opts = q.options(); + bool has_lse = true; +diff --git a/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_varlen_bwd_ck.hip b/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_varlen_bwd_ck.hip +index 5bb142400966e..e9b0c76cf96f9 100644 +--- a/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_varlen_bwd_ck.hip ++++ b/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_varlen_bwd_ck.hip +@@ -473,8 +473,7 @@ mha_varlen_bwd_ck(const at::Tensor &dout, // total_q x num_hea + dout_padded = dout; + } + +- // Cast to char to avoid compiler warning about narrowing +- at::cuda::CUDAGuard device_guard{(char)q.get_device()}; ++ at::cuda::CUDAGuard device_guard{q.get_device()}; + + auto opts = q.options(); + auto softmax_d = at::empty({batch_size, num_heads, max_seqlen_q}, opts.dtype(at::kFloat)); +diff --git a/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_varlen_fwd_ck.hip b/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_varlen_fwd_ck.hip +index eb4a3b841fa3e..4340cde7177a0 100644 +--- a/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_varlen_fwd_ck.hip ++++ b/aten/src/ATen/native/transformers/hip/flash_attn/ck/mha_varlen_fwd_ck.hip +@@ -303,8 +303,7 @@ mha_varlen_fwd_ck(const at::Tensor &q, // total_q x num_heads + const int head_size_8x = round_multiple(head_size_og, 8); + + // Otherwise the kernel will be launched from cuda:0 device +- // Cast to char to avoid compiler warning about narrowing +- at::cuda::CUDAGuard device_guard{(char)q.get_device()}; ++ at::cuda::CUDAGuard device_guard{q.get_device()}; + + auto opts = q.options(); + bool has_lse = true; \ No newline at end of file diff --git a/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch b/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch new file mode 100644 index 0000000000..452c238a22 --- /dev/null +++ b/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch @@ -0,0 +1,176 @@ +diff --git a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp b/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp +index 4e971649d..000000000 100644 +--- a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp ++++ b/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp +@@ -1154,94 +1154,97 @@ + + template + CK_TILE_DEVICE void operator()(Y& y, const X& x) const; ++}; + +- template <> +- CK_TILE_HOST void operator()(float& y, const float& x) const +- { ++#if defined(__HIP_DEVICE_COMPILE__) ++template <> ++CK_TILE_DEVICE void FastGeluAsm::operator()(float& y, const float& x) const ++{ ++ const uint32_t c1 = 0xbd92220c; // -2.0 * 0.035677f; ++ const float c2 = -2.0 * 0.797885f; ++ const uint32_t log2e_ = 0x3fb8aa3b; // log2e_v; ++ float tmp; ++ ++ asm volatile("v_mul_f32 %[v_tmp], %[v_x], %[v_x] ; x*x\n" ++ "v_fma_f32 %[v_tmp], %[v_tmp], %[s_c1], %[v_c2] ; c1*x*x+c2\n" ++ "v_mul_f32 %[v_tmp], %[v_tmp], %[v_x] ; x*(c1*x*x+c2)\n" ++ "v_mul_f32 %[v_tmp], %[v_tmp], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" ++ "v_exp_f32 %[v_tmp], %[v_tmp] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" ++ "s_nop 0 ; hazard for exp\n" ++ "v_add_f32 %[v_tmp], %[v_tmp], 1.0 ; emu+1.0f\n" ++ "v_rcp_f32 %[v_tmp], %[v_tmp] ; 1/(emu+1.0f)\n" ++ "s_nop 0 ; hazard for rcp \n" ++ "v_mul_f32 %[v_y], %[v_tmp], %[v_x] ; x * 1/(emu+1f)\n" ++ : [v_y] "=v"(y), [v_tmp] "+v"(tmp) ++ : [v_x] "v"(x), [s_c1] "s"(c1), [v_c2] "v"(c2), [s_log2e] "s"(log2e_) ++ :); ++} ++#else ++template <> ++CK_TILE_DEVICE void FastGeluAsm::operator()(float& y, const float& x) const ++{ + // const float u = -2.f * x * (0.035677f * x * x + 0.797885f); + const float c1 = -2.0 * 0.035677f; + const float c2 = -2.0 * 0.797885f; + const float u = x * (c1 * x * x + c2); + const float emu = exp(u); + y = x / (1.f + emu); +- } ++} ++#endif + +- // device code, use lower precision "__ocml_exp_f32" and "rcp" +- template <> +- CK_TILE_DEVICE void operator()(float& y, const float& x) const +- { +- const uint32_t c1 = 0xbd92220c; // -2.0 * 0.035677f; +- const float c2 = -2.0 * 0.797885f; +- const uint32_t log2e_ = 0x3fb8aa3b; // log2e_v; +- float tmp; ++#if defined(__HIP_DEVICE_COMPILE__) ++// Packed implementation interleaves the two lanes to reduce instruction hazards. ++template <> ++CK_TILE_DEVICE void FastGeluAsm::operator()(fp32x2_t& y, ++ const fp32x2_t& x) const ++{ ++ const uint32_t c1 = 0xbd92220c; // -2.0 * 0.035677f; ++ float c2 = -2.0 * 0.797885f; ++ const uint32_t log2e_ = 0x3fb8aa3b; // log2e_v; ++ float tmp0, tmp1; ++ float y0 = x.x, y1 = x.y; + +- asm volatile("v_mul_f32 %[v_tmp], %[v_x], %[v_x] ; x*x\n" +- "v_fma_f32 %[v_tmp], %[v_tmp], %[s_c1], %[v_c2] ; c1*x*x+c2\n" +- "v_mul_f32 %[v_tmp], %[v_tmp], %[v_x] ; x*(c1*x*x+c2)\n" +- "v_mul_f32 %[v_tmp], %[v_tmp], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" +- "v_exp_f32 %[v_tmp], %[v_tmp] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" +- "s_nop 0 ; hazard for exp\n" +- "v_add_f32 %[v_tmp], %[v_tmp], 1.0 ; emu+1.0f\n" +- "v_rcp_f32 %[v_tmp], %[v_tmp] ; 1/(emu+1.0f)\n" +- "s_nop 0 ; hazard for rcp \n" +- "v_mul_f32 %[v_y], %[v_tmp], %[v_x] ; x * 1/(emu+1f)\n" +- : [v_y] "=v"(y), [v_tmp] "+v"(tmp) +- : [v_x] "v"(x), [s_c1] "s"(c1), [v_c2] "v"(c2), [s_log2e] "s"(log2e_) +- :); +- } +- +- template <> +- CK_TILE_HOST void operator()(fp32x2_t& y, const fp32x2_t& x) const +- { ++ asm volatile( ++ "v_mul_f32 %[v_tmp0], %[v_y0], %[v_y0] ; x*x\n" ++ "v_mul_f32 %[v_tmp1], %[v_y1], %[v_y1] ; x*x\n" ++ "v_fma_f32 %[v_tmp0], %[v_tmp0], %[s_c1], %[v_c2] ; c1*x*x+c2\n" ++ "v_fma_f32 %[v_tmp1], %[v_tmp1], %[s_c1], %[v_c2] ; c1*x*x+c2\n" ++ "v_mul_f32 %[v_tmp0], %[v_tmp0], %[v_y0] ; x*(c1*x*x+c2)\n" ++ "v_mul_f32 %[v_tmp1], %[v_tmp1], %[v_y1] ; x*(c1*x*x+c2)\n" ++ "v_mul_f32 %[v_tmp0], %[v_tmp0], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" ++ "v_mul_f32 %[v_tmp1], %[v_tmp1], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" ++ "v_exp_f32 %[v_tmp0], %[v_tmp0] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" ++ "v_exp_f32 %[v_tmp1], %[v_tmp1] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" ++ "v_add_f32 %[v_tmp0], %[v_tmp0], 1.0 ; emu+1.0f\n" ++ "v_add_f32 %[v_tmp1], %[v_tmp1], 1.0 ; emu+1.0f\n" ++ "v_rcp_f32 %[v_tmp0], %[v_tmp0] ; 1/(emu+1.0f)\n" ++ "v_rcp_f32 %[v_tmp1], %[v_tmp1] ; 1/(emu+1.0f)\n" ++ "v_mul_f32 %[v_y0], %[v_tmp0], %[v_y0] ; x * 1/(emu+1f)\n" ++ "v_mul_f32 %[v_y1], %[v_tmp1], %[v_y1] ; x * 1/(emu+1f)\n" ++ : [v_y0] "+v"(y0), ++ [v_y1] "+v"(y1), ++ [v_c2] "+v"(c2), ++ // c2, y0, and y1 may otherwise share a register. Mark the temporary operands ++ // read-write so the compiler allocates distinct registers for the asm block. ++ [v_tmp0] "+v"(tmp0), ++ [v_tmp1] "+v"(tmp1) ++ : [s_c1] "s"(c1), [s_log2e] "s"(log2e_) ++ :); ++ y.x = y0; ++ y.y = y1; ++} ++#else ++template <> ++CK_TILE_DEVICE void FastGeluAsm::operator()(fp32x2_t& y, ++ const fp32x2_t& x) const ++{ + const float c1 = -2.0 * 0.035677f; + const float c2 = -2.0 * 0.797885f; + const float u0 = x.x * (c1 * x.x * x.x + c2); + const float emu1 = exp(u1); + y.y = x.y / (1.f + emu1); +- } +- +- // this is packed verion to remove data hazard for trans +- template <> +- CK_TILE_DEVICE void operator()(fp32x2_t& y, const fp32x2_t& x) const +- { +- const uint32_t c1 = 0xbd92220c; // -2.0 * 0.035677f; +- float c2 = -2.0 * 0.797885f; +- const uint32_t log2e_ = 0x3fb8aa3b; // log2e_v; +- float tmp0, tmp1; +- float y0 = x.x, y1 = x.y; +- +- asm volatile( +- "v_mul_f32 %[v_tmp0], %[v_y0], %[v_y0] ; x*x\n" +- "v_mul_f32 %[v_tmp1], %[v_y1], %[v_y1] ; x*x\n" +- "v_fma_f32 %[v_tmp0], %[v_tmp0], %[s_c1], %[v_c2] ; c1*x*x+c2\n" +- "v_fma_f32 %[v_tmp1], %[v_tmp1], %[s_c1], %[v_c2] ; c1*x*x+c2\n" +- "v_mul_f32 %[v_tmp0], %[v_tmp0], %[v_y0] ; x*(c1*x*x+c2)\n" +- "v_mul_f32 %[v_tmp1], %[v_tmp1], %[v_y1] ; x*(c1*x*x+c2)\n" +- "v_mul_f32 %[v_tmp0], %[v_tmp0], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" +- "v_mul_f32 %[v_tmp1], %[v_tmp1], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" +- "v_exp_f32 %[v_tmp0], %[v_tmp0] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" +- "v_exp_f32 %[v_tmp1], %[v_tmp1] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" +- "v_add_f32 %[v_tmp0], %[v_tmp0], 1.0 ; emu+1.0f\n" +- "v_add_f32 %[v_tmp1], %[v_tmp1], 1.0 ; emu+1.0f\n" +- "v_rcp_f32 %[v_tmp0], %[v_tmp0] ; 1/(emu+1.0f)\n" +- "v_rcp_f32 %[v_tmp1], %[v_tmp1] ; 1/(emu+1.0f)\n" +- "v_mul_f32 %[v_y0], %[v_tmp0], %[v_y0] ; x * 1/(emu+1f)\n" +- "v_mul_f32 %[v_y1], %[v_tmp1], %[v_y1] ; x * 1/(emu+1f)\n" +- : [v_y0] "+v"(y0), +- [v_y1] "+v"(y1), +- [v_c2] "+v"(c2), +- // NOTE! it is totally possible that c2/y0/y1 share same register, they are all local +- // tmp variables we need to expicitly hint compiler they may read+write, to allow +- // allocate different register , the side effect is c2=** may issue for every such +- // inline asm block +- [v_tmp0] "+v"(tmp0), +- [v_tmp1] "+v"(tmp1) +- : [s_c1] "s"(c1), [s_log2e] "s"(log2e_) +- :); +- y.x = y0; +- y.y = y1; +- } +-}; ++} ++#endif + + // https://paperswithcode.com/method/method/gelu + // y = 0.5*x*(1+erf(x/sqrt(2))) From b50e752cc551dda97c6c61d615ad25f1de1ef9f3 Mon Sep 17 00:00:00 2001 From: Ameil Kumar Date: Mon, 17 Aug 2026 11:49:19 -0700 Subject: [PATCH 2/7] Fix fastgeluasm patch: make required, fail loudly if missing --- p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh b/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh index 02e23656c4..f8577694d4 100644 --- a/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh +++ b/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh @@ -199,9 +199,9 @@ wget -q --spider "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patc git apply "$SCRIPT_DIR/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" || echo "pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch missing or already applied, skipped" # Fix FastGeluAsm explicit specializations rejected by AMD clang 23.0 in composable_kernel -wget -q --spider "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_fastgeluasm.patch" && \ - wget -q -O "$SCRIPT_DIR/pytorch_v2.13.0_rocm_fastgeluasm.patch" "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_fastgeluasm.patch" && \ - git apply --directory=third_party/composable_kernel "$SCRIPT_DIR/pytorch_v2.13.0_rocm_fastgeluasm.patch" || echo "pytorch_v2.13.0_rocm_fastgeluasm.patch missing or already applied, skipped" +# This patch is required — fail loudly if it cannot be downloaded or applied +wget -q -O "$SCRIPT_DIR/pytorch_v2.13.0_rocm_fastgeluasm.patch" "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_fastgeluasm.patch" +git apply --directory=third_party/composable_kernel "$SCRIPT_DIR/pytorch_v2.13.0_rocm_fastgeluasm.patch" # Build echo "Building PyTorch (this will take a while)" From 0a3a65c83cf4d601e61c715655f8bc188c7962c2 Mon Sep 17 00:00:00 2001 From: Ameil Kumar Date: Tue, 18 Aug 2026 01:21:15 -0500 Subject: [PATCH 3/7] Make cuda_guard_narrowing patch required, fail loudly if missing --- p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh b/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh index f8577694d4..14266f207c 100644 --- a/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh +++ b/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh @@ -194,9 +194,9 @@ echo "Applying patches" PATCH_BASE_URL=${PATCH_BASE_URL:-"https://raw.githubusercontent.com/ppc64le/build-scripts/refs/heads/master/p/pytorch"} # Fix CUDAGuard narrowing conversion errors under GCC 14 in ROCm HIP flash-attn files -wget -q --spider "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" && \ - wget -q -O "$SCRIPT_DIR/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" && \ - git apply "$SCRIPT_DIR/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" || echo "pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch missing or already applied, skipped" +# This patch is required — fail loudly if it cannot be downloaded or applied +wget -q -O "$SCRIPT_DIR/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" +git apply "$SCRIPT_DIR/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" # Fix FastGeluAsm explicit specializations rejected by AMD clang 23.0 in composable_kernel # This patch is required — fail loudly if it cannot be downloaded or applied From db60b435a7323f012a63ccb0c96b8321e6020261 Mon Sep 17 00:00:00 2001 From: Ameil Kumar Date: Tue, 18 Aug 2026 07:15:54 -0500 Subject: [PATCH 4/7] Fix corrupt index line in fastgeluasm patch --- p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch | 1 - 1 file changed, 1 deletion(-) diff --git a/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch b/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch index 452c238a22..0fbc8c346e 100644 --- a/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch +++ b/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch @@ -1,5 +1,4 @@ diff --git a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp b/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp -index 4e971649d..000000000 100644 --- a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp +++ b/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp @@ -1154,94 +1154,97 @@ From 124a4646245627117809653d618c3b539a091b5b Mon Sep 17 00:00:00 2001 From: Ameil Kumar Date: Tue, 18 Aug 2026 07:47:17 -0500 Subject: [PATCH 5/7] Replace patch files with correct diffs from actual v2.13.0 source --- ...ch_v2.13.0_rocm_cuda_guard_narrowing.patch | 2 +- .../pytorch_v2.13.0_rocm_fastgeluasm.patch | 206 ++++++++++-------- 2 files changed, 117 insertions(+), 91 deletions(-) diff --git a/p/pytorch/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch b/p/pytorch/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch index 550b32fb98..7be06d35fb 100644 --- a/p/pytorch/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch +++ b/p/pytorch/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch @@ -86,4 +86,4 @@ index eb4a3b841fa3e..4340cde7177a0 100644 + at::cuda::CUDAGuard device_guard{q.get_device()}; auto opts = q.options(); - bool has_lse = true; \ No newline at end of file + bool has_lse = true; diff --git a/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch b/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch index 0fbc8c346e..92d74285ec 100644 --- a/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch +++ b/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch @@ -1,7 +1,8 @@ diff --git a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp b/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp +index 4e971649d..7d76097cf 100644 --- a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp +++ b/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp -@@ -1154,94 +1154,97 @@ +@@ -1154,97 +1154,109 @@ struct FastGeluAsm template CK_TILE_DEVICE void operator()(Y& y, const X& x) const; @@ -10,42 +11,16 @@ diff --git a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp b/ - template <> - CK_TILE_HOST void operator()(float& y, const float& x) const - { -+#if defined(__HIP_DEVICE_COMPILE__) -+template <> -+CK_TILE_DEVICE void FastGeluAsm::operator()(float& y, const float& x) const -+{ -+ const uint32_t c1 = 0xbd92220c; // -2.0 * 0.035677f; -+ const float c2 = -2.0 * 0.797885f; -+ const uint32_t log2e_ = 0x3fb8aa3b; // log2e_v; -+ float tmp; -+ -+ asm volatile("v_mul_f32 %[v_tmp], %[v_x], %[v_x] ; x*x\n" -+ "v_fma_f32 %[v_tmp], %[v_tmp], %[s_c1], %[v_c2] ; c1*x*x+c2\n" -+ "v_mul_f32 %[v_tmp], %[v_tmp], %[v_x] ; x*(c1*x*x+c2)\n" -+ "v_mul_f32 %[v_tmp], %[v_tmp], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" -+ "v_exp_f32 %[v_tmp], %[v_tmp] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" -+ "s_nop 0 ; hazard for exp\n" -+ "v_add_f32 %[v_tmp], %[v_tmp], 1.0 ; emu+1.0f\n" -+ "v_rcp_f32 %[v_tmp], %[v_tmp] ; 1/(emu+1.0f)\n" -+ "s_nop 0 ; hazard for rcp \n" -+ "v_mul_f32 %[v_y], %[v_tmp], %[v_x] ; x * 1/(emu+1f)\n" -+ : [v_y] "=v"(y), [v_tmp] "+v"(tmp) -+ : [v_x] "v"(x), [s_c1] "s"(c1), [v_c2] "v"(c2), [s_log2e] "s"(log2e_) -+ :); -+} -+#else -+template <> -+CK_TILE_DEVICE void FastGeluAsm::operator()(float& y, const float& x) const -+{ - // const float u = -2.f * x * (0.035677f * x * x + 0.797885f); - const float c1 = -2.0 * 0.035677f; - const float c2 = -2.0 * 0.797885f; - const float u = x * (c1 * x * x + c2); - const float emu = exp(u); - y = x / (1.f + emu); +- // const float u = -2.f * x * (0.035677f * x * x + 0.797885f); +- const float c1 = -2.0 * 0.035677f; +- const float c2 = -2.0 * 0.797885f; +- const float u = x * (c1 * x * x + c2); +- const float emu = exp(u); +- y = x / (1.f + emu); - } -+} -+#endif ++// Explicit specializations must be defined outside the class body (C++ standard §14.7.3). ++// Defining them inside caused clang to treat enclosing-function locals as captures, ++// breaking the "s" inline-asm register constraints for c1 and log2e_. - // device code, use lower precision "__ocml_exp_f32" and "rcp" - template <> @@ -55,18 +30,7 @@ diff --git a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp b/ - const float c2 = -2.0 * 0.797885f; - const uint32_t log2e_ = 0x3fb8aa3b; // log2e_v; - float tmp; -+#if defined(__HIP_DEVICE_COMPILE__) -+// Packed implementation interleaves the two lanes to reduce instruction hazards. -+template <> -+CK_TILE_DEVICE void FastGeluAsm::operator()(fp32x2_t& y, -+ const fp32x2_t& x) const -+{ -+ const uint32_t c1 = 0xbd92220c; // -2.0 * 0.035677f; -+ float c2 = -2.0 * 0.797885f; -+ const uint32_t log2e_ = 0x3fb8aa3b; // log2e_v; -+ float tmp0, tmp1; -+ float y0 = x.x, y1 = x.y; - +- - asm volatile("v_mul_f32 %[v_tmp], %[v_x], %[v_x] ; x*x\n" - "v_fma_f32 %[v_tmp], %[v_tmp], %[s_c1], %[v_c2] ; c1*x*x+c2\n" - "v_mul_f32 %[v_tmp], %[v_tmp], %[v_x] ; x*(c1*x*x+c2)\n" @@ -81,51 +45,59 @@ diff --git a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp b/ - : [v_x] "v"(x), [s_c1] "s"(c1), [v_c2] "v"(c2), [s_log2e] "s"(log2e_) - :); - } -- ++template <> ++CK_TILE_HOST void FastGeluAsm::operator()(float& y, const float& x) const ++{ ++ // const float u = -2.f * x * (0.035677f * x * x + 0.797885f); ++ const float c1 = -2.0 * 0.035677f; ++ const float c2 = -2.0 * 0.797885f; ++ const float u = x * (c1 * x * x + c2); ++ const float emu = exp(u); ++ y = x / (1.f + emu); ++} + - template <> - CK_TILE_HOST void operator()(fp32x2_t& y, const fp32x2_t& x) const - { -+ asm volatile( -+ "v_mul_f32 %[v_tmp0], %[v_y0], %[v_y0] ; x*x\n" -+ "v_mul_f32 %[v_tmp1], %[v_y1], %[v_y1] ; x*x\n" -+ "v_fma_f32 %[v_tmp0], %[v_tmp0], %[s_c1], %[v_c2] ; c1*x*x+c2\n" -+ "v_fma_f32 %[v_tmp1], %[v_tmp1], %[s_c1], %[v_c2] ; c1*x*x+c2\n" -+ "v_mul_f32 %[v_tmp0], %[v_tmp0], %[v_y0] ; x*(c1*x*x+c2)\n" -+ "v_mul_f32 %[v_tmp1], %[v_tmp1], %[v_y1] ; x*(c1*x*x+c2)\n" -+ "v_mul_f32 %[v_tmp0], %[v_tmp0], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" -+ "v_mul_f32 %[v_tmp1], %[v_tmp1], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" -+ "v_exp_f32 %[v_tmp0], %[v_tmp0] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" -+ "v_exp_f32 %[v_tmp1], %[v_tmp1] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" -+ "v_add_f32 %[v_tmp0], %[v_tmp0], 1.0 ; emu+1.0f\n" -+ "v_add_f32 %[v_tmp1], %[v_tmp1], 1.0 ; emu+1.0f\n" -+ "v_rcp_f32 %[v_tmp0], %[v_tmp0] ; 1/(emu+1.0f)\n" -+ "v_rcp_f32 %[v_tmp1], %[v_tmp1] ; 1/(emu+1.0f)\n" -+ "v_mul_f32 %[v_y0], %[v_tmp0], %[v_y0] ; x * 1/(emu+1f)\n" -+ "v_mul_f32 %[v_y1], %[v_tmp1], %[v_y1] ; x * 1/(emu+1f)\n" -+ : [v_y0] "+v"(y0), -+ [v_y1] "+v"(y1), -+ [v_c2] "+v"(c2), -+ // c2, y0, and y1 may otherwise share a register. Mark the temporary operands -+ // read-write so the compiler allocates distinct registers for the asm block. -+ [v_tmp0] "+v"(tmp0), -+ [v_tmp1] "+v"(tmp1) -+ : [s_c1] "s"(c1), [s_log2e] "s"(log2e_) -+ :); -+ y.x = y0; -+ y.y = y1; -+} -+#else +- const float c1 = -2.0 * 0.035677f; +- const float c2 = -2.0 * 0.797885f; +- const float u0 = x.x * (c1 * x.x * x.x + c2); +- const float emu0 = exp(u0); +- y.x = x.x / (1.f + emu0); +- const float u1 = x.y * (c1 * x.y * x.y + c2); +- const float emu1 = exp(u1); +- y.y = x.y / (1.f + emu1); +- } ++// device code, use lower precision "__ocml_exp_f32" and "rcp" ++// Guard with __HIP_DEVICE_COMPILE__: in the host pass CK_TILE_DEVICE expands to plain ++// inline, making this a duplicate definition of the host specialization above. The ++// duplicate causes the compiler to treat the asm operands as references from an ++// "enclosing function", breaking the "s" (SGPR) register constraints. ++#ifdef __HIP_DEVICE_COMPILE__ +template <> -+CK_TILE_DEVICE void FastGeluAsm::operator()(fp32x2_t& y, -+ const fp32x2_t& x) const ++CK_TILE_DEVICE void FastGeluAsm::operator()(float& y, const float& x) const +{ - const float c1 = -2.0 * 0.035677f; - const float c2 = -2.0 * 0.797885f; - const float u0 = x.x * (c1 * x.x * x.x + c2); - const float emu1 = exp(u1); - y.y = x.y / (1.f + emu1); -- } -- ++ const uint32_t c1 = 0xbd92220c; // -2.0 * 0.035677f; ++ const float c2 = -2.0 * 0.797885f; ++ const uint32_t log2e_ = 0x3fb8aa3b; // log2e_v; ++ float tmp; ++ ++ asm volatile("v_mul_f32 %[v_tmp], %[v_x], %[v_x] ; x*x\n" ++ "v_fma_f32 %[v_tmp], %[v_tmp], %[s_c1], %[v_c2] ; c1*x*x+c2\n" ++ "v_mul_f32 %[v_tmp], %[v_tmp], %[v_x] ; x*(c1*x*x+c2)\n" ++ "v_mul_f32 %[v_tmp], %[v_tmp], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" ++ "v_exp_f32 %[v_tmp], %[v_tmp] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" ++ "s_nop 0 ; hazard for exp\n" ++ "v_add_f32 %[v_tmp], %[v_tmp], 1.0 ; emu+1.0f\n" ++ "v_rcp_f32 %[v_tmp], %[v_tmp] ; 1/(emu+1.0f)\n" ++ "s_nop 0 ; hazard for rcp \n" ++ "v_mul_f32 %[v_y], %[v_tmp], %[v_x] ; x * 1/(emu+1f)\n" ++ : [v_y] "=v"(y), [v_tmp] "+v"(tmp) ++ : [v_x] "v"(x), [s_c1] "s"(c1), [v_c2] "v"(c2), [s_log2e] "s"(log2e_) ++ :); ++} ++#endif // __HIP_DEVICE_COMPILE__ + - // this is packed verion to remove data hazard for trans - template <> - CK_TILE_DEVICE void operator()(fp32x2_t& y, const fp32x2_t& x) const @@ -168,8 +140,62 @@ diff --git a/include/ck_tile/ops/elementwise/unary_element_wise_operation.hpp b/ - y.y = y1; - } -}; ++template <> ++CK_TILE_HOST void FastGeluAsm::operator()(fp32x2_t& y, const fp32x2_t& x) const ++{ ++ const float c1 = -2.0 * 0.035677f; ++ const float c2 = -2.0 * 0.797885f; ++ const float u0 = x.x * (c1 * x.x * x.x + c2); ++ const float emu0 = exp(u0); ++ y.x = x.x / (1.f + emu0); ++ const float u1 = x.y * (c1 * x.y * x.y + c2); ++ const float emu1 = exp(u1); ++ y.y = x.y / (1.f + emu1); ++} ++ ++// this is packed verion to remove data hazard for trans ++#ifdef __HIP_DEVICE_COMPILE__ ++template <> ++CK_TILE_DEVICE void FastGeluAsm::operator()(fp32x2_t& y, const fp32x2_t& x) const ++{ ++ const uint32_t c1 = 0xbd92220c; // -2.0 * 0.035677f; ++ float c2 = -2.0 * 0.797885f; ++ const uint32_t log2e_ = 0x3fb8aa3b; // log2e_v; ++ float tmp0, tmp1; ++ float y0 = x.x, y1 = x.y; ++ ++ asm volatile( ++ "v_mul_f32 %[v_tmp0], %[v_y0], %[v_y0] ; x*x\n" ++ "v_mul_f32 %[v_tmp1], %[v_y1], %[v_y1] ; x*x\n" ++ "v_fma_f32 %[v_tmp0], %[v_tmp0], %[s_c1], %[v_c2] ; c1*x*x+c2\n" ++ "v_fma_f32 %[v_tmp1], %[v_tmp1], %[s_c1], %[v_c2] ; c1*x*x+c2\n" ++ "v_mul_f32 %[v_tmp0], %[v_tmp0], %[v_y0] ; x*(c1*x*x+c2)\n" ++ "v_mul_f32 %[v_tmp1], %[v_tmp1], %[v_y1] ; x*(c1*x*x+c2)\n" ++ "v_mul_f32 %[v_tmp0], %[v_tmp0], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" ++ "v_mul_f32 %[v_tmp1], %[v_tmp1], %[s_log2e] ; log2e*x*(c1*x*x+c2)\n" ++ "v_exp_f32 %[v_tmp0], %[v_tmp0] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" ++ "v_exp_f32 %[v_tmp1], %[v_tmp1] ; emu = exp2(log2e*x*(c1*x*x+c2))\n" ++ "v_add_f32 %[v_tmp0], %[v_tmp0], 1.0 ; emu+1.0f\n" ++ "v_add_f32 %[v_tmp1], %[v_tmp1], 1.0 ; emu+1.0f\n" ++ "v_rcp_f32 %[v_tmp0], %[v_tmp0] ; 1/(emu+1.0f)\n" ++ "v_rcp_f32 %[v_tmp1], %[v_tmp1] ; 1/(emu+1.0f)\n" ++ "v_mul_f32 %[v_y0], %[v_tmp0], %[v_y0] ; x * 1/(emu+1f)\n" ++ "v_mul_f32 %[v_y1], %[v_tmp1], %[v_y1] ; x * 1/(emu+1f)\n" ++ : [v_y0] "+v"(y0), ++ [v_y1] "+v"(y1), ++ [v_c2] "+v"(c2), ++ // NOTE! it is totally possible that c2/y0/y1 share same register, they are all local ++ // tmp variables we need to expicitly hint compiler they may read+write, to allow ++ // allocate different register , the side effect is c2=** may issue for every such ++ // inline asm block ++ [v_tmp0] "+v"(tmp0), ++ [v_tmp1] "+v"(tmp1) ++ : [s_c1] "s"(c1), [s_log2e] "s"(log2e_) ++ :); ++ y.x = y0; ++ y.y = y1; +} -+#endif ++#endif // __HIP_DEVICE_COMPILE__ - // https://paperswithcode.com/method/method/gelu + // https://paperswithcode.com/method/gelu // y = 0.5*x*(1+erf(x/sqrt(2))) From d2796e02f0b4349126d21a2dd387f08bd3c02795 Mon Sep 17 00:00:00 2001 From: Deepali Patel Date: Wed, 19 Aug 2026 06:22:42 +0000 Subject: [PATCH 6/7] update script to use patch from script_dir --- p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh b/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh index 14266f207c..a3d758e2f0 100644 --- a/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh +++ b/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh @@ -195,17 +195,17 @@ PATCH_BASE_URL=${PATCH_BASE_URL:-"https://raw.githubusercontent.com/ppc64le/buil # Fix CUDAGuard narrowing conversion errors under GCC 14 in ROCm HIP flash-attn files # This patch is required — fail loudly if it cannot be downloaded or applied -wget -q -O "$SCRIPT_DIR/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" +#wget -q -O "$SCRIPT_DIR/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" git apply "$SCRIPT_DIR/pytorch_v2.13.0_rocm_cuda_guard_narrowing.patch" # Fix FastGeluAsm explicit specializations rejected by AMD clang 23.0 in composable_kernel # This patch is required — fail loudly if it cannot be downloaded or applied -wget -q -O "$SCRIPT_DIR/pytorch_v2.13.0_rocm_fastgeluasm.patch" "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_fastgeluasm.patch" +#wget -q -O "$SCRIPT_DIR/pytorch_v2.13.0_rocm_fastgeluasm.patch" "$PATCH_BASE_URL/pytorch_v2.13.0_rocm_fastgeluasm.patch" git apply --directory=third_party/composable_kernel "$SCRIPT_DIR/pytorch_v2.13.0_rocm_fastgeluasm.patch" # Build echo "Building PyTorch (this will take a while)" -export PYTORCH_BUILD_VERSION=${PACKAGE_VERSION#v}+rocm +export PYTORCH_BUILD_VERSION=${PACKAGE_VERSION#v}+rocm+714 export PYTORCH_BUILD_NUMBER=1 if ! MAX_JOBS=$(nproc) $PYTHON -m pip install --no-build-isolation -v -e .; then From 0f4c9aa57bb5d4b0abda83a2f8b2a9506fd29cc0 Mon Sep 17 00:00:00 2001 From: Deepali Patel Date: Wed, 19 Aug 2026 11:20:39 +0000 Subject: [PATCH 7/7] fix version --- p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh b/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh index a3d758e2f0..0d8e0de35f 100644 --- a/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh +++ b/p/pytorch/pytorch_2.13.0_rocm_ubi_10.sh @@ -205,7 +205,7 @@ git apply --directory=third_party/composable_kernel "$SCRIPT_DIR/pytorch_v2.13.0 # Build echo "Building PyTorch (this will take a while)" -export PYTORCH_BUILD_VERSION=${PACKAGE_VERSION#v}+rocm+714 +export PYTORCH_BUILD_VERSION=${PACKAGE_VERSION#v}+rocm714 export PYTORCH_BUILD_NUMBER=1 if ! MAX_JOBS=$(nproc) $PYTHON -m pip install --no-build-isolation -v -e .; then