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..0d8e0de35f --- /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 +# 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 +#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}+rocm714 +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..7be06d35fb --- /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; 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..92d74285ec --- /dev/null +++ b/p/pytorch/pytorch_v2.13.0_rocm_fastgeluasm.patch @@ -0,0 +1,201 @@ +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,97 +1154,109 @@ struct FastGeluAsm + + template + CK_TILE_DEVICE void operator()(Y& y, const X& x) const; ++}; + +- template <> +- CK_TILE_HOST void 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); +- } ++// 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 <> +- 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; +- +- 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 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 +- { +- 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()(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_) ++ :); ++} ++#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 +- { +- 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; +- } +-}; ++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 // __HIP_DEVICE_COMPILE__ + + // https://paperswithcode.com/method/gelu + // y = 0.5*x*(1+erf(x/sqrt(2)))