diff --git a/.github/workflows/benchmark.yml b/.github/workflows/benchmark.yml new file mode 100644 index 0000000..8a5cdb0 --- /dev/null +++ b/.github/workflows/benchmark.yml @@ -0,0 +1,299 @@ +name: Benchmark + +on: + issue_comment: + types: [created] + workflow_dispatch: + inputs: + warmup: + description: "Warmup iterations" + default: "5" + iterations: + description: "Timed iterations" + default: "20" + sizes: + description: "Comma-separated square GEMM sizes (M=N=K)" + default: "256,512,1024,2048,4096" + +permissions: + contents: read + pull-requests: write + issues: write + +env: + BUILD_TYPE: Release + BENCH_WARMUP: ${{ github.event.inputs.warmup || '5' }} + BENCH_ITERS: ${{ github.event.inputs.iterations || '20' }} + BENCH_SIZES: ${{ github.event.inputs.sizes || '256,512,1024,2048,4096' }} + DEPS_CACHE: /tmp/sofieblas-cmake-deps + +jobs: + parse-command: + name: Parse trigger + runs-on: ubuntu-latest + if: | + github.event_name == 'workflow_dispatch' || + ( + github.event_name == 'issue_comment' && + github.event.issue.pull_request != null && + contains(fromJSON('["sanjibansg","lmoneta"]'), github.event.comment.user.login) && + startsWith(github.event.comment.body, '/runbenchmark') + ) + outputs: + runner_label: ${{ steps.parse.outputs.runner_label }} + valid: ${{ steps.parse.outputs.valid }} + flavor: ${{ steps.parse.outputs.flavor }} + valid_flavors: ${{ steps.parse.outputs.valid_flavors }} + backend: ${{ steps.parse.outputs.backend }} + steps: + - name: React to trigger comment + if: github.event_name == 'issue_comment' + uses: actions/github-script@v7 + with: + script: | + await github.rest.reactions.createForIssueComment({ + owner: context.repo.owner, + repo: context.repo.repo, + comment_id: context.payload.comment.id, + content: 'eyes' + }); + - name: Parse runner flavor + id: parse + uses: actions/github-script@v7 + with: + script: | + // default flavor keyed by name -> { self-hosted runner label, GPU backend } + const flavors = { + l40s: { runner: 'ml4ep-l40s', backend: 'cuda' }, + h100: { runner: 'ml4ep-h100', backend: 'cuda' }, + 'h100-47gb': { runner: 'ml4ep-h100-47gb', backend: 'cuda' }, + w7900: { runner: 'ml4ep-w7900', backend: 'hip' }, + mi300x: { runner: 'ml4ep-mi300x', backend: 'hip' }, + }; + let flavor = 'l40s'; + let valid = true; + if (context.eventName === 'issue_comment') { + const m = context.payload.comment.body.trim().match(/^\/runbenchmark(?:\s+(\S+))?/i); + if (m && m[1]) { + flavor = m[1].toLowerCase(); + valid = Object.prototype.hasOwnProperty.call(flavors, flavor); + } + } + core.setOutput('flavor', flavor); + core.setOutput('valid', String(valid)); + core.setOutput('runner_label', valid ? flavors[flavor].runner : ''); + core.setOutput('backend', valid ? flavors[flavor].backend : ''); + core.setOutput('valid_flavors', Object.keys(flavors).join(', ')); + - name: Post run link + if: github.event_name == 'issue_comment' + uses: actions/github-script@v7 + with: + script: | + const runUrl = `${context.serverUrl}/${context.repo.owner}/${context.repo.repo}/actions/runs/${context.runId}`; + await github.rest.issues.createComment({ + owner: context.repo.owner, + repo: context.repo.repo, + issue_number: context.payload.issue.number, + body: `\`/runbenchmark\` (${{ steps.parse.outputs.flavor }}): triggered - [view run](${runUrl})` + }); + + benchmark: + name: Benchmark Comparison + needs: parse-command + if: needs.parse-command.outputs.valid == 'true' + concurrency: + group: benchmark-${{ needs.parse-command.outputs.flavor }}-${{ github.event.issue.number || github.ref }} + cancel-in-progress: true + runs-on: ${{ needs.parse-command.outputs.runner_label }} + container: registry.cern.ch/ngt/lxplus-like:9 + timeout-minutes: 60 + + steps: + - name: GPU check (NVIDIA) + if: needs.parse-command.outputs.backend == 'cuda' + run: nvidia-smi + + - name: GPU check (AMD) + if: needs.parse-command.outputs.backend == 'hip' + run: rocm-smi + + - name: Resolve PR info + id: pr + if: github.event_name == 'issue_comment' + uses: actions/github-script@v7 + with: + script: | + const pr = await github.rest.pulls.get({ + owner: context.repo.owner, + repo: context.repo.repo, + pull_number: context.payload.issue.number + }); + core.setOutput('number', pr.data.number); + core.setOutput('merge_ref', `refs/pull/${pr.data.number}/merge`); + core.setOutput('base_ref', pr.data.base.ref); + + - name: Checkout PR branch + uses: actions/checkout@v4 + with: + path: sofieblas-pr + ref: ${{ github.event_name == 'issue_comment' && steps.pr.outputs.merge_ref || github.ref }} + + - name: Checkout base branch + if: github.event_name == 'issue_comment' + uses: actions/checkout@v4 + with: + ref: ${{ steps.pr.outputs.base_ref }} + path: sofieblas-main + + - name: Cache FetchContent dependencies + uses: actions/cache@v4 + with: + path: ${{ env.DEPS_CACHE }} + key: cmake-deps-bench-${{ needs.parse-command.outputs.backend }}-${{ hashFiles('sofieblas-pr/cmake/SofieBLASBackends.cmake') }} + restore-keys: cmake-deps-bench-${{ needs.parse-command.outputs.backend }}- + + - name: Configure PR build + env: + BACKEND: ${{ needs.parse-command.outputs.backend }} + run: | + cmake -B sofieblas-pr/build -S sofieblas-pr \ + -DCMAKE_BUILD_TYPE=${{ env.BUILD_TYPE }} \ + -DSOFIEBLAS_BUILD_BENCHMARKS=ON \ + -DCMAKE_CUDA_ARCHITECTURES=native \ + -DCMAKE_HIP_ARCHITECTURES=native \ + "-DSOFIEBLAS_ENABLE_CUDA=$([ "$BACKEND" = "cuda" ] && echo ON || echo OFF)" \ + "-DSOFIEBLAS_ENABLE_HIP=$([ "$BACKEND" = "hip" ] && echo ON || echo OFF)" \ + "-DFETCHCONTENT_BASE_DIR=${{ env.DEPS_CACHE }}" + + - name: Build PR benchmark + run: cmake --build sofieblas-pr/build --target "bench_${{ needs.parse-command.outputs.backend }}" -j"$(nproc)" + + - name: Run PR benchmark + working-directory: sofieblas-pr/build/benchmark + run: | + "./bench_${{ needs.parse-command.outputs.backend }}" \ + -w "${{ env.BENCH_WARMUP }}" \ + -n "${{ env.BENCH_ITERS }}" \ + --sizes "${{ env.BENCH_SIZES }}" \ + | tee benchmark_pr.txt + + - name: Check whether base branch supports the benchmark harness + id: base_support + if: github.event_name == 'issue_comment' + run: | + if grep -q "SOFIEBLAS_BUILD_BENCHMARKS" sofieblas-main/CMakeLists.txt 2>/dev/null; then + echo "supported=true" >> "$GITHUB_OUTPUT" + else + echo "supported=false" >> "$GITHUB_OUTPUT" + echo "::warning::${{ steps.pr.outputs.base_ref }} does not define the SOFIEBLAS_BUILD_BENCHMARKS option yet - skipping the base-branch comparison." + fi + + - name: Configure base build + if: github.event_name == 'issue_comment' && steps.base_support.outputs.supported == 'true' + env: + BACKEND: ${{ needs.parse-command.outputs.backend }} + run: | + cmake -B sofieblas-main/build -S sofieblas-main \ + -DCMAKE_BUILD_TYPE=${{ env.BUILD_TYPE }} \ + -DSOFIEBLAS_BUILD_BENCHMARKS=ON \ + -DCMAKE_CUDA_ARCHITECTURES=native \ + -DCMAKE_HIP_ARCHITECTURES=native \ + "-DSOFIEBLAS_ENABLE_CUDA=$([ "$BACKEND" = "cuda" ] && echo ON || echo OFF)" \ + "-DSOFIEBLAS_ENABLE_HIP=$([ "$BACKEND" = "hip" ] && echo ON || echo OFF)" \ + "-DFETCHCONTENT_BASE_DIR=${{ env.DEPS_CACHE }}" + + - name: Build base benchmark + if: github.event_name == 'issue_comment' && steps.base_support.outputs.supported == 'true' + run: cmake --build sofieblas-main/build --target "bench_${{ needs.parse-command.outputs.backend }}" -j"$(nproc)" + + - name: Run base benchmark + if: github.event_name == 'issue_comment' && steps.base_support.outputs.supported == 'true' + working-directory: sofieblas-main/build/benchmark + run: | + "./bench_${{ needs.parse-command.outputs.backend }}" \ + -w "${{ env.BENCH_WARMUP }}" \ + -n "${{ env.BENCH_ITERS }}" \ + --sizes "${{ env.BENCH_SIZES }}" \ + | tee benchmark_main.txt + + - name: Build comparison summary + id: summary + if: always() + run: | + { + echo "summary</dev/null || echo "(no results)" + if [ -f sofieblas-main/build/benchmark/benchmark_main.txt ]; then + echo "" + echo "── ${{ steps.pr.outputs.base_ref }} ──────────────────────────────────────────────" + cat sofieblas-main/build/benchmark/benchmark_main.txt + fi + echo '```' + echo "BENCH_EOF" + } >> "$GITHUB_OUTPUT" + + - name: Post results as PR comment + if: always() && github.event_name == 'issue_comment' + uses: marocchino/sticky-pull-request-comment@v2 + with: + number: ${{ steps.pr.outputs.number }} + header: benchmark + message: ${{ steps.summary.outputs.summary }} + + - name: Write job summary + if: always() + run: echo "${{ steps.summary.outputs.summary }}" >> "$GITHUB_STEP_SUMMARY" + + - name: Notify benchmark completion + if: always() && github.event_name == 'issue_comment' + uses: actions/github-script@v7 + with: + script: | + const status = '${{ job.status }}' === 'success' ? '✅ completed' : '❌ failed'; + const runUrl = `${context.serverUrl}/${context.repo.owner}/${context.repo.repo}/actions/runs/${context.runId}`; + await github.rest.issues.createComment({ + owner: context.repo.owner, + repo: context.repo.repo, + issue_number: context.payload.issue.number, + body: `\`/runbenchmark\` (${{ needs.parse-command.outputs.flavor }}): Benchmark ${status} - [view results](${runUrl})` + }); + + - name: Upload benchmark results + if: always() + uses: actions/upload-artifact@v4 + with: + name: benchmark-results-${{ github.run_id }} + path: | + sofieblas-pr/build/benchmark/benchmark_pr.txt + sofieblas-main/build/benchmark/benchmark_main.txt + if-no-files-found: ignore + + report-invalid-flavor: + name: Report invalid runner flavor + needs: parse-command + runs-on: ubuntu-latest + if: >- + always() && github.event_name == 'issue_comment' && + needs.parse-command.result == 'success' && + needs.parse-command.outputs.valid == 'false' + steps: + - uses: actions/github-script@v7 + with: + script: | + await github.rest.issues.createComment({ + owner: context.repo.owner, + repo: context.repo.repo, + issue_number: context.payload.issue.number, + body: `\`/runbenchmark\`: unknown runner flavor \`${{ needs.parse-command.outputs.flavor }}\`. ` + + `Valid options: ${{ needs.parse-command.outputs.valid_flavors }} (or omit for the default, l40s).` + }); diff --git a/.github/workflows/tests.yml b/.github/workflows/tests.yml new file mode 100644 index 0000000..e41ae81 --- /dev/null +++ b/.github/workflows/tests.yml @@ -0,0 +1,240 @@ +name: Unit Tests + +on: + push: + branches: [main, dev] + pull_request: + issue_comment: + types: [created] + workflow_dispatch: + +permissions: + contents: read + pull-requests: write + issues: write + +env: + BUILD_TYPE: Release + DEPS_CACHE: /tmp/sofieblas-cmake-deps + +jobs: + cpu-tests: + name: CPU Unit Tests + if: | + github.event_name != 'issue_comment' || + ( + github.event.issue.pull_request != null && + !startsWith(github.event.comment.body, '/runtest') && + !startsWith(github.event.comment.body, '/runbenchmark') + ) + concurrency: + group: cpu-tests-${{ github.event.issue.number || github.ref }} + cancel-in-progress: true + runs-on: ubuntu-latest + timeout-minutes: 30 + steps: + - name: Checkout + uses: actions/checkout@v4 + with: + ref: ${{ github.event_name == 'issue_comment' && format('refs/pull/{0}/merge', github.event.issue.number) || github.ref }} + + - name: Install OpenBLAS + run: | + sudo apt-get update + sudo apt-get install -y libopenblas-dev + + - name: Cache FetchContent dependencies + uses: actions/cache@v4 + with: + path: ${{ env.DEPS_CACHE }} + key: cmake-deps-cpu-${{ hashFiles('cmake/SofieBLASBackends.cmake') }} + restore-keys: cmake-deps-cpu- + + - name: Configure + run: | + cmake -B build -S . \ + -DCMAKE_BUILD_TYPE=${{ env.BUILD_TYPE }} \ + -DSOFIEBLAS_BUILD_TESTS=ON \ + -DSOFIEBLAS_BUILD_BENCHMARKS=ON \ + -DSOFIEBLAS_ENABLE_CUDA=OFF \ + -DSOFIEBLAS_ENABLE_HIP=OFF \ + "-DFETCHCONTENT_BASE_DIR=${{ env.DEPS_CACHE }}" + + - name: Build + run: cmake --build build -j"$(nproc)" + + - name: Run tests + working-directory: build + run: ctest --output-on-failure -j"$(nproc)" -R '\.cpu(\.|$)' + + - name: Upload test log + if: always() + uses: actions/upload-artifact@v4 + with: + name: cpu-test-log-${{ github.run_id }} + path: build/Testing/Temporary/LastTest.log + if-no-files-found: ignore + + parse-command: + name: Parse trigger + runs-on: ubuntu-latest + if: | + github.event_name == 'workflow_dispatch' || + ( + github.event_name == 'issue_comment' && + github.event.issue.pull_request != null && + contains(fromJSON('["sanjibansg","lmoneta"]'), github.event.comment.user.login) && + startsWith(github.event.comment.body, '/runtest') + ) + outputs: + runner_label: ${{ steps.parse.outputs.runner_label }} + valid: ${{ steps.parse.outputs.valid }} + flavor: ${{ steps.parse.outputs.flavor }} + valid_flavors: ${{ steps.parse.outputs.valid_flavors }} + backend: ${{ steps.parse.outputs.backend }} + steps: + - name: React to trigger comment + if: github.event_name == 'issue_comment' + uses: actions/github-script@v7 + with: + script: | + await github.rest.reactions.createForIssueComment({ + owner: context.repo.owner, + repo: context.repo.repo, + comment_id: context.payload.comment.id, + content: 'eyes' + }); + - name: Parse runner flavor + id: parse + uses: actions/github-script@v7 + with: + script: | + const flavors = { + l40s: { runner: 'ml4ep-l40s', backend: 'cuda' }, + h100: { runner: 'ml4ep-h100', backend: 'cuda' }, + 'h100-47gb': { runner: 'ml4ep-h100-47gb', backend: 'cuda' }, + w7900: { runner: 'ml4ep-w7900', backend: 'hip' }, + mi300x: { runner: 'ml4ep-mi300x', backend: 'hip' }, + }; + let flavor = 'l40s'; + let valid = true; + if (context.eventName === 'issue_comment') { + const m = context.payload.comment.body.trim().match(/^\/runtest(?:\s+(\S+))?/i); + if (m && m[1]) { + flavor = m[1].toLowerCase(); + valid = Object.prototype.hasOwnProperty.call(flavors, flavor); + } + } + core.setOutput('flavor', flavor); + core.setOutput('valid', String(valid)); + core.setOutput('runner_label', valid ? flavors[flavor].runner : ''); + core.setOutput('backend', valid ? flavors[flavor].backend : ''); + core.setOutput('valid_flavors', Object.keys(flavors).join(', ')); + - name: Post run link + if: github.event_name == 'issue_comment' + uses: actions/github-script@v7 + with: + script: | + const runUrl = `${context.serverUrl}/${context.repo.owner}/${context.repo.repo}/actions/runs/${context.runId}`; + await github.rest.issues.createComment({ + owner: context.repo.owner, + repo: context.repo.repo, + issue_number: context.payload.issue.number, + body: `\`/runtest\` (${{ steps.parse.outputs.flavor }}): triggered - [view run](${runUrl})` + }); + + gpu-tests: + name: GPU Unit Tests + needs: parse-command + if: needs.parse-command.outputs.valid == 'true' + concurrency: + group: gpu-tests-${{ needs.parse-command.outputs.flavor }}-${{ github.event.issue.number || github.ref }} + cancel-in-progress: true + runs-on: ${{ needs.parse-command.outputs.runner_label }} + container: + image: registry.cern.ch/ngt/lxplus-like:9 + options: --user 0:0 + timeout-minutes: 45 + steps: + - name: GPU check (NVIDIA) + if: needs.parse-command.outputs.backend == 'cuda' + run: nvidia-smi + + - name: GPU check (AMD) + if: needs.parse-command.outputs.backend == 'hip' + run: rocm-smi + + - name: Checkout + uses: actions/checkout@v4 + with: + ref: ${{ github.event_name == 'issue_comment' && format('refs/pull/{0}/merge', github.event.issue.number) || github.ref }} + + - name: Cache FetchContent dependencies + uses: actions/cache@v4 + with: + path: ${{ env.DEPS_CACHE }} + key: cmake-deps-gpu-${{ needs.parse-command.outputs.backend }}-${{ hashFiles('cmake/SofieBLASBackends.cmake') }} + restore-keys: cmake-deps-gpu-${{ needs.parse-command.outputs.backend }}- + + - name: Configure + env: + BACKEND: ${{ needs.parse-command.outputs.backend }} + run: | + cmake -B build -S . \ + -DCMAKE_BUILD_TYPE=${{ env.BUILD_TYPE }} \ + -DSOFIEBLAS_BUILD_TESTS=ON \ + -DSOFIEBLAS_BUILD_BENCHMARKS=ON \ + -DCMAKE_CUDA_ARCHITECTURES=native \ + -DCMAKE_HIP_ARCHITECTURES=native \ + "-DSOFIEBLAS_ENABLE_CUDA=$([ "$BACKEND" = "cuda" ] && echo ON || echo OFF)" \ + "-DSOFIEBLAS_ENABLE_HIP=$([ "$BACKEND" = "hip" ] && echo ON || echo OFF)" \ + "-DFETCHCONTENT_BASE_DIR=${{ env.DEPS_CACHE }}" + + - name: Build + run: cmake --build build -j"$(nproc)" + + - name: Run GPU tests + working-directory: build + run: ctest --output-on-failure -j"$(nproc)" -R "\.${{ needs.parse-command.outputs.backend }}(\.|\$)" + + - name: Upload test log + if: always() + uses: actions/upload-artifact@v4 + with: + name: gpu-test-log-${{ github.run_id }} + path: build/Testing/Temporary/LastTest.log + if-no-files-found: ignore + + - name: Report result on PR + if: always() && github.event_name == 'issue_comment' + uses: actions/github-script@v7 + with: + script: | + const status = '${{ job.status }}' === 'success' ? '✅ passed' : '❌ failed'; + const runUrl = `${context.serverUrl}/${context.repo.owner}/${context.repo.repo}/actions/runs/${context.runId}`; + await github.rest.issues.createComment({ + owner: context.repo.owner, + repo: context.repo.repo, + issue_number: context.payload.issue.number, + body: `\`/runtest\` (${{ needs.parse-command.outputs.flavor }}): GPU Unit Tests ${status} - [view run](${runUrl})` + }); + + report-invalid-flavor: + name: Report invalid runner flavor + needs: parse-command + runs-on: ubuntu-latest + if: >- + always() && github.event_name == 'issue_comment' && + needs.parse-command.result == 'success' && + needs.parse-command.outputs.valid == 'false' + steps: + - uses: actions/github-script@v7 + with: + script: | + await github.rest.issues.createComment({ + owner: context.repo.owner, + repo: context.repo.repo, + issue_number: context.payload.issue.number, + body: `\`/runtest\`: unknown runner flavor \`${{ needs.parse-command.outputs.flavor }}\`. ` + + `Valid options: ${{ needs.parse-command.outputs.valid_flavors }} (or omit for the default, l40s).` + }); diff --git a/.gitignore b/.gitignore index 475c359..2f9ffe1 100644 --- a/.gitignore +++ b/.gitignore @@ -41,4 +41,7 @@ *.dwo # build files -**/build/ \ No newline at end of file +**/build/ + +# vscode settings +.vscode/ \ No newline at end of file diff --git a/CMakeLists.txt b/CMakeLists.txt new file mode 100644 index 0000000..368a522 --- /dev/null +++ b/CMakeLists.txt @@ -0,0 +1,69 @@ +cmake_minimum_required(VERSION 3.21) +project(sofieBLAS + VERSION 0.1.0 + DESCRIPTION "Header-only C++ BLAS abstraction over heterogeneous backends" + LANGUAGES CXX) + +set(CMAKE_CXX_STANDARD 20) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CXX_EXTENSIONS OFF) + +include(GNUInstallDirs) +include(CMakePackageConfigHelpers) + +option(SOFIEBLAS_BUILD_TESTS "Build sofieBLAS unit tests" OFF) +option(SOFIEBLAS_BUILD_BENCHMARKS "Build sofieBLAS GEMM benchmarks" OFF) + +# --- Header-only interface library --- +add_library(sofieBLAS INTERFACE) +add_library(sofieBLAS::sofieBLAS ALIAS sofieBLAS) + +target_include_directories(sofieBLAS INTERFACE + $ + $) + +target_compile_features(sofieBLAS INTERFACE cxx_std_20) + +if(SOFIEBLAS_BUILD_TESTS OR SOFIEBLAS_BUILD_BENCHMARKS) + list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake") + include(SofieBLASBackends) +endif() + +if(SOFIEBLAS_BUILD_TESTS) + enable_testing() + add_subdirectory(tests) +endif() + +if(SOFIEBLAS_BUILD_BENCHMARKS) + add_subdirectory(benchmark) +endif() + +# --- Install (header-only package) --- +install(TARGETS sofieBLAS EXPORT sofieBLASTargets) +install(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include/ DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}) + +configure_package_config_file( + ${CMAKE_CURRENT_SOURCE_DIR}/cmake/sofieBLASConfig.cmake.in + ${CMAKE_CURRENT_BINARY_DIR}/sofieBLASConfig.cmake + INSTALL_DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/sofieBLAS +) + +write_basic_package_version_file( + ${CMAKE_CURRENT_BINARY_DIR}/sofieBLASConfigVersion.cmake + VERSION ${PROJECT_VERSION} + COMPATIBILITY AnyNewerVersion + ARCH_INDEPENDENT +) + +install( + EXPORT sofieBLASTargets + FILE sofieBLASTargets.cmake + NAMESPACE sofieBLAS:: + DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/sofieBLAS +) + +install(FILES + ${CMAKE_CURRENT_BINARY_DIR}/sofieBLASConfig.cmake + ${CMAKE_CURRENT_BINARY_DIR}/sofieBLASConfigVersion.cmake + DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/sofieBLAS +) diff --git a/README.md b/README.md index 93bfdf6..191fe2a 100644 --- a/README.md +++ b/README.md @@ -1,15 +1,54 @@ # sofieBLAS -sofieBLAS is an abstract C++ (header-only) interface for BLAS operations targeting heterogeneous architectures. It currently supports only ALPAKA buffers and the GEMM operation, acting as a thin, efficient wrapper over existing BLAS libraries such as OpenBLAS, MKL, cuBLAS, and others- allowing the actual backend to be selected through template-based dispatching using traits. +sofieBLAS is an abstract C++ (header-only) interface for BLAS operations targeting heterogeneous architectures. It currently supports only ALPAKA buffers and the GEMM operation, acting as a thin, efficient wrapper over existing BLAS libraries such as OpenBLAS, MKL, BLIS, Apple Accelerate, cuBLASLt, and hipBLASLt - allowing the actual backend to be selected through template-based dispatching using traits. We plan to extend support to more BLAS routines and buffer types in future releases. ## Features - Unified Interface: Common C++ API over multiple BLAS backends. -- Heterogeneous Support: CPU (OpenBLAS, MKL) and GPU (cuBLAS) support. -- Template-Based Dispatching: Backend selection via traits at compile-time. +- Heterogeneous Support: + - CPU: OpenBLAS, MKL, BLIS, Apple Accelerate (any CBLAS-compatible library). + - GPU: NVIDIA (cuBLASLt) and AMD (hipBLASLt). +- Template-Based Dispatching: Backend selection via traits at compile-time, keyed on the Alpaka accelerator tag (e.g. `alpaka::TagCpuSerial`, `alpaka::TagGpuCudaRt`, `alpaka::TagGpuHipRt`). - Header-Only: Lightweight, easy to integrate- no separate compilation required. - Minimal Dependency Overhead: Only depends on the backend BLAS libraries of choice. +- One File Per Backend: Each vendor library (CPU or GPU) lives in its own header under `include/sofieBLAS/backends//`, selected at compile time via a preprocessor macro (`ALPAKA_ACC_GPU_CUDA_ENABLED`, `ALPAKA_ACC_GPU_HIP_ENABLED`, `SOFIEBLAS_USE_OPENBLAS`, `SOFIEBLAS_USE_MKL`, `SOFIEBLAS_USE_BLIS`, `SOFIEBLAS_USE_ACCELERATE`) so only the code for the library you actually build against gets compiled. + +## Selecting a backend + +Backend selection happens entirely at compile time, in two steps: + +1. **Compiler-flag macros decide the backend implementations.** A macro (or macro pair) must be defined to compile in a given backend's header - see the table below. For CPU, if `ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED` is defined but none of the `SOFIEBLAS_USE_*` macros are, sofieBLAS defaults to OpenBLAS; the GPU backends have no such default (each is tied 1:1 to its `ALPAKA_ACC_GPU_*_ENABLED` macro, so there's nothing to default between). +2. **The Alpaka tag you instantiate `sofieBLAS` with decides which compiled-in backend a given call site actually uses**, via the `traits::sofieBLAS` specialization (e.g. `sofieBLAS` resolves to the hipBLASLt backend). There is no runtime dispatch - if the macro for that tag's backend wasn't defined, the code simply won't compile. + +| Accelerator | Macro(s) | Alpaka tag | +| --- | --- | --- | +| CPU (OpenBLAS) | `ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED` + `SOFIEBLAS_USE_OPENBLAS` | `alpaka::TagCpuSerial` (or other `TagCpu*`) | +| CPU (Intel MKL) | `ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED` + `SOFIEBLAS_USE_MKL` | `alpaka::TagCpuSerial` (or other `TagCpu*`) | +| CPU (BLIS) | `ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED` + `SOFIEBLAS_USE_BLIS` | `alpaka::TagCpuSerial` (or other `TagCpu*`) | +| CPU (Apple Accelerate) | `ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED` + `SOFIEBLAS_USE_ACCELERATE` | `alpaka::TagCpuSerial` (or other `TagCpu*`) | +| NVIDIA GPU (cuBLASLt) | `ALPAKA_ACC_GPU_CUDA_ENABLED` | `alpaka::TagGpuCudaRt` | +| AMD GPU (hipBLASLt) | `ALPAKA_ACC_GPU_HIP_ENABLED` | `alpaka::TagGpuHipRt` | + +`sofieBLAS/sofieBLAS.hpp` includes the matching backend header(s) for you based on these macros; nothing else needs to change in your source beyond picking the right tag. + +## Building tests and benchmarks + +sofieBLAS itself is header-only (`add_subdirectory`/`find_package(sofieBLAS)` gives you an `INTERFACE` target with no build step). Tests and benchmarks are opt-in via CMake options and auto-detect whichever backends are available on the machine for each backend target (`test_cpu`/`test_cuda`/`test_hip`, `bench_cpu`/`bench_cuda`/`bench_hip`) and is skipped if its dependency isn't found: + +```bash +cmake -B build -S . -DSOFIEBLAS_BUILD_TESTS=ON -DSOFIEBLAS_BUILD_BENCHMARKS=ON +cmake --build build -j"$(nproc)" + +# Run the correctness tests (matmul/gemm/gemmrelu/gemmgelu vs. a reference +# implementation, per available backend) +ctest --test-dir build --output-on-failure + +# Run the GEMM throughput benchmark +./build/benchmark/bench_cuda -w 5 -n 20 --sizes 256,512,1024,2048,4096 +``` + +`CPU_BLAS_LIB` (`OpenBLAS`/`MKL`/`BLIS`/`Accelerate`) and `CUDA_BASE`/`ROCM_BASE`/`ONEAPI_BASE`/`BLIS_BASE` are available as `-D` cache variables to point at non-default install locations; alpaka is picked up via `find_package(alpaka)` if already installed, otherwise fetched automatically. See `tests/CMakeLists.txt` and `benchmark/CMakeLists.txt` for the target definitions, and `cmake/SofieBLASBackends.cmake` for the shared backend-detection logic. ## Usage example @@ -32,11 +71,12 @@ int main() { auto C = alpaka::allocBuf(device, size * size); // (Initialize A and B here...) - // Create sofieBLAS instance for CPU backend + // Create sofieBLAS instance for the CPU backend selected via + // SOFIEBLAS_USE_OPENBLAS / SOFIEBLAS_USE_MKL / SOFIEBLAS_USE_BLIS / SOFIEBLAS_USE_ACCELERATE sofieBLAS blas(queue); - // Perform GEMM: C = alpha * A * B + beta * C - blas.gemm('n', 'n', size, size, size, 1.0f, A, size, B, size, 0.0f, C, size); + // C = alpha * op(A) * op(B) + beta * C (leading dimensions inferred from m, n, k) + blas.matmul('N', 'N', size, size, size, 1.0f, A, B, 0.0f, C); alpaka::wait(queue); std::cout << "GEMM completed on CPU backend.\n"; @@ -45,6 +85,19 @@ int main() { } ``` +Switching to a GPU is the same code shape, just a different Alpaka tag, queue/device type, and build-time macro. For example, on AMD (built with `-DALPAKA_ACC_GPU_HIP_ENABLED`): + +```cpp +alpaka::PlatformHipRt platform; +auto device = alpaka::getDevByIdx(platform, 0u); +alpaka::Queue queue{device}; +sofieBLAS blas(queue); + +blas.matmul('N', 'N', size, size, size, 1.0f, dA, dB, 0.0f, dC); +``` + +The GPU backends (`BlasCuda`, `BlasHip`) additionally expose `gemmrelu`/`gemmgelu` (fused bias + activation via cuBLASLt/hipBLASLt epilogues), `gemmStridedBatched`, and `addLayoutConfig` (used to pre-register cuBLASLt/hipBLASLt matrix layouts for a given shape before the first `matmul`/`gemm` call on that shape). + ## Contributing diff --git a/benchmark/CMakeLists.txt b/benchmark/CMakeLists.txt new file mode 100644 index 0000000..a30efd1 --- /dev/null +++ b/benchmark/CMakeLists.txt @@ -0,0 +1,71 @@ +set(CXXFLAGS -O3 -DALPAKA_HAS_STD_ATOMIC_REF) +set(CXX_HOST_FLAGS -fPIC -pthread) +set(CXX_CUDA_FLAGS -Wno-deprecated-gpu-targets --extended-lambda --expt-relaxed-constexpr) +set(XCOMPILER_FLAGS -Xcompiler=-fPIC,-pthread) + +if(SOFIEBLAS_CPU_BLAS_LIBS) + add_executable(bench_cpu bench.cc) + target_compile_features(bench_cpu PUBLIC cxx_std_20) + target_compile_options(bench_cpu PRIVATE ${CXXFLAGS} ${CXX_HOST_FLAGS}) + target_compile_definitions(bench_cpu PRIVATE + ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED ${SOFIEBLAS_CPU_BLAS_DEFINE}) + target_include_directories(bench_cpu PRIVATE ${SOFIEBLAS_CPU_BLAS_INCLUDE_DIR}) + target_link_libraries(bench_cpu PRIVATE + sofieBLAS::sofieBLAS alpaka::alpaka ${SOFIEBLAS_CPU_BLAS_LIBS}) + + if(SOFIEBLAS_BUILD_TESTS) + add_test(NAME sofieBLAS.bench.cpu.smoke + COMMAND bench_cpu -w 1 -n 1 --sizes 64,128) + endif() +else() + message(STATUS "sofieBLAS benchmark: skipping bench_cpu (no CPU BLAS library found)") +endif() + +if(SOFIEBLAS_CUDA_ENABLED) + configure_file(bench.cc bench_cuda_gen.cc COPYONLY) + add_executable(bench_cuda) + target_sources(bench_cuda PRIVATE ${CMAKE_CURRENT_BINARY_DIR}/bench_cuda_gen.cc) + set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/bench_cuda_gen.cc PROPERTIES LANGUAGE CUDA) + set_target_properties(bench_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + target_compile_features(bench_cuda PUBLIC cxx_std_20) + + target_compile_options(bench_cuda PRIVATE + ${CXXFLAGS} + ${CXX_CUDA_FLAGS} + ${XCOMPILER_FLAGS} + ) + + target_compile_definitions(bench_cuda PRIVATE ALPAKA_ACC_GPU_CUDA_ENABLED) + target_link_libraries(bench_cuda PRIVATE + sofieBLAS::sofieBLAS alpaka::alpaka CUDA::cudart CUDA::cublas CUDA::cublasLt) + + if(SOFIEBLAS_BUILD_TESTS) + add_test(NAME sofieBLAS.bench.cuda.smoke + COMMAND bench_cuda -w 1 -n 1 --sizes 64,128) + endif() +else() + message(VERBOSE "sofieBLAS benchmark: skipping bench_cuda (CUDA backend not enabled)") +endif() + +if(SOFIEBLAS_HIP_ENABLED) + configure_file(bench.cc bench_hip_gen.cc COPYONLY) + add_executable(bench_hip) + target_sources(bench_hip PRIVATE ${CMAKE_CURRENT_BINARY_DIR}/bench_hip_gen.cc) + set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/bench_hip_gen.cc PROPERTIES LANGUAGE HIP) + target_compile_features(bench_hip PUBLIC cxx_std_20) + + target_compile_options(bench_hip PRIVATE ${CXXFLAGS} ${CXX_HOST_FLAGS}) + + target_compile_definitions(bench_hip PRIVATE ALPAKA_ACC_GPU_HIP_ENABLED) + target_include_directories(bench_hip PRIVATE ${ROCM_BASE}/include) + target_link_directories(bench_hip PRIVATE ${ROCM_BASE}/lib) + target_link_libraries(bench_hip PRIVATE + sofieBLAS::sofieBLAS alpaka::alpaka hipblaslt hipblas amdhip64) + + if(SOFIEBLAS_BUILD_TESTS) + add_test(NAME sofieBLAS.bench.hip.smoke + COMMAND bench_hip -w 1 -n 1 --sizes 64,128) + endif() +else() + message(VERBOSE "sofieBLAS benchmark: skipping bench_hip (HIP backend not enabled)") +endif() diff --git a/benchmark/bench.cc b/benchmark/bench.cc new file mode 100644 index 0000000..3263dab --- /dev/null +++ b/benchmark/bench.cc @@ -0,0 +1,228 @@ +#include "sofieBLAS/sofieBLAS.hpp" +#include + +#include +#include +#include +#include +#include +#include +#include + +using Idx = uint32_t; + +struct BenchOptions { + int warmup = 5; + int iterations = 20; + std::vector sizes = {256, 512, 1024, 2048, 4096}; +}; + +static BenchOptions parseArgs(int argc, char **argv) { + BenchOptions opt; + for (int i = 1; i < argc; ++i) { + std::string arg = argv[i]; + auto nextVal = [&](void) -> std::string { + if (i + 1 >= argc) { + std::cerr << "Missing value for " << arg << "\n"; + std::exit(EXIT_FAILURE); + } + return argv[++i]; + }; + if (arg == "-w" || arg == "--warmup") { + opt.warmup = std::stoi(nextVal()); + } else if (arg == "-n" || arg == "--iterations") { + opt.iterations = std::stoi(nextVal()); + } else if (arg == "--sizes") { + opt.sizes.clear(); + std::stringstream ss(nextVal()); + std::string tok; + while (std::getline(ss, tok, ',')) + opt.sizes.push_back(std::stoi(tok)); + } else if (arg == "-h" || arg == "--help") { + std::cout << "Usage: " << argv[0] + << " [-w warmup] [-n iterations] [--sizes s1,s2,...]\n"; + std::exit(EXIT_SUCCESS); + } else { + std::cerr << "Unknown argument: " << arg << "\n"; + std::exit(EXIT_FAILURE); + } + } + return opt; +} + +static void fillSeq(float *M, int n, float start = 0.01f, float step = 0.001f) { + for (int i = 0; i < n; ++i) + M[i] = start + static_cast(i % 997) * step; +} + +static void printHeader() { + std::cout << std::left << std::setw(10) << "Backend" << std::right + << std::setw(10) << "M=N=K" << std::setw(14) << "avg (ms)" + << std::setw(14) << "GFLOP/s" << "\n"; + std::cout << std::string(48, '-') << "\n"; +} + +static void printRow(const std::string &backend, int size, double avgMs, + double gflops) { + std::cout << std::left << std::setw(10) << backend << std::right + << std::setw(10) << size << std::setw(14) << std::fixed + << std::setprecision(3) << avgMs << std::setw(14) + << std::setprecision(2) << gflops << "\n"; +} + +static double gflopsFor(int size, double avgSeconds) { + const double flops = 2.0 * static_cast(size) * + static_cast(size) * static_cast(size); + return flops / avgSeconds / 1e9; +} + +// --------------------------------------------------------------------------- +// CPU benchmark +// --------------------------------------------------------------------------- +#ifdef ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED + +static void runCpuBench(const BenchOptions &opt) { + alpaka::PlatformCpu platform{}; + auto dev = alpaka::getDevByIdx(platform, 0u); + alpaka::Queue queue{dev}; + sofieBLAS blas(queue); + + for (int size : opt.sizes) { + const int M = size, N = size, K = size; + auto hA = alpaka::allocBuf(dev, static_cast(M * K)); + auto hB = alpaka::allocBuf(dev, static_cast(K * N)); + auto hC = alpaka::allocBuf(dev, static_cast(M * N)); + fillSeq(alpaka::getPtrNative(hA), M * K); + fillSeq(alpaka::getPtrNative(hB), K * N, 0.02f, 0.0005f); + + for (int i = 0; i < opt.warmup; ++i) + blas.matmul('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hC); + + auto t0 = std::chrono::steady_clock::now(); + for (int i = 0; i < opt.iterations; ++i) + blas.matmul('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hC); + auto t1 = std::chrono::steady_clock::now(); + + double avgSeconds = + std::chrono::duration(t1 - t0).count() / opt.iterations; + printRow("cpu", size, avgSeconds * 1e3, gflopsFor(size, avgSeconds)); + } +} + +#endif // ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED + +// --------------------------------------------------------------------------- +// CUDA benchmark +// --------------------------------------------------------------------------- +#ifdef ALPAKA_ACC_GPU_CUDA_ENABLED + +static void runCudaBench(const BenchOptions &opt) { + alpaka::PlatformCudaRt platform{}; + auto dev = alpaka::getDevByIdx(platform, 0u); + alpaka::Queue queue{dev}; + sofieBLAS blas(queue); + + alpaka::PlatformCpu hostPlatform{}; + auto hostDev = alpaka::getDevByIdx(hostPlatform, 0u); + + for (int size : opt.sizes) { + const int M = size, N = size, K = size; + auto hA = alpaka::allocBuf(hostDev, static_cast(M * K)); + auto hB = alpaka::allocBuf(hostDev, static_cast(K * N)); + fillSeq(alpaka::getPtrNative(hA), M * K); + fillSeq(alpaka::getPtrNative(hB), K * N, 0.02f, 0.0005f); + + auto dA = alpaka::allocAsyncBuf(queue, static_cast(M * K)); + auto dB = alpaka::allocAsyncBuf(queue, static_cast(K * N)); + auto dC = alpaka::allocAsyncBuf(queue, static_cast(M * N)); + alpaka::memcpy(queue, dA, hA); + alpaka::memcpy(queue, dB, hB); + alpaka::wait(queue); + + blas.addLayoutConfig(M, N, K, M, K, M, 'N', 'N'); + + for (int i = 0; i < opt.warmup; ++i) + blas.matmul('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dC); + alpaka::wait(queue); + + auto t0 = std::chrono::steady_clock::now(); + for (int i = 0; i < opt.iterations; ++i) + blas.matmul('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dC); + alpaka::wait(queue); + auto t1 = std::chrono::steady_clock::now(); + + double avgSeconds = + std::chrono::duration(t1 - t0).count() / opt.iterations; + printRow("cuda", size, avgSeconds * 1e3, gflopsFor(size, avgSeconds)); + } +} + +#endif // ALPAKA_ACC_GPU_CUDA_ENABLED + +// --------------------------------------------------------------------------- +// HIP benchmark +// --------------------------------------------------------------------------- +#ifdef ALPAKA_ACC_GPU_HIP_ENABLED + +static void runHipBench(const BenchOptions &opt) { + alpaka::PlatformHipRt platform{}; + auto dev = alpaka::getDevByIdx(platform, 0u); + alpaka::Queue queue{dev}; + sofieBLAS blas(queue); + + alpaka::PlatformCpu hostPlatform{}; + auto hostDev = alpaka::getDevByIdx(hostPlatform, 0u); + + for (int size : opt.sizes) { + const int M = size, N = size, K = size; + auto hA = alpaka::allocBuf(hostDev, static_cast(M * K)); + auto hB = alpaka::allocBuf(hostDev, static_cast(K * N)); + fillSeq(alpaka::getPtrNative(hA), M * K); + fillSeq(alpaka::getPtrNative(hB), K * N, 0.02f, 0.0005f); + + auto dA = alpaka::allocAsyncBuf(queue, static_cast(M * K)); + auto dB = alpaka::allocAsyncBuf(queue, static_cast(K * N)); + auto dC = alpaka::allocAsyncBuf(queue, static_cast(M * N)); + alpaka::memcpy(queue, dA, hA); + alpaka::memcpy(queue, dB, hB); + alpaka::wait(queue); + + blas.addLayoutConfig(M, N, K, M, K, M, 'N', 'N'); + + for (int i = 0; i < opt.warmup; ++i) + blas.matmul('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dC); + alpaka::wait(queue); + + auto t0 = std::chrono::steady_clock::now(); + for (int i = 0; i < opt.iterations; ++i) + blas.matmul('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dC); + alpaka::wait(queue); + auto t1 = std::chrono::steady_clock::now(); + + double avgSeconds = + std::chrono::duration(t1 - t0).count() / opt.iterations; + printRow("hip", size, avgSeconds * 1e3, gflopsFor(size, avgSeconds)); + } +} + +#endif // ALPAKA_ACC_GPU_HIP_ENABLED + +int main(int argc, char **argv) { + BenchOptions opt = parseArgs(argc, argv); + + std::cout << "sofieBLAS benchmark (warmup=" << opt.warmup + << ", iterations=" << opt.iterations << ")\n\n"; + printHeader(); + +#ifdef ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED + runCpuBench(opt); +#endif +#ifdef ALPAKA_ACC_GPU_CUDA_ENABLED + runCudaBench(opt); +#endif +#ifdef ALPAKA_ACC_GPU_HIP_ENABLED + runHipBench(opt); +#endif + + return EXIT_SUCCESS; +} diff --git a/check_style.sh b/check_style.sh index 22a56e4..c937e86 100755 --- a/check_style.sh +++ b/check_style.sh @@ -4,11 +4,12 @@ set -e # Directories SRC_DIR="./include" TEST_DIR="./tests" +BENCH_DIR="./benchmark" echo "📝 Discovering source/header files..." -FILES=$(find "$SRC_DIR" "$TEST_DIR" \ - -path "$TEST_DIR/build" -prune -o \ +FILES=$(find "$SRC_DIR" "$TEST_DIR" "$BENCH_DIR" \ + -path "*/build" -prune -o \ -type f \( \ -name '*.cpp' -o -name '*.cc' -o -name '*.cxx' -o \ -name '*.h' -o -name '*.hpp' -o -name '*.hxx' -o -name '*.hh' \ diff --git a/cmake/SofieBLASBackends.cmake b/cmake/SofieBLASBackends.cmake new file mode 100644 index 0000000..d137d2b --- /dev/null +++ b/cmake/SofieBLASBackends.cmake @@ -0,0 +1,170 @@ +# Shared backend discovery for sofieBLAS tests and benchmarks. + +include(CheckLanguage) +include(FetchContent) + +set(CUDA_BASE "/usr/local/cuda" CACHE PATH "CUDA base path") +set(ROCM_BASE "/opt/rocm" CACHE PATH "ROCm base path") +set(ONEAPI_BASE "/opt/intel/oneapi" CACHE PATH "Intel oneAPI base path") +set(TBB_BASE "/usr" CACHE PATH "TBB base path") +set(BLIS_BASE "/usr" CACHE PATH "BLIS base path") + +set(SOFIEBLAS_ALPAKA_GIT_TAG "2fa91a34ed11b2076e474c5507d920e85cf9b79d" CACHE STRING + "alpaka commit to fetch when no installed alpaka is found") + +find_package(alpaka CONFIG QUIET) +if(NOT alpaka_FOUND) + message(STATUS "sofieBLAS: alpaka not found via find_package, fetching ${SOFIEBLAS_ALPAKA_GIT_TAG}") + FetchContent_Declare( + alpaka + GIT_REPOSITORY https://github.com/alpaka-group/alpaka + GIT_TAG ${SOFIEBLAS_ALPAKA_GIT_TAG} + ) + FetchContent_MakeAvailable(alpaka) +else() + message(STATUS "sofieBLAS: using alpaka ${alpaka_VERSION} from ${alpaka_DIR}") +endif() + +set(SOFIEBLAS_ENABLE_CUDA "AUTO" CACHE STRING "ON/OFF/AUTO: whether to enable the CUDA backend") +set(SOFIEBLAS_ENABLE_HIP "AUTO" CACHE STRING "ON/OFF/AUTO: whether to enable the HIP backend") +set_property(CACHE SOFIEBLAS_ENABLE_CUDA PROPERTY STRINGS ON OFF AUTO) +set_property(CACHE SOFIEBLAS_ENABLE_HIP PROPERTY STRINGS ON OFF AUTO) + +if(NOT SOFIEBLAS_ENABLE_CUDA STREQUAL "OFF") + if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES OR CMAKE_CUDA_ARCHITECTURES STREQUAL "") + set(CMAKE_CUDA_ARCHITECTURES native) + endif() + check_language(CUDA) + if(CMAKE_CUDA_COMPILER) + enable_language(CUDA) + find_package(CUDAToolkit QUIET) + if(CUDAToolkit_FOUND) + set(SOFIEBLAS_CUDA_ENABLED TRUE) + message(STATUS "sofieBLAS: CUDA toolkit ${CUDAToolkit_VERSION} found, CUDA targets enabled") + else() + message(STATUS "sofieBLAS: nvcc found but CUDAToolkit package not found, CUDA targets disabled") + endif() + elseif(SOFIEBLAS_ENABLE_CUDA STREQUAL "ON") + message(FATAL_ERROR "sofieBLAS: SOFIEBLAS_ENABLE_CUDA=ON but no CUDA compiler was found") + else() + message(STATUS "sofieBLAS: no CUDA compiler found, CUDA targets disabled") + endif() +else() + message(STATUS "sofieBLAS: CUDA backend disabled (SOFIEBLAS_ENABLE_CUDA=OFF)") +endif() + +if(NOT SOFIEBLAS_ENABLE_HIP STREQUAL "OFF") + if(NOT DEFINED CMAKE_HIP_ARCHITECTURES OR CMAKE_HIP_ARCHITECTURES STREQUAL "") + set(CMAKE_HIP_ARCHITECTURES gfx1100) + endif() + check_language(HIP) + if(CMAKE_HIP_COMPILER) + enable_language(HIP) + set(SOFIEBLAS_HIP_ENABLED TRUE) + message(STATUS "sofieBLAS: HIP compiler found, HIP targets enabled") + elseif(SOFIEBLAS_ENABLE_HIP STREQUAL "ON") + message(FATAL_ERROR "sofieBLAS: SOFIEBLAS_ENABLE_HIP=ON but no HIP compiler was found") + else() + message(STATUS "sofieBLAS: no HIP compiler found, HIP targets disabled") + endif() +else() + message(STATUS "sofieBLAS: HIP backend disabled (SOFIEBLAS_ENABLE_HIP=OFF)") +endif() + +# --- CPU BLAS selection --- +set(AVAILABLE_BLAS_LIBS OpenBLAS MKL BLIS Accelerate CACHE STRING "Choose CPU BLAS library") +set_property(CACHE AVAILABLE_BLAS_LIBS PROPERTY STRINGS OpenBLAS MKL BLIS Accelerate) + +if(NOT DEFINED CPU_BLAS_LIB) + set(CPU_BLAS_LIB "OpenBLAS" CACHE STRING "CPU BLAS library to use") +endif() + +function(sofieblas_find_openblas out_lib out_define out_include) + find_library(OPENBLAS_LIB NAMES openblas + PATHS /usr/lib/x86_64-linux-gnu/openblas-pthread + /usr/lib/x86_64-linux-gnu/openblas-serial + /usr/lib/x86_64-linux-gnu + /usr/local/lib) + find_path(OPENBLAS_INCLUDE_DIR NAMES cblas.h + PATHS /usr/include + /usr/include/x86_64-linux-gnu/openblas-pthread + /usr/include/x86_64-linux-gnu/openblas-serial + /usr/include/openblas + /usr/local/include) + if(OPENBLAS_LIB AND OPENBLAS_INCLUDE_DIR) + set(${out_lib} ${OPENBLAS_LIB} PARENT_SCOPE) + set(${out_define} SOFIEBLAS_USE_OPENBLAS PARENT_SCOPE) + set(${out_include} ${OPENBLAS_INCLUDE_DIR} PARENT_SCOPE) + else() + set(${out_lib} "" PARENT_SCOPE) + set(${out_define} "" PARENT_SCOPE) + set(${out_include} "" PARENT_SCOPE) + endif() +endfunction() + +function(sofieblas_find_mkl out_lib out_define out_include) + find_library(MKL_LIB NAMES mkl_rt PATHS ${ONEAPI_BASE}/mkl/latest/lib/intel64 ${ONEAPI_BASE}/mkl/latest/lib) + find_path(MKL_INCLUDE_DIR NAMES mkl.h PATHS ${ONEAPI_BASE}/mkl/latest/include) + if(MKL_LIB AND MKL_INCLUDE_DIR) + set(${out_lib} ${MKL_LIB} PARENT_SCOPE) + set(${out_define} SOFIEBLAS_USE_MKL PARENT_SCOPE) + set(${out_include} ${MKL_INCLUDE_DIR} PARENT_SCOPE) + else() + set(${out_lib} "" PARENT_SCOPE) + set(${out_define} "" PARENT_SCOPE) + set(${out_include} "" PARENT_SCOPE) + endif() +endfunction() + +function(sofieblas_find_blis out_lib out_define out_include) + find_library(BLIS_LIB NAMES blis PATHS ${BLIS_BASE}/lib ${BLIS_BASE}/lib64) + find_path(BLIS_INCLUDE_DIR NAMES blis/cblas.h PATHS ${BLIS_BASE}/include) + if(BLIS_LIB AND BLIS_INCLUDE_DIR) + set(${out_lib} ${BLIS_LIB} PARENT_SCOPE) + set(${out_define} SOFIEBLAS_USE_BLIS PARENT_SCOPE) + set(${out_include} ${BLIS_INCLUDE_DIR} PARENT_SCOPE) + else() + set(${out_lib} "" PARENT_SCOPE) + set(${out_define} "" PARENT_SCOPE) + set(${out_include} "" PARENT_SCOPE) + endif() +endfunction() + +function(sofieblas_find_accelerate out_lib out_define out_include) + set(${out_lib} "" PARENT_SCOPE) + set(${out_define} "" PARENT_SCOPE) + set(${out_include} "" PARENT_SCOPE) + if(APPLE) + find_library(ACCELERATE_LIB Accelerate) + if(ACCELERATE_LIB) + set(${out_lib} ${ACCELERATE_LIB} PARENT_SCOPE) + set(${out_define} SOFIEBLAS_USE_ACCELERATE PARENT_SCOPE) + endif() + endif() +endfunction() + +set(_search_order OpenBLAS MKL BLIS Accelerate) +if(NOT CPU_BLAS_LIB IN_LIST _search_order) + message(FATAL_ERROR "Unknown CPU_BLAS_LIB option: ${CPU_BLAS_LIB}") +endif() +list(REMOVE_ITEM _search_order ${CPU_BLAS_LIB}) +list(PREPEND _search_order ${CPU_BLAS_LIB}) + +foreach(_lib IN LISTS _search_order) + if(NOT SOFIEBLAS_CPU_BLAS_LIBS) + string(TOLOWER ${_lib} _lib_lower) + cmake_language(CALL sofieblas_find_${_lib_lower} + SOFIEBLAS_CPU_BLAS_LIBS SOFIEBLAS_CPU_BLAS_DEFINE SOFIEBLAS_CPU_BLAS_INCLUDE_DIR) + if(SOFIEBLAS_CPU_BLAS_LIBS) + set(SOFIEBLAS_CPU_BLAS_FOUND ${_lib}) + else() + message(STATUS "sofieBLAS: ${_lib} not found, trying next CPU BLAS option...") + endif() + endif() +endforeach() + +if(NOT SOFIEBLAS_CPU_BLAS_LIBS) + message(WARNING "sofieBLAS: no suitable CPU BLAS library found (tried: ${_search_order}). CPU targets disabled.") +else() + message(STATUS "sofieBLAS: using CPU BLAS library ${SOFIEBLAS_CPU_BLAS_FOUND} (${SOFIEBLAS_CPU_BLAS_LIBS})") +endif() diff --git a/cmake/sofieBLASConfig.cmake.in b/cmake/sofieBLASConfig.cmake.in new file mode 100644 index 0000000..78a6644 --- /dev/null +++ b/cmake/sofieBLASConfig.cmake.in @@ -0,0 +1,5 @@ +@PACKAGE_INIT@ + +include("${CMAKE_CURRENT_LIST_DIR}/sofieBLASTargets.cmake") + +check_required_components(sofieBLAS) diff --git a/include/.vscode/settings.json b/include/.vscode/settings.json deleted file mode 100644 index 64f45f4..0000000 --- a/include/.vscode/settings.json +++ /dev/null @@ -1,8 +0,0 @@ -{ - "files.associations": { - "array": "cpp", - "string": "cpp", - "string_view": "cpp", - "span": "cpp" - } -} diff --git a/include/sofieBLAS/backends/cpu/detail/sofieBLAS_cblas_common.hpp b/include/sofieBLAS/backends/cpu/detail/sofieBLAS_cblas_common.hpp new file mode 100644 index 0000000..ac2d1a2 --- /dev/null +++ b/include/sofieBLAS/backends/cpu/detail/sofieBLAS_cblas_common.hpp @@ -0,0 +1,272 @@ +#pragma once + +// Shared implementation for every CPU backend that exposes a standard CBLAS +// API. The vendor-specific header (cblas.h / mkl.h / blis/cblas.h / +// Accelerate.h) must already be included by the backend wrapper before this +// file, since they all provide the same CBLAS symbols. + +#ifdef ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED + +#include "sofieBLAS/core.hpp" +#include + +#include +#include +#include + +class BlasCpu { +public: + BlasCpu(alpaka::QueueCpuBlocking &queue [[maybe_unused]]) {} + + inline CBLAS_TRANSPOSE charToTranspose(char trans) { + switch (trans) { + case 'N': + case 'n': + return CblasNoTrans; + case 'T': + case 't': + return CblasTrans; + case 'C': + case 'c': + return CblasConjTrans; + default: + throw std::invalid_argument("Invalid transpose character."); + } + } + + // C = alpha * op(A) * op(B) + beta * C (no bias, leading dims inferred) + template + inline void matmul(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, + alpaka::BufCpu, TIdx> const &A, + alpaka::BufCpu, TIdx> const &B, + float beta, + alpaka::BufCpu, TIdx> &C) { + int lda = (transa == 'N' || transa == 'n') ? static_cast(m) + : static_cast(k); + int ldb = (transb == 'N' || transb == 'n') ? static_cast(k) + : static_cast(n); + cblas_sgemm(CblasColMajor, charToTranspose(transa), charToTranspose(transb), + static_cast(m), static_cast(n), static_cast(k), + alpha, alpaka::getPtrNative(A), lda, alpaka::getPtrNative(B), + ldb, beta, alpaka::getPtrNative(C), static_cast(m)); + } + + template + inline void + matmul(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> &C) { + int lda = (transa == 'N' || transa == 'n') ? static_cast(m) + : static_cast(k); + int ldb = (transb == 'N' || transb == 'n') ? static_cast(k) + : static_cast(n); + cblas_sgemm(CblasColMajor, charToTranspose(transa), charToTranspose(transb), + static_cast(m), static_cast(n), static_cast(k), + alpha, alpaka::getPtrNative(A), lda, alpaka::getPtrNative(B), + ldb, beta, alpaka::getPtrNative(C), static_cast(m)); + } + + // C = alpha * op(A) * op(B) + beta * bias + bias_vec (bias_vec broadcast per + // row) + template + inline void + gemm(char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, alpaka::BufCpu, TIdx> const &A, + alpaka::BufCpu, TIdx> const &B, float beta, + alpaka::BufCpu, TIdx> &bias, + alpaka::BufCpu, TIdx> &C) { + int lda = (transa == 'N' || transa == 'n') ? static_cast(m) + : static_cast(k); + int ldb = (transb == 'N' || transb == 'n') ? static_cast(k) + : static_cast(n); + cblas_sgemm(CblasColMajor, charToTranspose(transa), charToTranspose(transb), + static_cast(m), static_cast(n), static_cast(k), + alpha, alpaka::getPtrNative(A), lda, alpaka::getPtrNative(B), + ldb, 0.0f, alpaka::getPtrNative(C), static_cast(m)); + float *c = alpaka::getPtrNative(C); + const float *b = alpaka::getPtrNative(bias); + for (unsigned int j = 0; j < n; ++j) + for (unsigned int i = 0; i < m; ++i) + c[j * m + i] += beta * b[j * m + i] + b[i]; + } + + template + inline void + gemm(char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> &bias, + alpaka::ViewPlainPtr, TIdx> &C) { + int lda = (transa == 'N' || transa == 'n') ? static_cast(m) + : static_cast(k); + int ldb = (transb == 'N' || transb == 'n') ? static_cast(k) + : static_cast(n); + cblas_sgemm(CblasColMajor, charToTranspose(transa), charToTranspose(transb), + static_cast(m), static_cast(n), static_cast(k), + alpha, alpaka::getPtrNative(A), lda, alpaka::getPtrNative(B), + ldb, 0.0f, alpaka::getPtrNative(C), static_cast(m)); + float *c = alpaka::getPtrNative(C); + const float *b = alpaka::getPtrNative(bias); + for (unsigned int j = 0; j < n; ++j) + for (unsigned int i = 0; i < m; ++i) + c[j * m + i] += beta * b[j * m + i] + b[i]; + } + + // C = relu(alpha * op(A) * op(B) + beta * bias + bias_vec) + template + inline void gemmrelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, + alpaka::BufCpu, TIdx> const &A, + alpaka::BufCpu, TIdx> const &B, + float beta, + alpaka::BufCpu, TIdx> &bias, + alpaka::BufCpu, TIdx> &C) { + gemm(transa, transb, m, n, k, alpha, A, B, beta, bias, C); + float *c = alpaka::getPtrNative(C); + for (unsigned int i = 0; i < m * n; ++i) + c[i] = c[i] > 0.0f ? c[i] : 0.0f; + } + + template + inline void gemmrelu( + char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> &bias, + alpaka::ViewPlainPtr, TIdx> &C) { + gemm(transa, transb, m, n, k, alpha, A, B, beta, bias, C); + float *c = alpaka::getPtrNative(C); + for (unsigned int i = 0; i < m * n; ++i) + c[i] = c[i] > 0.0f ? c[i] : 0.0f; + } + + // C = gelu(alpha * op(A) * op(B) + beta * bias + bias_vec) + // Uses the standard GELU: x * 0.5 * (1 + erf(x / sqrt(2))) + template + inline void gemmgelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, + alpaka::BufCpu, TIdx> const &A, + alpaka::BufCpu, TIdx> const &B, + float beta, + alpaka::BufCpu, TIdx> &bias, + alpaka::BufCpu, TIdx> &C) { + gemm(transa, transb, m, n, k, alpha, A, B, beta, bias, C); + float *c = alpaka::getPtrNative(C); + constexpr float kInvSqrt2 = 0.7071067811865476f; + for (unsigned int i = 0; i < m * n; ++i) + c[i] *= 0.5f * (1.0f + std::erff(c[i] * kInvSqrt2)); + } + + template + inline void gemmgelu( + char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> &bias, + alpaka::ViewPlainPtr, TIdx> &C) { + gemm(transa, transb, m, n, k, alpha, A, B, beta, bias, C); + float *c = alpaka::getPtrNative(C); + constexpr float kInvSqrt2 = 0.7071067811865476f; + for (unsigned int i = 0; i < m * n; ++i) + c[i] *= 0.5f * (1.0f + std::erff(c[i] * kInvSqrt2)); + } + + // Raw-pointer overloads: accept T const*/T* from any BufXxx or ViewPlainPtr + // via getPtrNative() + template + inline void matmul(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *C) { + int lda = (transa == 'N' || transa == 'n') ? static_cast(m) + : static_cast(k); + int ldb = (transb == 'N' || transb == 'n') ? static_cast(k) + : static_cast(n); + cblas_sgemm(CblasColMajor, charToTranspose(transa), charToTranspose(transb), + static_cast(m), static_cast(n), static_cast(k), + alpha, A, lda, B, ldb, beta, C, static_cast(m)); + } + + template + inline void gemm(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *bias, T *C) { + int lda = (transa == 'N' || transa == 'n') ? static_cast(m) + : static_cast(k); + int ldb = (transb == 'N' || transb == 'n') ? static_cast(k) + : static_cast(n); + cblas_sgemm(CblasColMajor, charToTranspose(transa), charToTranspose(transb), + static_cast(m), static_cast(n), static_cast(k), + alpha, A, lda, B, ldb, 0.0f, C, static_cast(m)); + for (unsigned int j = 0; j < n; ++j) + for (unsigned int i = 0; i < m; ++i) + C[j * m + i] += beta * bias[j * m + i] + bias[i]; + } + + template + inline void gemmrelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *bias, T *C) { + gemm(transa, transb, m, n, k, alpha, A, B, beta, bias, C); + for (unsigned int i = 0; i < m * n; ++i) + C[i] = C[i] > 0.0f ? C[i] : 0.0f; + } + + template + inline void gemmgelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *bias, T *C) { + gemm(transa, transb, m, n, k, alpha, A, B, beta, bias, C); + constexpr float kInvSqrt2 = 0.7071067811865476f; + for (unsigned int i = 0; i < m * n; ++i) + C[i] *= 0.5f * (1.0f + std::erff(C[i] * kInvSqrt2)); + } +}; + +namespace traits { + +template <> class sofieBLAS { +public: + using Impl = BlasCpu; +}; + +template <> class sofieBLAS { +public: + using Impl = BlasCpu; +}; + +template <> class sofieBLAS { +public: + using Impl = BlasCpu; +}; + +template <> class sofieBLAS { +public: + using Impl = BlasCpu; +}; + +template <> class sofieBLAS { +public: + using Impl = BlasCpu; +}; + +} // namespace traits + +#endif // ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED diff --git a/include/sofieBLAS/backends/cpu/sofieBLAS_accelerate.hpp b/include/sofieBLAS/backends/cpu/sofieBLAS_accelerate.hpp new file mode 100644 index 0000000..9ba7b6b --- /dev/null +++ b/include/sofieBLAS/backends/cpu/sofieBLAS_accelerate.hpp @@ -0,0 +1,11 @@ +#pragma once + +#if defined(ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED) && \ + defined(SOFIEBLAS_USE_ACCELERATE) + +// Apple's Accelerate framework (macOS only) exposes the same CBLAS API. +#include + +#include "sofieBLAS/backends/cpu/detail/sofieBLAS_cblas_common.hpp" + +#endif // ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED && SOFIEBLAS_USE_ACCELERATE diff --git a/include/sofieBLAS/backends/cpu/sofieBLAS_blis.hpp b/include/sofieBLAS/backends/cpu/sofieBLAS_blis.hpp new file mode 100644 index 0000000..952bb51 --- /dev/null +++ b/include/sofieBLAS/backends/cpu/sofieBLAS_blis.hpp @@ -0,0 +1,11 @@ +#pragma once + +#if defined(ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED) && defined(SOFIEBLAS_USE_BLIS) + +// BLIS installs its CBLAS-compatible header under blis/cblas.h +// when configured with --enable-cblas. +#include + +#include "sofieBLAS/backends/cpu/detail/sofieBLAS_cblas_common.hpp" + +#endif // ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED && SOFIEBLAS_USE_BLIS diff --git a/include/sofieBLAS/backends/cpu/sofieBLAS_cpu.hpp b/include/sofieBLAS/backends/cpu/sofieBLAS_cpu.hpp deleted file mode 100644 index 263703c..0000000 --- a/include/sofieBLAS/backends/cpu/sofieBLAS_cpu.hpp +++ /dev/null @@ -1,83 +0,0 @@ -#pragma once - -#ifdef ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED - -#include "sofieBLAS/core.hpp" -#include - -#if defined(SOFIEBLAS_USE_MKL) -#include -#elif defined(SOFIEBLAS_USE_OPENBLAS) -#include -#else -#error \ - "No CPU BLAS backend selected. Define SOFIEBLAS_USE_MKL or SOFIEBLAS_USE_OPENBLAS." -#endif - -#include - -class BlasCpu { -public: - BlasCpu(alpaka::QueueCpuBlocking &queue [[maybe_unused]]) {} - - inline CBLAS_TRANSPOSE charToTranspose(char trans) { - switch (trans) { - case 'N': - case 'n': - return CblasNoTrans; - case 'T': - case 't': - return CblasTrans; - case 'C': - case 'c': - return CblasConjTrans; - default: - throw std::invalid_argument("Invalid transpose character."); - } - } - - template - inline void - gemm(char transa, char transb, const unsigned int m, const unsigned int n, - const unsigned int k, const float alpha, - alpaka::BufCpu, TIdx> const &A, const int lda, - alpaka::BufCpu, TIdx> const &B, const int ldb, - const float beta, alpaka::BufCpu, TIdx> &C, - const int ldc) { - CBLAS_TRANSPOSE TransA = charToTranspose(transa); - CBLAS_TRANSPOSE TransB = charToTranspose(transb); - cblas_sgemm(CblasColMajor, TransA, TransB, m, n, k, alpha, A.data(), lda, - B.data(), ldb, beta, C.data(), ldc); - } -}; - -namespace traits { - -template <> class sofieBLAS { -public: - using Impl = BlasCpu; -}; - -template <> class sofieBLAS { -public: - using Impl = BlasCpu; -}; - -template <> class sofieBLAS { -public: - using Impl = BlasCpu; -}; - -template <> class sofieBLAS { -public: - using Impl = BlasCpu; -}; - -template <> class sofieBLAS { -public: - using Impl = BlasCpu; -}; - -} // namespace traits - -#endif // ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED diff --git a/include/sofieBLAS/backends/cpu/sofieBLAS_mkl.hpp b/include/sofieBLAS/backends/cpu/sofieBLAS_mkl.hpp new file mode 100644 index 0000000..66e7d1f --- /dev/null +++ b/include/sofieBLAS/backends/cpu/sofieBLAS_mkl.hpp @@ -0,0 +1,9 @@ +#pragma once + +#if defined(ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED) && defined(SOFIEBLAS_USE_MKL) + +#include + +#include "sofieBLAS/backends/cpu/detail/sofieBLAS_cblas_common.hpp" + +#endif // ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED && SOFIEBLAS_USE_MKL diff --git a/include/sofieBLAS/backends/cpu/sofieBLAS_openblas.hpp b/include/sofieBLAS/backends/cpu/sofieBLAS_openblas.hpp new file mode 100644 index 0000000..3f32cb9 --- /dev/null +++ b/include/sofieBLAS/backends/cpu/sofieBLAS_openblas.hpp @@ -0,0 +1,10 @@ +#pragma once + +#if defined(ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED) && \ + defined(SOFIEBLAS_USE_OPENBLAS) + +#include + +#include "sofieBLAS/backends/cpu/detail/sofieBLAS_cblas_common.hpp" + +#endif // ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED && SOFIEBLAS_USE_OPENBLAS diff --git a/include/sofieBLAS/backends/cuda/sofieBLAS_cublas.hpp b/include/sofieBLAS/backends/cuda/sofieBLAS_cublas.hpp index cbf18e0..ded6ba8 100644 --- a/include/sofieBLAS/backends/cuda/sofieBLAS_cublas.hpp +++ b/include/sofieBLAS/backends/cuda/sofieBLAS_cublas.hpp @@ -15,7 +15,7 @@ #include #define CHECK_CUDA(err) \ - if (err != cudaSuccess) { \ + if ((err) != cudaSuccess) { \ std::cerr << "CUDA error: " << cudaGetErrorString(err) << " at line " \ << __LINE__ << "\n"; \ exit(EXIT_FAILURE); \ @@ -23,10 +23,9 @@ #define CHECK_CUBLAS(status) \ do { \ - cublasStatus_t s = (status); \ - if (s != CUBLAS_STATUS_SUCCESS) { \ - std::cerr << "cuBLAS error " << s << " at line " << __LINE__ \ - << std::endl; \ + cublasStatus_t _s = (status); \ + if (_s != CUBLAS_STATUS_SUCCESS) { \ + std::cerr << "cuBLAS error " << _s << " at line " << __LINE__ << "\n"; \ exit(EXIT_FAILURE); \ } \ } while (0) @@ -47,35 +46,80 @@ struct PairEq { } }; +struct DescKey { + int transA; // CUBLAS_OP_N / CUBLAS_OP_T encoded as int + int transB; + int epilogue; // cublasLtEpilogue_t encoded as int + bool operator==(const DescKey &o) const noexcept { + return transA == o.transA && transB == o.transB && epilogue == o.epilogue; + } +}; + +struct DescKeyHash { + std::size_t operator()(const DescKey &k) const noexcept { + std::size_t h = static_cast(k.transA) * 97u + + static_cast(k.transB) * 31u + + static_cast(k.epilogue); + return h ^ (h >> 16); + } +}; + +struct AlgoKey { + DescKey dk; + std::size_t rowsA, colsA; // physical dimensions of A in layoutStore + std::size_t rowsB, colsB; // physical dimensions of B in layoutStore + bool operator==(const AlgoKey &o) const noexcept { + return dk == o.dk && rowsA == o.rowsA && colsA == o.colsA && + rowsB == o.rowsB && colsB == o.colsB; + } +}; + +struct AlgoKeyHash { + std::size_t operator()(const AlgoKey &k) const noexcept { + std::size_t h = DescKeyHash{}(k.dk); + auto mix = [&](std::size_t v) { + h ^= std::hash{}(v) + 0x9e3779b97f4a7c15ULL + (h << 6) + + (h >> 2); + }; + mix(k.rowsA); + mix(k.colsA); + mix(k.rowsB); + mix(k.colsB); + return h; + } +}; + class BlasCuda { cublasLtHandle_t ltHandle = nullptr; cublasHandle_t handle = nullptr; - cublasLtMatmulDesc_t operationDesc = nullptr; cublasLtMatmulPreference_t preference = nullptr; void *d_workspace = nullptr; - size_t workspaceSize = 1 << 22; // 4MB + size_t workspaceSize = 1u << 25; // 32 MB cudaStream_t stream = nullptr; - cublasLtMatmulHeuristicResult_t heuristic; - cublasLtEpilogue_t epilogue = CUBLASLT_EPILOGUE_DEFAULT; - int error_flag = 0; std::unordered_map, cublasLtMatrixLayout_t, PairHash, PairEq> - LayoutStore; + layoutStore; + + std::unordered_map descStore; + + std::unordered_map + algoCache; public: - BlasCuda(const BlasCuda&) = delete; - BlasCuda& operator=(const BlasCuda&) = delete; - BlasCuda(BlasCuda&&) = delete; - BlasCuda& operator=(BlasCuda&&) = delete; + BlasCuda(const BlasCuda &) = delete; + BlasCuda &operator=(const BlasCuda &) = delete; + BlasCuda(BlasCuda &&) = delete; + BlasCuda &operator=(BlasCuda &&) = delete; BlasCuda(alpaka::QueueCudaRtNonBlocking &queue) : m_queue{queue} { stream = static_cast(m_queue.getNativeHandle()); + CHECK_CUBLAS(cublasLtCreate(<Handle)); + CHECK_CUBLAS(cublasCreate(&handle)); - heuristic = {}; - CHECK_CUBLAS(cublasLtMatmulDescCreate(&operationDesc, CUBLAS_COMPUTE_32F, - CUDA_R_32F)); + CHECK_CUBLAS(cublasSetStream(handle, stream)); + CHECK_CUBLAS(cublasLtMatmulPreferenceCreate(&preference)); CHECK_CUDA(cudaMalloc(&d_workspace, workspaceSize)); CHECK_CUBLAS(cublasLtMatmulPreferenceSetAttribute( @@ -84,22 +128,20 @@ class BlasCuda { } ~BlasCuda() { - for (auto& [key, layout] : LayoutStore) { - if (layout) { + for (auto &[key, layout] : layoutStore) + if (layout) cublasLtMatrixLayoutDestroy(layout); - } - } - LayoutStore.clear(); - + for (auto &[key, desc] : descStore) + if (desc) + cublasLtMatmulDescDestroy(desc); if (preference) cublasLtMatmulPreferenceDestroy(preference); - if (operationDesc) - cublasLtMatmulDescDestroy(operationDesc); if (ltHandle) cublasLtDestroy(ltHandle); + if (handle) + cublasDestroy(handle); if (d_workspace) cudaFree(d_workspace); - } inline cublasOperation_t charToCuBlasTranspose(char trans) { @@ -118,231 +160,321 @@ class BlasCuda { } } - void AddLayoutConfig(std::size_t m, std::size_t n, std::size_t k) { - CheckAndAddLayout(k, m); - CheckAndAddLayout(k, n); - CheckAndAddLayout(m, n); + void addLayoutConfig(std::size_t m, std::size_t n, std::size_t k, + std::size_t lda, std::size_t ldb, std::size_t ldc, + char transa, char transb) { + // Physical A: (m×k) if NoTrans, (k×m) if Trans + if (transa == 'N' || transa == 'n') + checkAndAddLayout(m, k, lda); + else + checkAndAddLayout(k, m, lda); + // Physical B: (k×n) if NoTrans, (n×k) if Trans + if (transb == 'N' || transb == 'n') + checkAndAddLayout(k, n, ldb); + else + checkAndAddLayout(n, k, ldb); + // C is always (m×n) + checkAndAddLayout(m, n, ldc); } -template -inline void -gemm(char transa, char transb, const unsigned int m, - const unsigned int n, const unsigned int k, - const float alpha, - alpaka::BufCudaRt, TIdx> const &A, - alpaka::BufCudaRt, TIdx> const &B, - const float beta, - alpaka::BufCudaRt, TIdx> &bias, - alpaka::BufCudaRt, TIdx> &C) -{ - cublasLtMatmulDesc_t localDesc = nullptr; - CHECK_CUBLAS(cublasLtMatmulDescCreate(&localDesc, CUBLAS_COMPUTE_32F, CUDA_R_32F)); - - cublasOperation_t transB_op = charToCuBlasTranspose(transb); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_TRANSB, &transB_op, sizeof(transB_op))); - - cublasOperation_t transA_op = charToCuBlasTranspose(transa); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_TRANSA, &transA_op, sizeof(transA_op))); - - void *bias_ptr = reinterpret_cast(alpaka::getPtrNative(bias)); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_BIAS_POINTER, &bias_ptr, sizeof(bias_ptr))); - - cublasLtEpilogue_t ep = CUBLASLT_EPILOGUE_BIAS; - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, - CUBLASLT_MATMUL_DESC_EPILOGUE, - &ep, - sizeof(ep))); - - - cublasLtMatmulHeuristicResult_t localHeuristic{}; - int returnedResults = 0; - CHECK_CUBLAS(cublasLtMatmulAlgoGetHeuristic( - ltHandle, - localDesc, - LayoutStore.at({k, m}), - LayoutStore.at({k, n}), - LayoutStore.at({m, n}), - LayoutStore.at({m, n}), - preference, - 1, - &localHeuristic, - &returnedResults)); - - if (returnedResults == 0) { - cublasLtMatmulDescDestroy(localDesc); - std::cerr << "No suitable cuBLASLt algorithm found!\n"; - exit(EXIT_FAILURE); - } - - CHECK_CUBLAS(cublasLtMatmul( - ltHandle, - localDesc, - &alpha, - alpaka::getPtrNative(A), LayoutStore.at({k, m}), - alpaka::getPtrNative(B), LayoutStore.at({k, n}), - &beta, - alpaka::getPtrNative(bias), LayoutStore.at({m, n}), - alpaka::getPtrNative(C), LayoutStore.at({m, n}), - &(localHeuristic.algo), - d_workspace, - workspaceSize, - stream)); - - cudaDeviceSynchronize(); - CHECK_CUBLAS(cublasLtMatmulDescDestroy(localDesc)); -} - -template -inline void -gemmrelu(char transa, char transb, const unsigned int m, - const unsigned int n, const unsigned int k, - const float alpha, - alpaka::BufCudaRt, TIdx> const &A, - alpaka::BufCudaRt, TIdx> const &B, - const float beta, - alpaka::BufCudaRt, TIdx> &bias, - alpaka::BufCudaRt, TIdx> &C) -{ - cublasLtMatmulDesc_t localDesc = nullptr; - CHECK_CUBLAS(cublasLtMatmulDescCreate(&localDesc, CUBLAS_COMPUTE_32F, CUDA_R_32F)); - - cublasOperation_t transB_op = charToCuBlasTranspose(transb); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_TRANSB, &transB_op, sizeof(transB_op))); + template + inline void + gemm(char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, alpaka::BufCudaRt, TIdx> const &A, + alpaka::BufCudaRt, TIdx> const &B, float beta, + alpaka::BufCudaRt, TIdx> &bias, + alpaka::BufCudaRt, TIdx> &C) { + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_BIAS, alpha, alpaka::getPtrNative(A), + alpaka::getPtrNative(B), beta, alpaka::getPtrNative(bias), + alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } - cublasOperation_t transA_op = charToCuBlasTranspose(transa); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_TRANSA, &transA_op, sizeof(transA_op))); + template + inline void gemm( + char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> + &bias, + alpaka::ViewPlainPtr, TIdx> &C) { + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_BIAS, alpha, alpaka::getPtrNative(A), + alpaka::getPtrNative(B), beta, alpaka::getPtrNative(bias), + alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } - void *bias_ptr = reinterpret_cast(alpaka::getPtrNative(bias)); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_BIAS_POINTER, &bias_ptr, sizeof(bias_ptr))); + template + inline void gemm(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *bias, T *C) { + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_BIAS, alpha, A, B, beta, bias, C, + static_cast(bias), layoutKeyA(transa, m, k), + layoutKeyB(transb, k, n), {m, n}); + } - cublasLtEpilogue_t ep = CUBLASLT_EPILOGUE_RELU_BIAS; - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_EPILOGUE, &ep, sizeof(ep))); + template + inline void gemmrelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, + alpaka::BufCudaRt, TIdx> const &A, + alpaka::BufCudaRt, TIdx> const &B, + float beta, + alpaka::BufCudaRt, TIdx> &bias, + alpaka::BufCudaRt, TIdx> &C) { + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_RELU_BIAS, alpha, alpaka::getPtrNative(A), + alpaka::getPtrNative(B), beta, alpaka::getPtrNative(bias), + alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } - cublasLtMatmulHeuristicResult_t localHeuristic{}; - CHECK_CUBLAS(cublasLtMatmulAlgoGetHeuristic( - ltHandle, - localDesc, - LayoutStore.at({k, m}), - LayoutStore.at({k, n}), - LayoutStore.at({m, n}), - LayoutStore.at({m, n}), - preference, - 1, - &localHeuristic, - &error_flag)); - - if (error_flag == 0) { - cublasLtMatmulDescDestroy(localDesc); - std::cerr << "No suitable cuBLASLt algorithm found!\n"; - exit(EXIT_FAILURE); - } + template + inline void gemmrelu( + char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> + &bias, + alpaka::ViewPlainPtr, TIdx> &C) { + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_RELU_BIAS, alpha, alpaka::getPtrNative(A), + alpaka::getPtrNative(B), beta, alpaka::getPtrNative(bias), + alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } - CHECK_CUBLAS(cublasLtMatmul( - ltHandle, - localDesc, - &alpha, - alpaka::getPtrNative(A), LayoutStore.at({k, m}), - alpaka::getPtrNative(B), LayoutStore.at({k, n}), - &beta, - alpaka::getPtrNative(bias), LayoutStore.at({m, n}), - alpaka::getPtrNative(C), LayoutStore.at({m, n}), - &(localHeuristic.algo), - d_workspace, - workspaceSize, - stream)); - - cudaDeviceSynchronize(); - CHECK_CUBLAS(cublasLtMatmulDescDestroy(localDesc)); -} + template + inline void gemmrelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *bias, T *C) { + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_RELU_BIAS, alpha, A, B, beta, bias, C, + static_cast(bias), layoutKeyA(transa, m, k), + layoutKeyB(transb, k, n), {m, n}); + } template - inline void gemmgelu(char transa, char transb, const unsigned int m, - const unsigned int n, const unsigned int k, - const float alpha, + inline void gemmgelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, alpaka::BufCudaRt, TIdx> const &A, alpaka::BufCudaRt, TIdx> const &B, - const float beta, + float beta, alpaka::BufCudaRt, TIdx> &bias, alpaka::BufCudaRt, TIdx> &C) { + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_GELU_BIAS, alpha, alpaka::getPtrNative(A), + alpaka::getPtrNative(B), beta, alpaka::getPtrNative(bias), + alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } - cublasLtMatmulDesc_t localDesc = nullptr; - CHECK_CUBLAS(cublasLtMatmulDescCreate(&localDesc, CUBLAS_COMPUTE_32F, CUDA_R_32F)); - - void *bias_ptr = reinterpret_cast(alpaka::getPtrNative(bias)); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_BIAS_POINTER, &bias_ptr, - sizeof(bias_ptr))); + template + inline void gemmgelu( + char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> + &bias, + alpaka::ViewPlainPtr, TIdx> &C) { + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_GELU_BIAS, alpha, alpaka::getPtrNative(A), + alpaka::getPtrNative(B), beta, alpaka::getPtrNative(bias), + alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } - cublasOperation_t transB = charToCuBlasTranspose(transb); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_TRANSB, &transB, sizeof(transB))); + template + inline void gemmgelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *bias, T *C) { + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_GELU_BIAS, alpha, A, B, beta, bias, C, + static_cast(bias), layoutKeyA(transa, m, k), + layoutKeyB(transb, k, n), {m, n}); + } - cublasOperation_t transA = charToCuBlasTranspose(transa); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_TRANSA, &transA, sizeof(transA))); - SetGeluActivation(); + template + inline void matmul(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, + alpaka::BufCudaRt, TIdx> const &A, + alpaka::BufCudaRt, TIdx> const &B, + float beta, + alpaka::BufCudaRt, TIdx> &C) { + float *c = alpaka::getPtrNative(C); + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_DEFAULT, alpha, alpaka::getPtrNative(A), + alpaka::getPtrNative(B), beta, c, c, nullptr, + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } - cublasLtMatmulHeuristicResult_t localHeuristic{}; - CHECK_CUBLAS(cublasLtMatmulAlgoGetHeuristic( - ltHandle, localDesc, - LayoutStore.at({k, m}), - LayoutStore.at({k, n}), - LayoutStore.at({m, n}), - LayoutStore.at({m, n}), - preference, 1, &localHeuristic, &error_flag)); - if (error_flag == 0) { - std::cerr << "No suitable cuBLASLt algorithm found!\n"; - exit(EXIT_FAILURE); - } + template + inline void matmul( + char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> &C) { + T *c = alpaka::getPtrNative(C); + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_DEFAULT, alpha, alpaka::getPtrNative(A), + alpaka::getPtrNative(B), beta, c, c, nullptr, + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } - CHECK_CUBLAS(cublasLtMatmul( - ltHandle, localDesc, &alpha, alpaka::getPtrNative(A), LayoutStore.at({k, m}), - alpaka::getPtrNative(B), LayoutStore.at({k, n}), &beta, alpaka::getPtrNative(bias), LayoutStore.at({m, n}), - alpaka::getPtrNative(C), LayoutStore.at({m, n}), &(localHeuristic.algo), d_workspace, - workspaceSize, stream)); + // Raw-pointer overload + template + inline void matmul(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *C) { + executeMatmul(charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), + CUBLASLT_EPILOGUE_DEFAULT, alpha, A, B, beta, C, C, nullptr, + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } + + inline void gemmStridedBatched(char transa, char transb, int m, int n, int k, + float alpha, const float *A, int lda, + long long strideA, const float *B, int ldb, + long long strideB, float beta, float *C, + int ldc, long long strideC, int batchCount) { + CHECK_CUBLAS(cublasSgemmStridedBatched( + handle, charToCuBlasTranspose(transa), charToCuBlasTranspose(transb), m, + n, k, &alpha, A, lda, strideA, B, ldb, strideB, &beta, C, ldc, strideC, + batchCount)); } private: alpaka::QueueCudaRtNonBlocking m_queue; - void CheckAndAddLayout(size_t rows, size_t cols) { + static std::pair + layoutKeyA(char trans, std::size_t m, std::size_t k) { + return (trans == 'N' || trans == 'n') ? std::make_pair(m, k) + : std::make_pair(k, m); + } + + static std::pair + layoutKeyB(char trans, std::size_t k, std::size_t n) { + return (trans == 'N' || trans == 'n') ? std::make_pair(k, n) + : std::make_pair(n, k); + } + + void checkAndAddLayout(std::size_t rows, std::size_t cols, std::size_t ld) { auto key = std::make_pair(rows, cols); - if (LayoutStore.find(key) == LayoutStore.end()) { - cublasLtMatrixLayout_t temp = nullptr; - size_t ld = rows; + if (layoutStore.find(key) == layoutStore.end()) { + cublasLtMatrixLayout_t layout = nullptr; CHECK_CUBLAS( - cublasLtMatrixLayoutCreate(&temp, CUDA_R_32F, rows, cols, ld)); - LayoutStore.emplace(key, temp); + cublasLtMatrixLayoutCreate(&layout, CUDA_R_32F, rows, cols, ld)); + layoutStore.emplace(key, layout); } } - void ResetActivation() { - epilogue = CUBLASLT_EPILOGUE_BIAS; - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute(operationDesc, - CUBLASLT_MATMUL_DESC_EPILOGUE, - &epilogue, sizeof(epilogue))); + cublasLtMatmulDesc_t &getOrCreateDesc(cublasOperation_t transA, + cublasOperation_t transB, + cublasLtEpilogue_t epilogue) { + DescKey key{(int)transA, (int)transB, (int)epilogue}; + auto it = descStore.find(key); + if (it != descStore.end()) + return it->second; + + cublasLtMatmulDesc_t desc = nullptr; + CHECK_CUBLAS( + cublasLtMatmulDescCreate(&desc, CUBLAS_COMPUTE_32F, CUDA_R_32F)); + CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( + desc, CUBLASLT_MATMUL_DESC_TRANSA, &transA, sizeof(transA))); + CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( + desc, CUBLASLT_MATMUL_DESC_TRANSB, &transB, sizeof(transB))); + CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( + desc, CUBLASLT_MATMUL_DESC_EPILOGUE, &epilogue, sizeof(epilogue))); + // For bias epilogues: set a non-null dummy pointer so the descriptor is + // valid for cublasLtMatmulAlgoGetHeuristic. + if (epilogue != CUBLASLT_EPILOGUE_DEFAULT) { + const void *dummy = d_workspace; + CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( + desc, CUBLASLT_MATMUL_DESC_BIAS_POINTER, &dummy, sizeof(dummy))); + } + descStore.emplace(key, desc); + return descStore.at(key); } - void SetReluActivation() { - epilogue = CUBLASLT_EPILOGUE_RELU_BIAS; - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute(operationDesc, - CUBLASLT_MATMUL_DESC_EPILOGUE, - &epilogue, sizeof(epilogue))); + cublasLtMatmulHeuristicResult_t & + getOrComputeAlgo(cublasOperation_t transA, cublasOperation_t transB, + cublasLtEpilogue_t epilogue, + const std::pair &kA, + const std::pair &kB, + const std::pair &kC) { + AlgoKey key{{(int)transA, (int)transB, (int)epilogue}, + kA.first, + kA.second, + kB.first, + kB.second}; + auto it = algoCache.find(key); + if (it != algoCache.end()) + return it->second; + + auto &desc = getOrCreateDesc(transA, transB, epilogue); + cublasLtMatmulHeuristicResult_t h{}; + int returnedResults = 0; + CHECK_CUBLAS(cublasLtMatmulAlgoGetHeuristic( + ltHandle, desc, layoutStore.at(kA), layoutStore.at(kB), + layoutStore.at(kC), layoutStore.at(kC), preference, 1, &h, + &returnedResults)); + if (returnedResults == 0) { + std::cerr << "[sofieBLAS] No suitable cuBLASLt algorithm found for " + << "transA=" << transA << " transB=" << transB + << " epilogue=" << epilogue << " A=[" << kA.first << "x" + << kA.second << "]" + << " B=[" << kB.first << "x" << kB.second << "]\n"; + exit(EXIT_FAILURE); + } + algoCache.emplace(key, h); + return algoCache.at(key); } - void SetGeluActivation() { - epilogue = CUBLASLT_EPILOGUE_GELU; - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute(operationDesc, - CUBLASLT_MATMUL_DESC_EPILOGUE, - &epilogue, sizeof(epilogue))); + void executeMatmul(cublasOperation_t transA, cublasOperation_t transB, + cublasLtEpilogue_t epilogue, float alpha, const float *A, + const float *B, float beta, const float *D_in, + float *C_out, const void *bias_ptr, + const std::pair &kA, + const std::pair &kB, + const std::pair &kC) { + // Retrieve (or lazily compute) the cached algorithm for this shape + auto &h = getOrComputeAlgo(transA, transB, epilogue, kA, kB, kC); + + // Retrieve the cached descriptor and patch the real bias pointer in-place + auto &desc = getOrCreateDesc(transA, transB, epilogue); + if (bias_ptr) { + CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( + desc, CUBLASLT_MATMUL_DESC_BIAS_POINTER, &bias_ptr, + sizeof(bias_ptr))); + } + + CHECK_CUBLAS(cublasLtMatmul(ltHandle, desc, &alpha, A, layoutStore.at(kA), + B, layoutStore.at(kB), &beta, D_in, + layoutStore.at(kC), C_out, layoutStore.at(kC), + &h.algo, d_workspace, workspaceSize, stream)); } }; diff --git a/include/sofieBLAS/backends/hip/sofieBLAS_hipblaslt.hpp b/include/sofieBLAS/backends/hip/sofieBLAS_hipblaslt.hpp new file mode 100644 index 0000000..0e41e70 --- /dev/null +++ b/include/sofieBLAS/backends/hip/sofieBLAS_hipblaslt.hpp @@ -0,0 +1,483 @@ +#pragma once + +#ifdef ALPAKA_ACC_GPU_HIP_ENABLED + +#include +#include +#include +#include +#include +#include + +#include "sofieBLAS/core.hpp" +#include +#include +#include + +#define CHECK_HIP(err) \ + if ((err) != hipSuccess) { \ + std::cerr << "HIP error: " << hipGetErrorString(err) << " at line " \ + << __LINE__ << "\n"; \ + exit(EXIT_FAILURE); \ + } + +#define CHECK_HIPBLAS(status) \ + do { \ + hipblasStatus_t _s = (status); \ + if (_s != HIPBLAS_STATUS_SUCCESS) { \ + std::cerr << "hipBLAS error " << _s << " at line " << __LINE__ << "\n"; \ + exit(EXIT_FAILURE); \ + } \ + } while (0) + +struct PairHash { + std::size_t + operator()(const std::pair &p) const noexcept { + std::size_t h1 = std::hash{}(p.first); + std::size_t h2 = std::hash{}(p.second); + return h1 ^ (h2 + 0x9e3779b97f4a7c15ULL + (h1 << 6) + (h1 >> 2)); + } +}; + +struct PairEq { + bool operator()(const std::pair &a, + const std::pair &b) const noexcept { + return a.first == b.first && a.second == b.second; + } +}; + +struct DescKey { + int transA; // HIPBLAS_OP_N / HIPBLAS_OP_T encoded as int + int transB; + int epilogue; // hipblasLtEpilogue_t encoded as int + bool operator==(const DescKey &o) const noexcept { + return transA == o.transA && transB == o.transB && epilogue == o.epilogue; + } +}; + +struct DescKeyHash { + std::size_t operator()(const DescKey &k) const noexcept { + std::size_t h = static_cast(k.transA) * 97u + + static_cast(k.transB) * 31u + + static_cast(k.epilogue); + return h ^ (h >> 16); + } +}; + +struct AlgoKey { + DescKey dk; + std::size_t rowsA, colsA; // physical dimensions of A in layoutStore + std::size_t rowsB, colsB; // physical dimensions of B in layoutStore + bool operator==(const AlgoKey &o) const noexcept { + return dk == o.dk && rowsA == o.rowsA && colsA == o.colsA && + rowsB == o.rowsB && colsB == o.colsB; + } +}; + +struct AlgoKeyHash { + std::size_t operator()(const AlgoKey &k) const noexcept { + std::size_t h = DescKeyHash{}(k.dk); + auto mix = [&](std::size_t v) { + h ^= std::hash{}(v) + 0x9e3779b97f4a7c15ULL + (h << 6) + + (h >> 2); + }; + mix(k.rowsA); + mix(k.colsA); + mix(k.rowsB); + mix(k.colsB); + return h; + } +}; + +class BlasHip { + hipblasLtHandle_t ltHandle = nullptr; + hipblasHandle_t handle = nullptr; + hipblasLtMatmulPreference_t preference = nullptr; + void *d_workspace = nullptr; + size_t workspaceSize = 1u << 25; // 32 MB + hipStream_t stream = nullptr; + + std::unordered_map, + hipblasLtMatrixLayout_t, PairHash, PairEq> + layoutStore; + + std::unordered_map descStore; + + std::unordered_map + algoCache; + +public: + BlasHip(const BlasHip &) = delete; + BlasHip &operator=(const BlasHip &) = delete; + BlasHip(BlasHip &&) = delete; + BlasHip &operator=(BlasHip &&) = delete; + + BlasHip(alpaka::QueueHipRtNonBlocking &queue) : m_queue{queue} { + stream = static_cast(m_queue.getNativeHandle()); + + CHECK_HIPBLAS(hipblasLtCreate(<Handle)); + + CHECK_HIPBLAS(hipblasCreate(&handle)); + CHECK_HIPBLAS(hipblasSetStream(handle, stream)); + + CHECK_HIPBLAS(hipblasLtMatmulPreferenceCreate(&preference)); + CHECK_HIP(hipMalloc(&d_workspace, workspaceSize)); + CHECK_HIPBLAS(hipblasLtMatmulPreferenceSetAttribute( + preference, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &workspaceSize, + sizeof(workspaceSize))); + } + + ~BlasHip() { + for (auto &[key, layout] : layoutStore) + if (layout) + hipblasLtMatrixLayoutDestroy(layout); + for (auto &[key, desc] : descStore) + if (desc) + hipblasLtMatmulDescDestroy(desc); + if (preference) + hipblasLtMatmulPreferenceDestroy(preference); + if (ltHandle) + hipblasLtDestroy(ltHandle); + if (handle) + hipblasDestroy(handle); + if (d_workspace) + hipFree(d_workspace); + } + + inline hipblasOperation_t charToHipBlasTranspose(char trans) { + switch (trans) { + case 'N': + case 'n': + return HIPBLAS_OP_N; + case 'T': + case 't': + return HIPBLAS_OP_T; + case 'C': + case 'c': + return HIPBLAS_OP_C; + default: + throw std::invalid_argument("Invalid transpose character for hipBLAS."); + } + } + + void addLayoutConfig(std::size_t m, std::size_t n, std::size_t k, + std::size_t lda, std::size_t ldb, std::size_t ldc, + char transa, char transb) { + if (transa == 'N' || transa == 'n') + checkAndAddLayout(m, k, lda); + else + checkAndAddLayout(k, m, lda); + if (transb == 'N' || transb == 'n') + checkAndAddLayout(k, n, ldb); + else + checkAndAddLayout(n, k, ldb); + checkAndAddLayout(m, n, ldc); + } + + template + inline void + gemm(char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, alpaka::BufHipRt, TIdx> const &A, + alpaka::BufHipRt, TIdx> const &B, float beta, + alpaka::BufHipRt, TIdx> &bias, + alpaka::BufHipRt, TIdx> &C) { + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_BIAS, + alpha, alpaka::getPtrNative(A), alpaka::getPtrNative(B), beta, + alpaka::getPtrNative(bias), alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void gemm( + char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> &bias, + alpaka::ViewPlainPtr, TIdx> &C) { + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_BIAS, + alpha, alpaka::getPtrNative(A), alpaka::getPtrNative(B), beta, + alpaka::getPtrNative(bias), alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void gemm(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *bias, T *C) { + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_BIAS, + alpha, A, B, beta, bias, C, static_cast(bias), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void gemmrelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, + alpaka::BufHipRt, TIdx> const &A, + alpaka::BufHipRt, TIdx> const &B, + float beta, + alpaka::BufHipRt, TIdx> &bias, + alpaka::BufHipRt, TIdx> &C) { + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_RELU_BIAS, + alpha, alpaka::getPtrNative(A), alpaka::getPtrNative(B), beta, + alpaka::getPtrNative(bias), alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void gemmrelu( + char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> &bias, + alpaka::ViewPlainPtr, TIdx> &C) { + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_RELU_BIAS, + alpha, alpaka::getPtrNative(A), alpaka::getPtrNative(B), beta, + alpaka::getPtrNative(bias), alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void gemmrelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *bias, T *C) { + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_RELU_BIAS, + alpha, A, B, beta, bias, C, static_cast(bias), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void gemmgelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, + alpaka::BufHipRt, TIdx> const &A, + alpaka::BufHipRt, TIdx> const &B, + float beta, + alpaka::BufHipRt, TIdx> &bias, + alpaka::BufHipRt, TIdx> &C) { + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_GELU_BIAS, + alpha, alpaka::getPtrNative(A), alpaka::getPtrNative(B), beta, + alpaka::getPtrNative(bias), alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void gemmgelu( + char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> &bias, + alpaka::ViewPlainPtr, TIdx> &C) { + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_GELU_BIAS, + alpha, alpaka::getPtrNative(A), alpaka::getPtrNative(B), beta, + alpaka::getPtrNative(bias), alpaka::getPtrNative(C), + static_cast(alpaka::getPtrNative(bias)), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void gemmgelu(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *bias, T *C) { + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_GELU_BIAS, + alpha, A, B, beta, bias, C, static_cast(bias), + layoutKeyA(transa, m, k), layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void matmul(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, + alpaka::BufHipRt, TIdx> const &A, + alpaka::BufHipRt, TIdx> const &B, + float beta, + alpaka::BufHipRt, TIdx> &C) { + float *c = alpaka::getPtrNative(C); + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_DEFAULT, + alpha, alpaka::getPtrNative(A), alpaka::getPtrNative(B), beta, + c, c, nullptr, layoutKeyA(transa, m, k), + layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void matmul( + char transa, char transb, unsigned int m, unsigned int n, unsigned int k, + float alpha, + alpaka::ViewPlainPtr, TIdx> const + &A, + alpaka::ViewPlainPtr, TIdx> const + &B, + float beta, + alpaka::ViewPlainPtr, TIdx> &C) { + T *c = alpaka::getPtrNative(C); + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_DEFAULT, + alpha, alpaka::getPtrNative(A), alpaka::getPtrNative(B), beta, + c, c, nullptr, layoutKeyA(transa, m, k), + layoutKeyB(transb, k, n), {m, n}); + } + + template + inline void matmul(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, T const *A, T const *B, + float beta, T *C) { + executeMatmul(charToHipBlasTranspose(transa), + charToHipBlasTranspose(transb), HIPBLASLT_EPILOGUE_DEFAULT, + alpha, A, B, beta, C, C, nullptr, layoutKeyA(transa, m, k), + layoutKeyB(transb, k, n), {m, n}); + } + + inline void gemmStridedBatched(char transa, char transb, int m, int n, int k, + float alpha, const float *A, int lda, + long long strideA, const float *B, int ldb, + long long strideB, float beta, float *C, + int ldc, long long strideC, int batchCount) { + CHECK_HIPBLAS(hipblasSgemmStridedBatched( + handle, charToHipBlasTranspose(transa), charToHipBlasTranspose(transb), + m, n, k, &alpha, A, lda, strideA, B, ldb, strideB, &beta, C, ldc, + strideC, batchCount)); + } + +private: + alpaka::QueueHipRtNonBlocking m_queue; + + static std::pair + layoutKeyA(char trans, std::size_t m, std::size_t k) { + return (trans == 'N' || trans == 'n') ? std::make_pair(m, k) + : std::make_pair(k, m); + } + + static std::pair + layoutKeyB(char trans, std::size_t k, std::size_t n) { + return (trans == 'N' || trans == 'n') ? std::make_pair(k, n) + : std::make_pair(n, k); + } + + void checkAndAddLayout(std::size_t rows, std::size_t cols, std::size_t ld) { + auto key = std::make_pair(rows, cols); + if (layoutStore.find(key) == layoutStore.end()) { + hipblasLtMatrixLayout_t layout = nullptr; + CHECK_HIPBLAS( + hipblasLtMatrixLayoutCreate(&layout, HIP_R_32F, rows, cols, ld)); + layoutStore.emplace(key, layout); + } + } + + hipblasLtMatmulDesc_t &getOrCreateDesc(hipblasOperation_t transA, + hipblasOperation_t transB, + hipblasLtEpilogue_t epilogue) { + DescKey key{(int)transA, (int)transB, (int)epilogue}; + auto it = descStore.find(key); + if (it != descStore.end()) + return it->second; + + hipblasLtMatmulDesc_t desc = nullptr; + CHECK_HIPBLAS( + hipblasLtMatmulDescCreate(&desc, HIPBLAS_COMPUTE_32F, HIP_R_32F)); + CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute( + desc, HIPBLASLT_MATMUL_DESC_TRANSA, &transA, sizeof(transA))); + CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute( + desc, HIPBLASLT_MATMUL_DESC_TRANSB, &transB, sizeof(transB))); + CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute( + desc, HIPBLASLT_MATMUL_DESC_EPILOGUE, &epilogue, sizeof(epilogue))); + + if (epilogue != HIPBLASLT_EPILOGUE_DEFAULT) { + const void *dummy = d_workspace; + CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute( + desc, HIPBLASLT_MATMUL_DESC_BIAS_POINTER, &dummy, sizeof(dummy))); + } + descStore.emplace(key, desc); + return descStore.at(key); + } + + hipblasLtMatmulHeuristicResult_t & + getOrComputeAlgo(hipblasOperation_t transA, hipblasOperation_t transB, + hipblasLtEpilogue_t epilogue, + const std::pair &kA, + const std::pair &kB, + const std::pair &kC) { + AlgoKey key{{(int)transA, (int)transB, (int)epilogue}, + kA.first, + kA.second, + kB.first, + kB.second}; + auto it = algoCache.find(key); + if (it != algoCache.end()) + return it->second; + + auto &desc = getOrCreateDesc(transA, transB, epilogue); + hipblasLtMatmulHeuristicResult_t h{}; + int returnedResults = 0; + CHECK_HIPBLAS(hipblasLtMatmulAlgoGetHeuristic( + ltHandle, desc, layoutStore.at(kA), layoutStore.at(kB), + layoutStore.at(kC), layoutStore.at(kC), preference, 1, &h, + &returnedResults)); + if (returnedResults == 0) { + std::cerr << "[sofieBLAS] No suitable hipBLASLt algorithm found for " + << "transA=" << transA << " transB=" << transB + << " epilogue=" << epilogue << " A=[" << kA.first << "x" + << kA.second << "]" + << " B=[" << kB.first << "x" << kB.second << "]\n"; + exit(EXIT_FAILURE); + } + algoCache.emplace(key, h); + return algoCache.at(key); + } + + void executeMatmul(hipblasOperation_t transA, hipblasOperation_t transB, + hipblasLtEpilogue_t epilogue, float alpha, const float *A, + const float *B, float beta, const float *D_in, + float *C_out, const void *bias_ptr, + const std::pair &kA, + const std::pair &kB, + const std::pair &kC) { + auto &h = getOrComputeAlgo(transA, transB, epilogue, kA, kB, kC); + + auto &desc = getOrCreateDesc(transA, transB, epilogue); + if (bias_ptr) { + CHECK_HIPBLAS(hipblasLtMatmulDescSetAttribute( + desc, HIPBLASLT_MATMUL_DESC_BIAS_POINTER, &bias_ptr, + sizeof(bias_ptr))); + } + + CHECK_HIPBLAS(hipblasLtMatmul(ltHandle, desc, &alpha, A, layoutStore.at(kA), + B, layoutStore.at(kB), &beta, D_in, + layoutStore.at(kC), C_out, layoutStore.at(kC), + &h.algo, d_workspace, workspaceSize, stream)); + } +}; + +namespace traits { + +template <> class sofieBLAS { +public: + using Impl = BlasHip; +}; + +} // namespace traits + +#endif // ALPAKA_ACC_GPU_HIP_ENABLED diff --git a/include/sofieBLAS/sofieBLAS.hpp b/include/sofieBLAS/sofieBLAS.hpp index 0d23c33..2ad40f8 100644 --- a/include/sofieBLAS/sofieBLAS.hpp +++ b/include/sofieBLAS/sofieBLAS.hpp @@ -3,9 +3,28 @@ #include "sofieBLAS/core.hpp" #if defined(ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED) -#include "sofieBLAS/backends/cpu/sofieBLAS_cpu.hpp" +// If no CPU BLAS backend was selected on the compiler command line, +// default to OpenBLAS +#if !defined(SOFIEBLAS_USE_OPENBLAS) && !defined(SOFIEBLAS_USE_MKL) && \ + !defined(SOFIEBLAS_USE_BLIS) && !defined(SOFIEBLAS_USE_ACCELERATE) +#define SOFIEBLAS_USE_OPENBLAS +#endif + +#if defined(SOFIEBLAS_USE_OPENBLAS) +#include "sofieBLAS/backends/cpu/sofieBLAS_openblas.hpp" +#elif defined(SOFIEBLAS_USE_MKL) +#include "sofieBLAS/backends/cpu/sofieBLAS_mkl.hpp" +#elif defined(SOFIEBLAS_USE_BLIS) +#include "sofieBLAS/backends/cpu/sofieBLAS_blis.hpp" +#elif defined(SOFIEBLAS_USE_ACCELERATE) +#include "sofieBLAS/backends/cpu/sofieBLAS_accelerate.hpp" +#endif #endif #if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) #include "sofieBLAS/backends/cuda/sofieBLAS_cublas.hpp" #endif + +#if defined(ALPAKA_ACC_GPU_HIP_ENABLED) +#include "sofieBLAS/backends/hip/sofieBLAS_hipblaslt.hpp" +#endif diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 2ebded2..2c55d06 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1,113 +1,62 @@ -cmake_minimum_required(VERSION 3.21) -project(AlpakaTest LANGUAGES CXX) - -# --- User-configurable options (override via -D) --- -set(CUDA_BASE "/usr/local/cuda" CACHE PATH "CUDA base path") -set(ROCM_BASE "/opt/rocm" CACHE PATH "ROCm base path") -set(ONEAPI_BASE "/opt/intel/oneapi" CACHE PATH "Intel oneAPI base path") -set(TBB_BASE "/usr" CACHE PATH "TBB base path") -set(ALPAKA_BASE "$ENV{HOME}/src/alpaka-group/alpaka" CACHE PATH "Alpaka base path") - -# --- CPU BLAS selection --- -set(AVAILABLE_BLAS_LIBS OpenBLAS MKL CACHE STRING "Choose CPU BLAS library") -set_property(CACHE AVAILABLE_BLAS_LIBS PROPERTY STRINGS OpenBLAS MKL) - -if(NOT DEFINED CPU_BLAS_LIB) - set(CPU_BLAS_LIB "OpenBLAS" CACHE STRING "CPU BLAS library to use") -endif() - -# --- Compiler flags --- set(CXXFLAGS -O2 -g -DALPAKA_HAS_STD_ATOMIC_REF) set(CXX_HOST_FLAGS -fPIC -pthread) -set(CUDA_ARCH "sm_86") -set(CXX_CUDA_FLAGS -arch=${CUDA_ARCH} -Wno-deprecated-gpu-targets --extended-lambda --expt-relaxed-constexpr) +set(CXX_CUDA_FLAGS -Wno-deprecated-gpu-targets --extended-lambda --expt-relaxed-constexpr) set(XCOMPILER_FLAGS -Xcompiler=-fPIC,-pthread) -# --- Include directories --- -include_directories(${ALPAKA_BASE}/include "../include") - -# --- Functions to find BLAS libraries --- -function(find_openblas blas_lib blas_define) - find_path(OPENBLAS_PATH NAMES libopenblas.a - PATHS /usr/lib/x86_64-linux-gnu/openblas-serial /usr/lib/x86_64-linux-gnu) - if(OPENBLAS_PATH) - find_library(OPENBLAS_LIB openblas PATHS ${OPENBLAS_PATH}) - set(${blas_lib} ${OPENBLAS_LIB} PARENT_SCOPE) - set(${blas_define} SOFIEBLAS_USE_OPENBLAS PARENT_SCOPE) - else() - set(${blas_lib} "" PARENT_SCOPE) - set(${blas_define} "" PARENT_SCOPE) - endif() -endfunction() - -function(find_mkl blas_lib blas_define) - find_path(MKL_PATH NAMES libmkl_rt.a PATHS ${ONEAPI_BASE}/mkl/latest/lib/intel64) - if(MKL_PATH) - find_library(MKL_LIB mkl_rt PATHS ${MKL_PATH}) - set(${blas_lib} ${MKL_LIB} PARENT_SCOPE) - set(${blas_define} SOFIEBLAS_USE_MKL PARENT_SCOPE) - else() - set(${blas_lib} "" PARENT_SCOPE) - set(${blas_define} "" PARENT_SCOPE) - endif() -endfunction() - -# --- BLAS detection and selection --- -if(CPU_BLAS_LIB STREQUAL "OpenBLAS") - find_openblas(BLAS_LIBS SOFIEBLAS_DEFINES) - if(NOT BLAS_LIBS) - message(WARNING "OpenBLAS not found, trying MKL...") - find_mkl(BLAS_LIBS SOFIEBLAS_DEFINES) - endif() -elseif(CPU_BLAS_LIB STREQUAL "MKL") - find_mkl(BLAS_LIBS SOFIEBLAS_DEFINES) - if(NOT BLAS_LIBS) - message(WARNING "MKL not found, trying OpenBLAS...") - find_openblas(BLAS_LIBS SOFIEBLAS_DEFINES) - endif() +if(SOFIEBLAS_CPU_BLAS_LIBS) + add_executable(test_cpu test.cc) + target_compile_features(test_cpu PUBLIC cxx_std_20) + target_compile_options(test_cpu PRIVATE ${CXXFLAGS} ${CXX_HOST_FLAGS}) + target_compile_definitions(test_cpu PRIVATE + ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED ${SOFIEBLAS_CPU_BLAS_DEFINE}) + target_include_directories(test_cpu PRIVATE ${SOFIEBLAS_CPU_BLAS_INCLUDE_DIR}) + target_link_libraries(test_cpu PRIVATE + sofieBLAS::sofieBLAS alpaka::alpaka ${SOFIEBLAS_CPU_BLAS_LIBS}) + + add_test(NAME sofieBLAS.cpu COMMAND test_cpu) else() - message(FATAL_ERROR "Unknown CPU_BLAS_LIB option: ${CPU_BLAS_LIB}") + message(STATUS "sofieBLAS tests: skipping test_cpu (no CPU BLAS library found)") endif() -if(NOT BLAS_LIBS) - message(FATAL_ERROR "No suitable CPU BLAS library found!") +if(SOFIEBLAS_CUDA_ENABLED) + configure_file(test.cc test_cuda_gen.cc COPYONLY) + add_executable(test_cuda) + target_sources(test_cuda PRIVATE ${CMAKE_CURRENT_BINARY_DIR}/test_cuda_gen.cc) + set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/test_cuda_gen.cc PROPERTIES LANGUAGE CUDA) + set_target_properties(test_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + target_compile_features(test_cuda PUBLIC cxx_std_20) + + target_compile_options(test_cuda PRIVATE + ${CXXFLAGS} + ${CXX_CUDA_FLAGS} + ${XCOMPILER_FLAGS} + ) + + target_compile_definitions(test_cuda PRIVATE ALPAKA_ACC_GPU_CUDA_ENABLED) + target_link_libraries(test_cuda PRIVATE + sofieBLAS::sofieBLAS alpaka::alpaka CUDA::cudart CUDA::cublas CUDA::cublasLt) + + add_test(NAME sofieBLAS.cuda COMMAND test_cuda) +else() + message(VERBOSE "sofieBLAS tests: skipping test_cuda (CUDA backend not enabled)") endif() -message(STATUS "Using CPU BLAS library: ${CPU_BLAS_LIB}") -message(STATUS "BLAS libraries: ${BLAS_LIBS}") -message(STATUS "Compile defines: ${SOFIEBLAS_DEFINES}") - -# --- test_cpu target --- -add_executable(test_cpu test.cc) -target_compile_features(test_cpu PUBLIC cxx_std_20) -target_compile_options(test_cpu PRIVATE ${CXXFLAGS} ${CXX_HOST_FLAGS}) -target_compile_definitions(test_cpu PRIVATE ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED ${SOFIEBLAS_DEFINES}) -target_include_directories(test_cpu PRIVATE "../sofieBLAS/include" ${ALPAKA_BASE}/include) -target_link_libraries(test_cpu PRIVATE ${BLAS_LIBS}) - - -set(TEST_SRC test.cc) +if(SOFIEBLAS_HIP_ENABLED) + configure_file(test.cc test_hip_gen.cc COPYONLY) + add_executable(test_hip) + target_sources(test_hip PRIVATE ${CMAKE_CURRENT_BINARY_DIR}/test_hip_gen.cc) + set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/test_hip_gen.cc PROPERTIES LANGUAGE HIP) + target_compile_features(test_hip PUBLIC cxx_std_20) -add_executable(test_cuda) -target_sources(test_cuda PRIVATE test.cc) -set_source_files_properties(test.cc PROPERTIES LANGUAGE CUDA) -enable_language(CUDA) -set_target_properties(test_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON) -target_compile_features(test_cuda PUBLIC cxx_std_20) + target_compile_options(test_hip PRIVATE ${CXXFLAGS} ${CXX_HOST_FLAGS}) -target_compile_options(test_cuda PRIVATE - ${CXXFLAGS} - ${CXX_CUDA_FLAGS} - ${XCOMPILER_FLAGS} -) + target_compile_definitions(test_hip PRIVATE ALPAKA_ACC_GPU_HIP_ENABLED) + target_include_directories(test_hip PRIVATE ${ROCM_BASE}/include) + target_link_directories(test_hip PRIVATE ${ROCM_BASE}/lib) + target_link_libraries(test_hip PRIVATE + sofieBLAS::sofieBLAS alpaka::alpaka hipblaslt hipblas amdhip64) -target_compile_definitions(test_cuda PRIVATE ALPAKA_ACC_GPU_CUDA_ENABLED) -target_include_directories(test_cuda PRIVATE ${ALPAKA_BASE}/include "../sofieBLAS/include" ${CUDA_BASE}/include) -target_link_directories(test_cuda PRIVATE ${CUDA_BASE}/lib64) -target_link_libraries(test_cuda PRIVATE cublas cudart) - -# --- clean target equivalent --- -add_custom_target(clean-all - COMMAND ${CMAKE_COMMAND} -E rm -f test_cpu test_cuda *.d *.o *.so - COMMENT "Cleaning all generated files" -) + add_test(NAME sofieBLAS.hip COMMAND test_hip) +else() + message(VERBOSE "sofieBLAS tests: skipping test_hip (HIP backend not enabled)") +endif() diff --git a/tests/test.cc b/tests/test.cc index 5812bc9..11d6e4c 100644 --- a/tests/test.cc +++ b/tests/test.cc @@ -1,121 +1,807 @@ #include "sofieBLAS/sofieBLAS.hpp" #include +#include +#include #include #include -#include +#include +#include -// index and size type using Idx = uint32_t; - -// dimensions -using Dim0D = alpaka::DimInt<0u>; using Dim1D = alpaka::DimInt<1u>; -using Dim2D = alpaka::DimInt<2u>; -using Dim3D = alpaka::DimInt<3u>; - -// Print a column-major matrix -template -void print(alpaka::BufCpu const &M, TIdx size) { - assert(alpaka::getExtentProduct(M) == size * size); - - for (TIdx row = 0; row < size; ++row) { - for (TIdx col = 0; col < size; ++col) { - std::cout << std::fixed << std::setprecision(2) << std::setw(7) - << M[col * size + row] << " "; + +// --------------------------------------------------------------------------- +// Reference implementations (column-major, float) +// --------------------------------------------------------------------------- +static inline float cm(const float *M, int row, int col, int ld) { + return M[col * ld + row]; +} + +// C = alpha * op(A) * op(B) + beta * C (in-place, column-major) +static void refMatmul(float *C, const float *A, const float *B, int m, int n, + int k, float alpha, float beta, bool transA, + bool transB) { + int lda = transA ? k : m; + int ldb = transB ? n : k; + for (int j = 0; j < n; ++j) { + for (int i = 0; i < m; ++i) { + float sum = 0.f; + for (int p = 0; p < k; ++p) { + float a = transA ? cm(A, p, i, lda) : cm(A, i, p, lda); + float b = transB ? cm(B, j, p, ldb) : cm(B, p, j, ldb); + sum += a * b; + } + C[j * m + i] = alpha * sum + beta * C[j * m + i]; } - std::cout << "\n"; } } -int main() { - constexpr Idx size = 4; +// C = alpha * op(A) * op(B) + beta * bias_matrix + bias_vec (per-row broadcast) +static void refGemm(float *C, const float *A, const float *B, const float *bias, + int m, int n, int k, float alpha, float beta, bool transA, + bool transB) { + int lda = transA ? k : m; + int ldb = transB ? n : k; + for (int j = 0; j < n; ++j) { + for (int i = 0; i < m; ++i) { + float sum = 0.f; + for (int p = 0; p < k; ++p) { + float a = transA ? cm(A, p, i, lda) : cm(A, i, p, lda); + float b = transB ? cm(B, j, p, ldb) : cm(B, p, j, ldb); + sum += a * b; + } + C[j * m + i] = alpha * sum + beta * bias[j * m + i] + bias[i]; + } + } +} - // Host platform and device - alpaka::PlatformCpu host_platform{}; - auto host = alpaka::getDevByIdx(host_platform, 0u); +static void refGemmRelu(float *C, const float *A, const float *B, + const float *bias, int m, int n, int k, float alpha, + float beta, bool transA, bool transB) { + refGemm(C, A, B, bias, m, n, k, alpha, beta, transA, transB); + for (int i = 0; i < m * n; ++i) + C[i] = C[i] > 0.f ? C[i] : 0.f; +} - // Allocate matrices (column-major) - auto A = alpaka::allocBuf(host, size * size); - auto B = alpaka::allocBuf(host, size * size); - auto C = alpaka::allocBuf(host, size * size); +static void refGemmGelu(float *C, const float *A, const float *B, + const float *bias, int m, int n, int k, float alpha, + float beta, bool transA, bool transB) { + refGemm(C, A, B, bias, m, n, k, alpha, beta, transA, transB); + constexpr float kInvSqrt2 = 0.7071067811865476f; + for (int i = 0; i < m * n; ++i) + C[i] *= 0.5f * (1.f + std::erff(C[i] * kInvSqrt2)); +} - // Fill A and B with random floats centered around 0 - std::random_device rd; - std::mt19937 gen(rd()); - std::normal_distribution dist(0.0f, 10.0f); +static int gFailures = 0; - for (int i = 0; i < size * size; ++i) { - A[i] = dist(gen); - B[i] = dist(gen); +static void checkClose(const float *got, const float *expected, int n, + const std::string &name, float rtol = 1e-4f, + float atol = 1e-4f) { + bool pass = true; + for (int i = 0; i < n; ++i) { + float diff = std::abs(got[i] - expected[i]); + float thr = atol + rtol * std::abs(expected[i]); + if (diff > thr) { + std::cerr << " FAIL [" << name << "] idx=" << i << " got=" << got[i] + << " expected=" << expected[i] << " diff=" << diff << "\n"; + pass = false; + } } - std::cout << "Matrix A:\n"; - print(A, size); - std::cout << '\n'; - std::cout << "Matrix B:\n"; - print(B, size); - std::cout << '\n'; + if (pass) + std::cout << " PASS " << name << "\n"; + else + ++gFailures; +} + +static void fillSeq(float *M, int n, float start = 1.f, float step = 1.f) { + for (int i = 0; i < n; ++i) + M[i] = start + static_cast(i) * step; +} + +static void fillVal(float *M, int n, float v) { + for (int i = 0; i < n; ++i) + M[i] = v; +} + +// --------------------------------------------------------------------------- +// CPU tests +// --------------------------------------------------------------------------- + +#ifdef ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED + +static void runCpuTests() { + std::cout << "\n=== CPU Tests ===\n"; + + alpaka::PlatformCpu platform{}; + auto dev = alpaka::getDevByIdx(platform, 0u); + alpaka::Queue queue{dev}; + sofieBLAS blas(queue); + + constexpr int M = 4, N = 3, K = 5; + + // Allocate host buffers + auto hA = alpaka::allocBuf(dev, static_cast(M * K)); + auto hB = alpaka::allocBuf(dev, static_cast(K * N)); + auto hC = alpaka::allocBuf(dev, static_cast(M * N)); + auto hBias = alpaka::allocBuf(dev, static_cast(M * N)); + + float *A = alpaka::getPtrNative(hA); + float *B = alpaka::getPtrNative(hB); + float *C = alpaka::getPtrNative(hC); + float *bias = alpaka::getPtrNative(hBias); + + fillSeq(A, M * K); + fillSeq(B, K * N, 1.f, 0.5f); + fillSeq(bias, M * N, 0.1f, 0.1f); + + std::vector ref(M * N); + + // --- matmul NN --- + fillVal(C, M * N, 0.f); + blas.matmul('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hC); + std::copy(C, C + M * N, ref.data()); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::matmul NN"); + + // --- matmul TN (A^T: K×M physical → M×K logical) --- + { + auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + fillVal(C, M * N, 0.f); + blas.matmul('T', 'N', M, N, K, 1.f, hAt, hB, 0.f, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), At, B, M, N, K, 1.f, 0.f, true, false); + checkClose(C, ref.data(), M * N, "cpu::matmul TN"); + } + + // --- matmul NT --- + { + auto hBt = alpaka::allocBuf(dev, static_cast(N * K)); + float *Bt = alpaka::getPtrNative(hBt); + fillSeq(Bt, N * K, 1.f, 0.5f); + fillVal(C, M * N, 0.f); + blas.matmul('N', 'T', M, N, K, 1.f, hA, hBt, 0.f, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, Bt, M, N, K, 1.f, 0.f, false, true); + checkClose(C, ref.data(), M * N, "cpu::matmul NT"); + } + + // --- matmul TT --- + { + auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); + auto hBt = alpaka::allocBuf(dev, static_cast(N * K)); + float *At = alpaka::getPtrNative(hAt); + float *Bt = alpaka::getPtrNative(hBt); + fillSeq(At, K * M); + fillSeq(Bt, N * K, 1.f, 0.5f); + fillVal(C, M * N, 0.f); + blas.matmul('T', 'T', M, N, K, 1.f, hAt, hBt, 0.f, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), At, Bt, M, N, K, 1.f, 0.f, true, true); + checkClose(C, ref.data(), M * N, "cpu::matmul TT"); + } + + // --- matmul: alpha scaling --- + fillVal(C, M * N, 0.f); + blas.matmul('N', 'N', M, N, K, 2.5f, hA, hB, 0.f, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, B, M, N, K, 2.5f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::matmul alpha=2.5"); + + // --- matmul: beta accumulation --- + fillSeq(C, M * N, 10.f); // pre-fill C + blas.matmul('N', 'N', M, N, K, 1.f, hA, hB, 0.5f, hC); + { + std::vector C0(M * N); + fillSeq(C0.data(), M * N, 10.f); + refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.5f, false, false); + std::copy(C0.begin(), C0.end(), ref.data()); + refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.5f, false, false); + } + checkClose(C, ref.data(), M * N, "cpu::matmul beta=0.5"); + + // --- gemm NN (beta=0, no prior accumulation) --- + fillVal(C, M * N, 0.f); + fillSeq(bias, M * N, 0.1f, 0.1f); + blas.gemm('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemm NN beta=0"); + + // --- gemm NN (beta=1 accumulation) --- + fillVal(C, M * N, 0.f); + blas.gemm('N', 'N', M, N, K, 1.f, hA, hB, 1.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 1.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemm NN beta=1"); + + // --- gemm TN --- + { + auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + fillVal(C, M * N, 0.f); + blas.gemm('T', 'N', M, N, K, 1.f, hAt, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); + checkClose(C, ref.data(), M * N, "cpu::gemm TN"); + } + + // --- gemmrelu: all-positive matmul result stays unchanged --- + { + // A and B with positive values ensure result is positive before bias + auto hAp = alpaka::allocBuf(dev, static_cast(M * K)); + auto hBp = alpaka::allocBuf(dev, static_cast(K * N)); + auto hBiasp = alpaka::allocBuf(dev, static_cast(M * N)); + float *Ap = alpaka::getPtrNative(hAp); + float *Bp = alpaka::getPtrNative(hBp); + float *biasp = alpaka::getPtrNative(hBiasp); + fillSeq(Ap, M * K, 0.1f, 0.1f); + fillSeq(Bp, K * N, 0.1f, 0.1f); + fillVal(biasp, M * N, 0.f); + fillVal(C, M * N, 0.f); + blas.gemmrelu('N', 'N', M, N, K, 1.f, hAp, hBp, 0.f, hBiasp, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), Ap, Bp, biasp, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemmrelu all-positive"); + } + + // --- gemmrelu: negative values clamped to zero --- + { + // Use alpha=-1 to force negative results + auto hBiasz = alpaka::allocBuf(dev, static_cast(M * N)); + fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); + fillVal(C, M * N, 0.f); + blas.gemmrelu('N', 'N', M, N, K, -1.f, hA, hB, 0.f, hBiasz, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), A, B, alpaka::getPtrNative(hBiasz), M, N, K, -1.f, + 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemmrelu alpha=-1 (clamped)"); + } + + // --- gemmrelu with bias --- + fillVal(C, M * N, 0.f); + fillSeq(bias, M * N, -5.f, 2.f); + blas.gemmrelu('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemmrelu with mixed bias"); + + // --- gemmgelu NN --- + fillVal(C, M * N, 0.f); + fillVal(bias, M * N, 0.f); + blas.gemmgelu('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemmgelu NN"); + + // --- gemmgelu with bias --- + fillVal(C, M * N, 0.f); + fillSeq(bias, M * N, -2.f, 0.5f); + blas.gemmgelu('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemmgelu with bias"); + + // --- gemmgelu TN --- + { + auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + fillVal(C, M * N, 0.f); + fillVal(bias, M * N, 0.f); + blas.gemmgelu('T', 'N', M, N, K, 1.f, hAt, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmGelu(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); + checkClose(C, ref.data(), M * N, "cpu::gemmgelu TN"); + } + + // --- edge: zero matrix --- + { + auto hZ = alpaka::allocBuf(dev, static_cast(M * K)); + fillVal(alpaka::getPtrNative(hZ), M * K, 0.f); + fillVal(C, M * N, 99.f); + fillVal(bias, M * N, 0.f); + blas.matmul('N', 'N', M, N, K, 1.f, hZ, hB, 0.f, hC); + std::fill(ref.begin(), ref.end(), 0.f); + checkClose(C, ref.data(), M * N, "cpu::matmul zero-A"); + } + + // --- edge: identity-like (square, known result) --- + { + constexpr int S = 3; + auto hI = alpaka::allocBuf(dev, static_cast(S * S)); + auto hX = alpaka::allocBuf(dev, static_cast(S * S)); + auto hY = alpaka::allocBuf(dev, static_cast(S * S)); + float *I = alpaka::getPtrNative(hI); + float *X = alpaka::getPtrNative(hX); + float *Y = alpaka::getPtrNative(hY); + fillVal(I, S * S, 0.f); + for (int i = 0; i < S; ++i) + I[i * S + i] = 1.f; + fillSeq(X, S * S); + fillVal(Y, S * S, 0.f); + blas.matmul('N', 'N', S, S, S, 1.f, hI, hX, 0.f, hY); + checkClose(Y, X, S * S, "cpu::matmul identity×X=X"); + } +} + +#endif // ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED + +#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) || defined(ALPAKA_ACC_GPU_HIP_ENABLED) +static int ldaFor(char trans, int m, int k) { + return (trans == 'N' || trans == 'n') ? m : k; +} +static int ldbFor(char trans, int k, int n) { + return (trans == 'N' || trans == 'n') ? k : n; +} +#endif + +// --------------------------------------------------------------------------- +// CUDA tests +// --------------------------------------------------------------------------- #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED + +static void runCudaTests() { + std::cout << "\n=== CUDA Tests ===\n"; + + alpaka::PlatformCudaRt platform{}; + auto dev = alpaka::getDevByIdx(platform, 0u); + alpaka::Queue queue{dev}; + sofieBLAS blas(queue); + + alpaka::PlatformCpu hostPlatform{}; + auto hostDev = alpaka::getDevByIdx(hostPlatform, 0u); + + constexpr int M = 4, N = 3, K = 5; + + auto hA = alpaka::allocBuf(hostDev, static_cast(M * K)); + auto hB = alpaka::allocBuf(hostDev, static_cast(K * N)); + auto hC = alpaka::allocBuf(hostDev, static_cast(M * N)); + auto hBias = alpaka::allocBuf(hostDev, static_cast(M * N)); + + float *A = alpaka::getPtrNative(hA); + float *B = alpaka::getPtrNative(hB); + float *bias = alpaka::getPtrNative(hBias); + + fillSeq(A, M * K); + fillSeq(B, K * N, 1.f, 0.5f); + fillVal(bias, M * N, 0.f); + + auto dA = alpaka::allocAsyncBuf(queue, static_cast(M * K)); + auto dB = alpaka::allocAsyncBuf(queue, static_cast(K * N)); + auto dC = alpaka::allocAsyncBuf(queue, static_cast(M * N)); + auto dBias = + alpaka::allocAsyncBuf(queue, static_cast(M * N)); + + alpaka::memcpy(queue, dA, hA); + alpaka::memcpy(queue, dB, hB); + alpaka::memcpy(queue, dBias, hBias); + alpaka::wait(queue); + + std::vector ref(M * N); + float *C = alpaka::getPtrNative(hC); + + auto verify = [&](const std::string &name) { + alpaka::memcpy(queue, hC, dC); + alpaka::wait(queue); + checkClose(C, ref.data(), M * N, name); + }; + + // ---- matmul NN ---- + blas.addLayoutConfig(M, N, K, ldaFor('N', M, K), ldbFor('N', K, N), M, 'N', + 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.f, false, false); + blas.matmul('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dC); + verify("cuda::matmul NN"); + + // ---- matmul TN ---- { - alpaka::PlatformCudaRt platform; - alpaka::DevCudaRt device = alpaka::getDevByIdx(platform, 0u); - alpaka::Queue queue{device}; + auto hAt = alpaka::allocBuf(hostDev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + auto dAt = + alpaka::allocAsyncBuf(queue, static_cast(K * M)); + alpaka::memcpy(queue, dAt, hAt); + alpaka::wait(queue); + blas.addLayoutConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, 'T', + 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), At, B, M, N, K, 1.f, 0.f, true, false); + blas.matmul('T', 'N', M, N, K, 1.f, dAt, dB, 0.f, dC); + verify("cuda::matmul TN"); + } - const Idx m = size; // rows of A and C - const Idx n = size; // columns of B and C - const Idx k = size; // columns of A and rows of B + // ---- matmul NT ---- + { + auto hBt = alpaka::allocBuf(hostDev, static_cast(N * K)); + float *Bt = alpaka::getPtrNative(hBt); + fillSeq(Bt, N * K, 1.f, 0.5f); + auto dBt = + alpaka::allocAsyncBuf(queue, static_cast(N * K)); + alpaka::memcpy(queue, dBt, hBt); + alpaka::wait(queue); + blas.addLayoutConfig(M, N, K, ldaFor('N', M, K), ldbFor('T', K, N), M, 'N', + 'T'); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, Bt, M, N, K, 1.f, 0.f, false, true); + blas.matmul('N', 'T', M, N, K, 1.f, dA, dBt, 0.f, dC); + verify("cuda::matmul NT"); + } - const float alpha = 1.0f; - const float beta = 0.0f; + // ---- matmul alpha=2.5 ---- + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, B, M, N, K, 2.5f, 0.f, false, false); + blas.matmul('N', 'N', M, N, K, 2.5f, dA, dB, 0.f, dC); + verify("cuda::matmul alpha=2.5"); - const Idx lda = size; // leading dimension of A - const Idx ldb = size; // leading dimension of B - const Idx ldc = size; + // ---- gemm NN beta=0 ---- + fillSeq(bias, M * N, 0.1f, 0.1f); + alpaka::memcpy(queue, dBias, hBias); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + blas.gemm('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); + verify("cuda::gemm NN beta=0"); - auto A_d = alpaka::allocAsyncBuf(queue, size * size); - auto B_d = alpaka::allocAsyncBuf(queue, size * size); - auto C_d = alpaka::allocAsyncBuf(queue, size * size); - alpaka::memcpy(queue, A_d, A); - alpaka::memcpy(queue, B_d, B); + // ---- gemm NN beta=1 ---- + // D_in = bias, so result = A*B + 1*bias_matrix + bias_vec + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 1.f, false, false); + blas.gemm('N', 'N', M, N, K, 1.f, dA, dB, 1.f, dBias, dC); + verify("cuda::gemm NN beta=1"); - sofieBLAS blas(queue); - blas.gemm('n', 'n', m, n, k, alpha, A_d, lda, B_d, ldb, beta, C_d, ldc); - alpaka::memcpy(queue, C, C_d); + // ---- gemm TN ---- + { + auto hAt = alpaka::allocBuf(hostDev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + auto dAt = + alpaka::allocAsyncBuf(queue, static_cast(K * M)); + alpaka::memcpy(queue, dAt, hAt); + alpaka::wait(queue); + blas.addLayoutConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, 'T', + 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); + blas.gemm('T', 'N', M, N, K, 1.f, dAt, dB, 0.f, dBias, dC); + verify("cuda::gemm TN"); + } + // ---- gemmrelu: all-positive (relu is identity) ---- + { + auto hAp = alpaka::allocBuf(hostDev, static_cast(M * K)); + auto hBp = alpaka::allocBuf(hostDev, static_cast(K * N)); + auto hBiasz = + alpaka::allocBuf(hostDev, static_cast(M * N)); + float *Ap = alpaka::getPtrNative(hAp); + float *Bp = alpaka::getPtrNative(hBp); + fillSeq(Ap, M * K, 0.1f, 0.1f); + fillSeq(Bp, K * N, 0.1f, 0.1f); + fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); + auto dAp = + alpaka::allocAsyncBuf(queue, static_cast(M * K)); + auto dBp = + alpaka::allocAsyncBuf(queue, static_cast(K * N)); + auto dBiasz = + alpaka::allocAsyncBuf(queue, static_cast(M * N)); + alpaka::memcpy(queue, dAp, hAp); + alpaka::memcpy(queue, dBp, hBp); + alpaka::memcpy(queue, dBiasz, hBiasz); alpaka::wait(queue); - std::cout << "CUDA Matrix C = A × B:\n"; - print(C, size); - std::cout << '\n'; + blas.addLayoutConfig(M, N, K, M, K, M, 'N', 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), Ap, Bp, alpaka::getPtrNative(hBiasz), M, N, K, 1.f, + 0.f, false, false); + blas.gemmrelu('N', 'N', M, N, K, 1.f, dAp, dBp, 0.f, dBiasz, dC); + verify("cuda::gemmrelu all-positive"); } -#endif -#ifdef ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED + // ---- gemmrelu: alpha=-1 forces negatives -> clamped to zero ---- { - alpaka::PlatformCpu platform; - alpaka::DevCpu device = alpaka::getDevByIdx(platform, 0u); - alpaka::Queue queue{device}; + auto hBiasz = + alpaka::allocBuf(hostDev, static_cast(M * N)); + fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); + auto dBiasz = + alpaka::allocAsyncBuf(queue, static_cast(M * N)); + alpaka::memcpy(queue, dBiasz, hBiasz); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), A, B, alpaka::getPtrNative(hBiasz), M, N, K, -1.f, + 0.f, false, false); + blas.gemmrelu('N', 'N', M, N, K, -1.f, dA, dB, 0.f, dBiasz, dC); + verify("cuda::gemmrelu alpha=-1 (clamped)"); + } - const Idx m = size; // rows of A and C - const Idx n = size; // columns of B and C - const Idx k = size; // columns of A and rows of B + // ---- gemmrelu with mixed bias ---- + fillSeq(bias, M * N, -5.f, 2.f); + alpaka::memcpy(queue, dBias, hBias); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + blas.gemmrelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); + verify("cuda::gemmrelu with mixed bias"); - const float alpha = 1.0f; - const float beta = 0.0f; + // ---- gemmgelu NN ---- + fillVal(bias, M * N, 0.f); + alpaka::memcpy(queue, dBias, hBias); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + blas.gemmgelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); + verify("cuda::gemmgelu NN"); - const Idx lda = size; // leading dimension of A - const Idx ldb = size; // leading dimension of B - const Idx ldc = size; // leading dimension of C + // ---- gemmgelu with bias ---- + fillSeq(bias, M * N, -2.f, 0.5f); + alpaka::memcpy(queue, dBias, hBias); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + blas.gemmgelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); + verify("cuda::gemmgelu with bias"); - sofieBLAS blas(queue); - blas.gemm('n', 'n', m, n, k, alpha, A, lda, B, ldb, beta, C, ldc); + // ---- edge: zero A ---- + { + auto hZero = alpaka::allocBuf(hostDev, static_cast(M * K)); + fillVal(alpaka::getPtrNative(hZero), M * K, 0.f); + auto dZero = + alpaka::allocAsyncBuf(queue, static_cast(M * K)); + alpaka::memcpy(queue, dZero, hZero); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + blas.matmul('N', 'N', M, N, K, 1.f, dZero, dB, 0.f, dC); + verify("cuda::matmul zero-A"); + } +} + +#endif // ALPAKA_ACC_GPU_CUDA_ENABLED + +// --------------------------------------------------------------------------- +// HIP tests +// --------------------------------------------------------------------------- + +#ifdef ALPAKA_ACC_GPU_HIP_ENABLED + +static void runHipTests() { + std::cout << "\n=== HIP Tests ===\n"; + + alpaka::PlatformHipRt platform{}; + auto dev = alpaka::getDevByIdx(platform, 0u); + alpaka::Queue queue{dev}; + sofieBLAS blas(queue); + + alpaka::PlatformCpu hostPlatform{}; + auto hostDev = alpaka::getDevByIdx(hostPlatform, 0u); + constexpr int M = 4, N = 3, K = 5; + + auto hA = alpaka::allocBuf(hostDev, static_cast(M * K)); + auto hB = alpaka::allocBuf(hostDev, static_cast(K * N)); + auto hC = alpaka::allocBuf(hostDev, static_cast(M * N)); + auto hBias = alpaka::allocBuf(hostDev, static_cast(M * N)); + + float *A = alpaka::getPtrNative(hA); + float *B = alpaka::getPtrNative(hB); + float *bias = alpaka::getPtrNative(hBias); + + fillSeq(A, M * K); + fillSeq(B, K * N, 1.f, 0.5f); + fillVal(bias, M * N, 0.f); + + auto dA = alpaka::allocAsyncBuf(queue, static_cast(M * K)); + auto dB = alpaka::allocAsyncBuf(queue, static_cast(K * N)); + auto dC = alpaka::allocAsyncBuf(queue, static_cast(M * N)); + auto dBias = + alpaka::allocAsyncBuf(queue, static_cast(M * N)); + + alpaka::memcpy(queue, dA, hA); + alpaka::memcpy(queue, dB, hB); + alpaka::memcpy(queue, dBias, hBias); + alpaka::wait(queue); + + std::vector ref(M * N); + float *C = alpaka::getPtrNative(hC); + + auto verify = [&](const std::string &name) { + alpaka::memcpy(queue, hC, dC); alpaka::wait(queue); - std::cout << "CPU Matrix C = A × B:\n"; - print(C, size); - std::cout << '\n'; + checkClose(C, ref.data(), M * N, name); + }; + + // ---- matmul NN ---- + blas.addLayoutConfig(M, N, K, ldaFor('N', M, K), ldbFor('N', K, N), M, 'N', + 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.f, false, false); + blas.matmul('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dC); + verify("hip::matmul NN"); + + // ---- matmul TN ---- + { + auto hAt = alpaka::allocBuf(hostDev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + auto dAt = + alpaka::allocAsyncBuf(queue, static_cast(K * M)); + alpaka::memcpy(queue, dAt, hAt); + alpaka::wait(queue); + blas.addLayoutConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, 'T', + 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), At, B, M, N, K, 1.f, 0.f, true, false); + blas.matmul('T', 'N', M, N, K, 1.f, dAt, dB, 0.f, dC); + verify("hip::matmul TN"); + } + + // ---- matmul NT ---- + { + auto hBt = alpaka::allocBuf(hostDev, static_cast(N * K)); + float *Bt = alpaka::getPtrNative(hBt); + fillSeq(Bt, N * K, 1.f, 0.5f); + auto dBt = + alpaka::allocAsyncBuf(queue, static_cast(N * K)); + alpaka::memcpy(queue, dBt, hBt); + alpaka::wait(queue); + blas.addLayoutConfig(M, N, K, ldaFor('N', M, K), ldbFor('T', K, N), M, 'N', + 'T'); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, Bt, M, N, K, 1.f, 0.f, false, true); + blas.matmul('N', 'T', M, N, K, 1.f, dA, dBt, 0.f, dC); + verify("hip::matmul NT"); + } + + // ---- matmul alpha=2.5 ---- + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, B, M, N, K, 2.5f, 0.f, false, false); + blas.matmul('N', 'N', M, N, K, 2.5f, dA, dB, 0.f, dC); + verify("hip::matmul alpha=2.5"); + + // ---- gemm NN beta=0 ---- + fillSeq(bias, M * N, 0.1f, 0.1f); + alpaka::memcpy(queue, dBias, hBias); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + blas.gemm('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); + verify("hip::gemm NN beta=0"); + + // ---- gemm NN beta=1 ---- + // D_in = bias, so result = A*B + 1*bias_matrix + bias_vec + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 1.f, false, false); + blas.gemm('N', 'N', M, N, K, 1.f, dA, dB, 1.f, dBias, dC); + verify("hip::gemm NN beta=1"); + + // ---- gemm TN ---- + { + auto hAt = alpaka::allocBuf(hostDev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + auto dAt = + alpaka::allocAsyncBuf(queue, static_cast(K * M)); + alpaka::memcpy(queue, dAt, hAt); + alpaka::wait(queue); + blas.addLayoutConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, 'T', + 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); + blas.gemm('T', 'N', M, N, K, 1.f, dAt, dB, 0.f, dBias, dC); + verify("hip::gemm TN"); + } + + // ---- gemmrelu: all-positive ---- + { + auto hAp = alpaka::allocBuf(hostDev, static_cast(M * K)); + auto hBp = alpaka::allocBuf(hostDev, static_cast(K * N)); + auto hBiasz = + alpaka::allocBuf(hostDev, static_cast(M * N)); + float *Ap = alpaka::getPtrNative(hAp); + float *Bp = alpaka::getPtrNative(hBp); + fillSeq(Ap, M * K, 0.1f, 0.1f); + fillSeq(Bp, K * N, 0.1f, 0.1f); + fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); + auto dAp = + alpaka::allocAsyncBuf(queue, static_cast(M * K)); + auto dBp = + alpaka::allocAsyncBuf(queue, static_cast(K * N)); + auto dBiasz = + alpaka::allocAsyncBuf(queue, static_cast(M * N)); + alpaka::memcpy(queue, dAp, hAp); + alpaka::memcpy(queue, dBp, hBp); + alpaka::memcpy(queue, dBiasz, hBiasz); + alpaka::wait(queue); + blas.addLayoutConfig(M, N, K, M, K, M, 'N', 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), Ap, Bp, alpaka::getPtrNative(hBiasz), M, N, K, 1.f, + 0.f, false, false); + blas.gemmrelu('N', 'N', M, N, K, 1.f, dAp, dBp, 0.f, dBiasz, dC); + verify("hip::gemmrelu all-positive"); + } + + // ---- gemmrelu: alpha=-1 forces negatives ---- + { + auto hBiasz = + alpaka::allocBuf(hostDev, static_cast(M * N)); + fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); + auto dBiasz = + alpaka::allocAsyncBuf(queue, static_cast(M * N)); + alpaka::memcpy(queue, dBiasz, hBiasz); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), A, B, alpaka::getPtrNative(hBiasz), M, N, K, -1.f, + 0.f, false, false); + blas.gemmrelu('N', 'N', M, N, K, -1.f, dA, dB, 0.f, dBiasz, dC); + verify("hip::gemmrelu alpha=-1 (clamped)"); + } + + // ---- gemmrelu with mixed bias ---- + fillSeq(bias, M * N, -5.f, 2.f); + alpaka::memcpy(queue, dBias, hBias); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + blas.gemmrelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); + verify("hip::gemmrelu with mixed bias"); + + // ---- gemmgelu NN ---- + fillVal(bias, M * N, 0.f); + alpaka::memcpy(queue, dBias, hBias); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + blas.gemmgelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); + verify("hip::gemmgelu NN"); + + // ---- gemmgelu with bias ---- + fillSeq(bias, M * N, -2.f, 0.5f); + alpaka::memcpy(queue, dBias, hBias); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + blas.gemmgelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); + verify("hip::gemmgelu with bias"); + + // ---- edge: zero A ---- + { + auto hZero = alpaka::allocBuf(hostDev, static_cast(M * K)); + fillVal(alpaka::getPtrNative(hZero), M * K, 0.f); + auto dZero = + alpaka::allocAsyncBuf(queue, static_cast(M * K)); + alpaka::memcpy(queue, dZero, hZero); + alpaka::wait(queue); + std::fill(ref.begin(), ref.end(), 0.f); + blas.matmul('N', 'N', M, N, K, 1.f, dZero, dB, 0.f, dC); + verify("hip::matmul zero-A"); } +} + +#endif // ALPAKA_ACC_GPU_HIP_ENABLED + +// --------------------------------------------------------------------------- +// main +// --------------------------------------------------------------------------- + +int main() { +#ifdef ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED + runCpuTests(); #endif +#ifdef ALPAKA_ACC_GPU_CUDA_ENABLED + runCudaTests(); +#endif +#ifdef ALPAKA_ACC_GPU_HIP_ENABLED + runHipTests(); +#endif + + std::cout << "\n"; + if (gFailures == 0) + std::cout << "All tests passed.\n"; + else + std::cout << gFailures << " test(s) FAILED.\n"; - return 0; + return gFailures > 0 ? EXIT_FAILURE : EXIT_SUCCESS; }