Freebuff/rdna2 native extension optimization a36e947c 679f 48b5 8290 8a6f1c09ab4f

.github/workflows/rdna2_level3.yml

@@ -0,0 +1,120 @@
+name: rdna2 level-3 bridge test
+
+# Builds the library and runs gpu/rdna2/test_oblas_gpu.c, which checks the CBLAS
+# level-3 entry points the rdna2 bridges serve -- SGEMM, SYMM, SYRK, SYR2K, TRMM
+# and TRSM -- against double-precision references the test computes itself. It
+# covers every order/side/uplo/trans/diag combination, alpha and beta including
+# zero, negative and large values, and the contracts the bridges owe the caller:
+# the triangle they were not asked to store, the stored diagonal of a
+# unit-diagonal factor and the ldb/ldc padding have to come back untouched.
+# The per-op comparisons treat a non-finite result as a failure instead of
+# letting it past a tolerance test that is always false for NaN. Nothing here
+# needs a GPU: on a CPU-only build the same cases still check the values and
+# those contracts. The dispatch itself is only exercised by a build made with
+# USE_HIP=1 on a machine with an AMD card, which hosted runners do not have --
+# the "gpu" job at the bottom is the place to add that on a self-hosted runner.
+#
+# Two gaps are written down here rather than left implicit. No case feeds a NaN
+# or an infinity into an input yet, so the non-finite handling above guards a
+# result that arithmetic produced; it is not a tested scenario, and the SGEMM
+# block still judges by max |diff| alone, where a NaN result passes silently.
+# The path filters below also omit interface/symm.c, interface/trsm.c (TRMM and
+# TRSM) and interface/gemm.c, so an edit to one of those bridge files does not
+# start this workflow on its own.
+
+on:
+  push:
+    paths:
+      - 'gpu/rdna2/**'
+      - 'kernel/rdna2/**'
+      - 'interface/syrk.c'
+      - 'interface/syr2k.c'
+      - 'interface/gemm_i8.c'
+      - 'cblas.h'
+      - 'Makefile'
+      - 'Makefile.system'
+      - 'Makefile.rule'
+      - '.github/workflows/rdna2_level3.yml'
+  pull_request:
+    paths:
+      - 'gpu/rdna2/**'
+      - 'kernel/rdna2/**'
+      - 'interface/syrk.c'
+      - 'interface/syr2k.c'
+      - 'interface/gemm_i8.c'
+      - 'cblas.h'
+      - 'Makefile'
+      - 'Makefile.system'
+      - 'Makefile.rule'
+      - '.github/workflows/rdna2_level3.yml'
+  workflow_dispatch:
+
+concurrency:
+  group: ${{ github.workflow }}-${{ github.head_ref || github.run_id }}
+  cancel-in-progress: true
+
+permissions:
+  contents: read
+
+jobs:
+  linux:
+    name: bridge test (linux, CPU build)
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@v5
+
+      - name: Install toolchain
+        run: |
+          sudo apt-get update
+          sudo apt-get install -y build-essential gfortran
+
+      # NOFORTRAN/NO_LAPACK keep the job to what the bridge test needs: CBLAS
+      # level-3 entry points and the kernels behind them.
+      - name: Build library
+        run: make -j"$(nproc)" NOFORTRAN=1 NO_LAPACK=1 NO_LAPACKE=1 libs
+
+      - name: Level-3 bridge test
+        run: make NOFORTRAN=1 NO_LAPACK=1 NO_LAPACKE=1 level3-test
+
+  windows:
+    name: bridge test (windows, CPU build)
+    runs-on: windows-latest
+    defaults:
+      run:
+        shell: msys2 {0}
+    steps:
+      - uses: actions/checkout@v5
+
+      - uses: msys2/setup-msys2@v2
+        with:
+          msystem: UCRT64
+          update: false
+          install: >-
+            base-devel
+            make
+            mingw-w64-ucrt-x86_64-gcc
+
+      - name: Build library
+        run: make -j"$(nproc)" CC=gcc NOFORTRAN=1 NO_LAPACK=1 NO_LAPACKE=1 libs
+
+      - name: Level-3 bridge test
+        run: make CC=gcc NOFORTRAN=1 NO_LAPACK=1 NO_LAPACKE=1 level3-test
+
+  # Device coverage needs an AMD GPU and a ROCm install, which GitHub-hosted
+  # runners do not provide. On a self-hosted runner carrying one, this job is
+  # the same two commands with the HIP build switched on:
+  #
+  # gpu:
+  #   name: bridge test (self-hosted, USE_HIP)
+  #   runs-on: [self-hosted, gpu, rocm]
+  #   steps:
+  #     - uses: actions/checkout@v5
+  #     - name: Build library with the device kernels
+  #       run: |
+  #         make clean
+  #         make -j"$(nproc)" USE_HIP=1 GPU_ARCH=gfx1201 libs
+  #     - name: Level-3 bridge test (asserts the dispatch actually fires)
+  #       run: |
+  #         make USE_HIP=1 GPU_ARCH=gfx1201 level3-test
+  #         make -C gpu/rdna2 oblas-test OPENBLAS_RDNA2_GPU_DEBUG=1 2>&1 | tee gpu.log
+  #         grep -q 'handled=1' gpu.log

.gitignore

@@ -8,6 +8,7 @@
 *.wasm
 *.out
 *.tmp
+*.stamp
 lapack-3.1.1
 lapack-3.1.1.tgz
 lapack-3.4.1

Makefile

@@ -52,7 +52,7 @@ endif
 
 SUBDIRS_ALL = $(SUBDIRS) test ctest utest exports benchmark ../laswp ../bench cpp_thread_test
 
-.PHONY : all libs netlib $(RELA) test ctest shared install
+.PHONY : all libs netlib $(RELA) test ctest shared install level3-test
 .NOTPARALLEL : shared
 
 all :: tests
@@ -171,6 +171,15 @@ ifeq ($(OSNAME), AIX)
 endif
 endif
 
+# Level-3 bridge test: gpu/rdna2/test_oblas_gpu.c checks the CBLAS level-3
+# entry points against its own double-precision reference, including the SYRK
+# contract that the unstored triangle and the ldc padding are left alone. It
+# needs nothing beyond a built library, so it also runs on CPU-only builds; when
+# the library was built with USE_HIP=1 it additionally exercises the device
+# dispatch, which is the case it exists for.
+level3-test : libs
+	@$(MAKE) -C gpu/rdna2 oblas-test
+
 tests : shared
 ifeq ($(NOFORTRAN), $(filter 0,$(NOFORTRAN)))
 	touch $(LIBNAME)

Makefile.gfx1030

@@ -0,0 +1 @@
+include Makefile.rdna2

Makefile.gfx1031

@@ -0,0 +1 @@
+include Makefile.rdna2

Makefile.gfx1200

@@ -0,0 +1 @@
+include Makefile.rdna2

Makefile.gfx1201

@@ -0,0 +1 @@
+include Makefile.rdna2

Makefile.prebuild

@@ -103,7 +103,7 @@ endif
 getarch : getarch.c cpuid.S dummy $(CPUIDEMU)
 	avx512=$$(./c_check$(SCRIPTSUFFIX) - - "$(CC)" $(TARGET_FLAGS) $(CFLAGS) | grep NO_AVX512); \
 	rv64gv=$$(./c_check$(SCRIPTSUFFIX) - - "$(CC)" $(TARGET_FLAGS) $(CFLAGS) | grep NO_RV64GV); \
-	$(HOSTCC) $(HOST_CFLAGS) $(EXFLAGS) $${avx512:+-D$${avx512}} $${rv64gv:+-D$${rv64gv}} -o $(@F) getarch.c cpuid.S $(CPUIDEMU)
+	$(HOSTCC) $(HOST_CFLAGS) $(EXFLAGS) $${avx512:+-D$${avx512}} $${rv64gv:+-D$${rv64gv}} $(if $(TARGET),-DFORCE_$(TARGET),) -o $(@F) getarch.c cpuid.S $(CPUIDEMU)
 
 getarch_2nd : getarch_2nd.c $(TARGET_CONF) dummy
 ifndef TARGET_CORE

Makefile.rdna2

@@ -0,0 +1,6 @@
+ifeq ($(C_COMPILER), GCC)
+CCOMMON_OPT += -O2
+endif
+ifeq ($(C_COMPILER), CLANG)
+CCOMMON_OPT += -O2
+endif

Makefile.rule

@@ -19,6 +19,20 @@ VERSION = 0.3.34.dev
 # automatically detected.
 # TARGET = PENRYN
 
+# Build the RDNA2/RDNA4 GPU kernels (HIP) alongside the CPU ones. TARGET still
+# selects the CPU kernels exactly as it does without USE_HIP, so USE_HIP=1 on an
+# ordinary target yields one library that is complete on both: large level-3
+# calls go to the device, everything else (small calls, complex, double, the
+# level 1/2 routines and all of LAPACK) stays on the CPU. Set
+# OPENBLAS_RDNA2_GPU=0 in the environment to force the CPU path at run time.
+# USE_HIP = 1
+
+# Device architecture(s) for the HIP kernels, independent of TARGET. Defaults to
+# the architecture of a GFX* TARGET, otherwise gfx1201. A space separated list
+# builds one fat object covering several cards, so a single library serves both
+# RDNA2 and RDNA4, e.g. GPU_ARCH = gfx1031 gfx1201
+# GPU_ARCH = gfx1201
+
 # If you want to support multiple architecture in one binary
 # DYNAMIC_ARCH = 1
 

Makefile.system

@@ -2014,3 +2014,36 @@ SUNPATH		= /opt/sunstudio12.1
 else
 SUNPATH		= /opt/SUNWspro
 endif
+
+# HIP/ROCm support for RDNA2/RDNA4
+ifdef USE_HIP
+HIPCC ?= hipcc
+HIPFLAGS ?= -O2 -std=c++17
+ifeq ($(HIP_PLATFORM),)
+HIP_PLATFORM := amd
+endif
+export HIP_PLATFORM
+CXX ?= $(HIPCC)
+# Device architecture(s) for the RDNA2/RDNA4 kernels, independent of the CPU
+# core, so an ordinary CPU build can carry the GPU side too. Unset means "take
+# it from a GFX* TARGET, else gfx1201".
+export GPU_ARCH
+CCOMMON_OPT += -DOPENBLAS_USE_HIP
+# Locate the ROCm install instead of hardcoding /opt/rocm, the way the CUDA
+# support uses CUDADIR: HIP_PATH from the environment (ROCm sets it) wins, then
+# whatever hipcc reports about itself, then the usual install prefix.
+HIP_PATH ?= $(patsubst %/,%,$(shell $(HIPCC) --path 2>/dev/null))
+ifeq ($(strip $(HIP_PATH)),)
+HIP_PATH := /opt/rocm
+endif
+# Backslashes in a Windows HIP_PATH would be eaten by the shell in the link
+# line, so normalise them here.
+# HIP_PATH and HIP_LIB_PATH can also arrive from the environment (ROCm sets
+# them), so normalise whatever ends up in them, not just the default.
+HIP_LIB_PATH ?= $(HIP_PATH)/lib
+HIP_LIB_PATH := $(subst \,/,$(HIP_LIB_PATH))
+export HIP_LIB_PATH
+# Only the HIP runtime is needed; hiprtc/hipblas are not used by the kernels.
+# The kernels are C++ via hipcc, so the C++ runtime must be linked in as well.
+EXTRALIB += -L$(HIP_LIB_PATH) -lamdhip64 -lstdc++
+endif

README.md

@@ -387,3 +387,7 @@ Please see Changelog.txt.
 ## Donation
 
 Please see [the donations section](http://www.openmathlib.org/OpenBLAS/docs/about/#donations) in the docs.
+
+## AMD RDNA2/RDNA4 GPU Support (HIP)
+
+Full GPU support for AMD RDNA2 (gfx1030/gfx1031) and RDNA4 (gfx1200/gfx1201) via HIP kernels. Build with `USE_HIP=1 TARGET=gfx1201` (or appropriate target). Kernels are in `gpu/rdna2/` and `kernel/rdna2/`. Quantized weights are kept byte-exact (no f32 dequantization of stored weights). Tested on RX 9070 XT (gfx1201).

TargetList.txt

@@ -158,3 +158,9 @@ CK860FV
 15. WebAssembly/Emscripten:
 WASM128_GENERIC
 
+
+# GPU Targets (RDNA2/RDNA4)
+GFX1030
+GFX1031
+GFX1200
+GFX1201

cblas.h

@@ -520,6 +520,69 @@ OPENBLAS_EXPORT void cblas_sbgemm_batch(OPENBLAS_CONST enum CBLAS_ORDER Order, O
 		       OPENBLAS_CONST float * alpha_array, OPENBLAS_CONST bfloat16 ** A_array, OPENBLAS_CONST blasint * lda_array, OPENBLAS_CONST bfloat16 ** B_array, OPENBLAS_CONST blasint * ldb_array, OPENBLAS_CONST float * beta_array, float ** C_array, OPENBLAS_CONST blasint * ldc_array, OPENBLAS_CONST blasint group_count, OPENBLAS_CONST blasint * group_size);
 
 OPENBLAS_EXPORT void cblas_sbgemm_batch_strided(OPENBLAS_CONST enum CBLAS_ORDER Order, OPENBLAS_CONST enum CBLAS_TRANSPOSE TransA, OPENBLAS_CONST enum CBLAS_TRANSPOSE TransB, OPENBLAS_CONST blasint M, OPENBLAS_CONST blasint N, OPENBLAS_CONST blasint K, OPENBLAS_CONST float alpha, OPENBLAS_CONST bfloat16 * A, OPENBLAS_CONST blasint lda, OPENBLAS_CONST blasint stridea, OPENBLAS_CONST bfloat16 * B, OPENBLAS_CONST blasint ldb, OPENBLAS_CONST blasint strideb, OPENBLAS_CONST float beta, float * C, OPENBLAS_CONST blasint ldc, OPENBLAS_CONST blasint stridec, OPENBLAS_CONST blasint group_size);
+/*** INT8 extensions (packed int8 GEMM) ***/
+/* C := op(A) * op(B), with 8-bit integer operands packed four per int32 word.
+ *
+ * These entry points expose the RDNA2/RDNA4 int8 GEMM kernels
+ * (gpu/rdna2/rdna2_kernels.hip) through the normal library, and are also
+ * available in builds without a GPU, where the same product is computed in
+ * portable C. Both paths are exact: accumulation is 32-bit two's complement
+ * (v_dot4_i32_i8), so the two agree bit for bit and no tolerance applies.
+ *
+ * Data contract
+ *   - A and B are arrays of int32 words, each holding four signed int8 values:
+ *     value k of a word is bits [8k, 8k+8), i.e. little-endian byte order. This
+ *     is the layout v_dot4_* consumes, so packed quantized weights are used
+ *     byte for byte as stored.
+ *   - K4 is the number of words per row, so the inner dimension is K = 4*K4
+ *     int8 values. K must be a multiple of 4: a partial trailing word is not
+ *     accepted, and this call performs no repacking of the stored values.
+ *   - lda and ldb are the leading dimensions in words, in the usual BLAS
+ *     sense for the requested order. A is stored M x K4 (TransA ==
+ *     CblasNoTrans) or K4 x M (transposed), B is stored K4 x N or N x K4, and:
+ *       CblasColMajor: lda is the row count of the stored array, so
+ *                      lda >= max(1, K4) when transposed and >= max(1, M)
+ *                      otherwise; ldb >= max(1, N) when transposed and
+ *                      >= max(1, K4) otherwise; ldc >= max(1, M).
+ *       CblasRowMajor: lda is the row stride, so lda >= max(1, M) when
+ *                      transposed and >= max(1, K4) otherwise; ldb >=
+ *                      max(1, K4) when transposed and >= max(1, N)
+ *                      otherwise; ldc >= max(1, N).
+ *   - C is int32 for cblas_gemm_i8 and float for cblas_gemm_i8_scale, with
+ *     ldc in C elements: ldc >= max(1, M) for CblasColMajor and
+ *     ldc >= max(1, N) for CblasRowMajor.
+ *   - Order == CblasRowMajor is supported and gives the same result as the
+ *     column-major call on the transposed problem; CblasConjTrans and
+ *     CblasConjNoTrans are accepted as CblasTrans and CblasNoTrans, since
+ *     conjugation is the identity on integers.
+ *
+ * Scaling (cblas_gemm_i8_scale only)
+ *   C[i][j] = (acc[i][j] - zp_b * rowsum_j) * sa[i] * sb[j]
+ *   where rowsum_j is the sum of the int8 values of column j of op(B) - the
+ *   exact integer form of asymmetric storage, q representing (q - zp_b) - and
+ *   the correction is computed in 64 bits. sa scales rows of op(A) and has M
+ *   entries, sb scales columns of op(B) and has N entries; either may be NULL
+ *   for 1.0. zp_b is the zero point of op(B), i.e. the weight side; 0 skips the
+ *   row-sum pass. No stored operand is converted to floating point: the only
+ *   float produced per output element is the finished dot product.
+ *
+ * There is no alpha/beta in either call. Scale the result on output, or use
+ * cblas_gemm_i8_scale, which applies the per-row/per-column scale for free.
+ * 4-bit formats (Q4_0, Q4_K, ...) are not supported: dp4a consumes whole
+ * bytes, so they would need an in-kernel nibble unpack first.
+ *
+ * Dispatch
+ *   In a build with USE_HIP=1 the HIP kernels run the product;
+ *   OPENBLAS_RDNA2_GPU=0 (or =n) forces the CPU implementation, which is also
+ *   what a build without USE_HIP and any failing or device-less launch uses.
+ *
+ * Returns 0 on success, -1 on an invalid argument or layout, -2 when a host
+ * workspace allocation failed. C is left untouched on error. Both calls are
+ * reentrant; all workspaces are allocated per call.
+ */
+OPENBLAS_EXPORT int cblas_gemm_i8(OPENBLAS_CONST enum CBLAS_ORDER Order, OPENBLAS_CONST enum CBLAS_TRANSPOSE TransA, OPENBLAS_CONST enum CBLAS_TRANSPOSE TransB, OPENBLAS_CONST blasint M, OPENBLAS_CONST blasint N, OPENBLAS_CONST blasint K4, OPENBLAS_CONST int *A, OPENBLAS_CONST blasint lda, OPENBLAS_CONST int *B, OPENBLAS_CONST blasint ldb, int *C, OPENBLAS_CONST blasint ldc);
+OPENBLAS_EXPORT int cblas_gemm_i8_scale(OPENBLAS_CONST enum CBLAS_ORDER Order, OPENBLAS_CONST enum CBLAS_TRANSPOSE TransA, OPENBLAS_CONST enum CBLAS_TRANSPOSE TransB, OPENBLAS_CONST blasint M, OPENBLAS_CONST blasint N, OPENBLAS_CONST blasint K4, OPENBLAS_CONST int *A, OPENBLAS_CONST blasint lda, OPENBLAS_CONST int *B, OPENBLAS_CONST blasint ldb, float *C, OPENBLAS_CONST blasint ldc, OPENBLAS_CONST float *sa, OPENBLAS_CONST float *sb, OPENBLAS_CONST int zp_b);
+
 /*** FLOAT16 extensions ***/
 OPENBLAS_EXPORT void cblas_shgemm(OPENBLAS_CONST enum CBLAS_ORDER Order, OPENBLAS_CONST enum CBLAS_TRANSPOSE TransA, OPENBLAS_CONST enum CBLAS_TRANSPOSE TransB, OPENBLAS_CONST blasint M, OPENBLAS_CONST blasint N, OPENBLAS_CONST blasint K,
 		    OPENBLAS_CONST float alpha, OPENBLAS_CONST hfloat16 *A, OPENBLAS_CONST blasint lda, OPENBLAS_CONST hfloat16 *B, OPENBLAS_CONST blasint ldb, OPENBLAS_CONST float beta, float *C, OPENBLAS_CONST blasint ldc);

driver/others/Makefile

@@ -59,7 +59,7 @@ endif
 endif
 
 ifeq ($(USE_CUDA), 1)
-COMMONOBJS	+= cuda_init.$(SUFFIX)
+$(error USE_CUDA=1 is not supported: driver/others/cuda_init.c is not part of this tree)
 endif
 
 ifdef FUNCTION_PROFILE

exports/gensymbol

@@ -130,7 +130,7 @@ cblasobjsz="
     cblas_zaxpyc cblas_zdrot cblas_zrotg cblas_dzamax cblas_dzamin cblas_zgemm_batch cblas_zgemm_batch_strided
 "
 
-cblasobjs="cblas_xerbla"
+cblasobjs="cblas_xerbla cblas_gemm_i8 cblas_gemm_i8_scale"
 
 bfcblasobjs="cblas_bgemm cblas_bgemv cblas_bscal cblas_sbgemm cblas_sbgemv cblas_sbdot 
     cblas_sbstobf16 cblas_sbdtobf16 cblas_sbf16tos cblas_dbf16tod cblas_sbgemm_batch cblas_sbgemm_batch_strided

exports/gensymbol.pl

@@ -126,7 +126,7 @@
     cblas_izamax, cblas_izamin, cblas_dzsum,cblas_zimatcopy,cblas_zomatcopy,
     cblas_zaxpyc, cblas_zdrot, cblas_zrotg, cblas_dzamax, cblas_dzamin, cblas_zgemm_batch, cblas_zgemm_batch_strided);
 
-@cblasobjs = (  cblas_xerbla );
+@cblasobjs = (  cblas_xerbla, cblas_gemm_i8, cblas_gemm_i8_scale );
 
 @bfcblasobjs = (cblas_bgemm, cblas_bgemv, cblas_bscal, cblas_sbgemm, cblas_sbgemmt, cblas_sbgemmtr, cblas_sbgemv, cblas_sbdot, cblas_sbstobf16, cblas_sbdtobf16, cblas_sbf16tos, cblas_dbf16tod, cblas_sbgemm_batch, cblas_sbgemm_batch_strided);
 @hfcblasobjs = (cblas_shgemm);