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..bf3fa5d 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,57 @@ 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` for batched gemm operations through strides +- `addOperationConfig` that creates the matrix layouts and resolves the multiply algorithm for a call site's shape ahead of its first call (see below). + +## GEMM call instantiation and the algorithm cache + +A GEMM call computes `C = alpha * op(A) * op(B) + beta * C`, where A and B are the input matrices, C the output, and `op` an optional transpose. To run one, cuBLASLt and hipBLASLt need three kinds of objects besides the data: + +- a **matrix layout** per matrix: a descriptor holding its rows, columns and leading dimension; +- a **matmul descriptor**: the operation settings (the transposes and the epilogue); +- an **algorithm**: the concrete GEMM kernel the library selects for the given settings and dimensions, obtained by querying its heuristic (`cublasLtMatmulAlgoGetHeuristic` / `hipblasLtMatmulAlgoGetHeuristic`). The query runs on the host and is not free. + +The CUDA backend (`BlasCuda`, over cuBLASLt) and the HIP backend (`BlasHip`, over hipBLASLt) behave identically: all three objects are created the first time a combination appears and cached, keyed by the exact dimensions plus, for descriptors and algorithms, the transposes and the epilogue. One instance therefore serves GEMM calls at sizes that vary at runtime: a size seen for the first time creates and caches its objects, and a repeated size reuses them without another heuristic query. + +### addOperationConfig + +`addOperationConfig(m, n, k, lda, ldb, ldc, transa, transb, epilogue)` creates all three objects for one operation (the matrix layouts, the matmul descriptor and the algorithm) for the given dimensions, transposes and epilogue, before the corresponding call is made. It is optional: a combination that was never configured is created and cached on its first call. The `epilogue` argument is the `Epilogue` enum from `sofieBLAS/core.hpp` and names which call the site will make, because the fused epilogue is part of the selected kernel: + +| `Epilogue` value | call it configures | +| --- | --- | +| `Epilogue::Default` | `matmul` (no bias) | +| `Epilogue::Bias` | `gemm` (adds the bias vector) | +| `Epilogue::ReluBias` | `gemmrelu` (bias, then ReLU) | +| `Epilogue::GeluBias` | `gemmgelu` (bias, then GELU) | + +### Initializing the cache limit + +The algorithm cache is unbounded by default. Passing a limit as the second constructor argument caps the number of cached algorithms; when an insertion would exceed the limit, the least recently used entries are evicted. Choose a limit at least as large as the number of distinct shapes the workload uses regularly, or leave it unbounded. `algoCacheSize()` returns the current number of entries. + +```cpp +sofieBLAS blas(queue); // unbounded algorithm cache (default) +sofieBLAS capped(queue, 32); // at most 32 entries, LRU eviction + +blas.addOperationConfig(64, 3, 5, 64, 5, 64, 'N', 'N', Epilogue::Default); +blas.matmul('N', 'N', 64, 3, 5, 1.0f, dA, dB, 0.0f, dC); // created by addOperationConfig: cache hit +blas.matmul('N', 'N', 37, 3, 5, 1.0f, dA, dB, 0.0f, dC); // new size: created on first use +blas.gemmrelu('N', 'N', 64, 3, 5, 1.0f, dA, dB, 0.0f, dBias, dC); // same size, other epilogue: layouts reused, descriptor and algorithm created on first use +``` + ## 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..687c312 --- /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.addOperationConfig(M, N, K, M, K, M, 'N', 'N', Epilogue::Default); + + 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.addOperationConfig(M, N, K, M, K, M, 'N', 'N', Epilogue::Default); + + 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..5b3b92a 100644 --- a/include/sofieBLAS/backends/cuda/sofieBLAS_cublas.hpp +++ b/include/sofieBLAS/backends/cuda/sofieBLAS_cublas.hpp @@ -5,7 +5,9 @@ #include #include #include +#include #include +#include #include #include @@ -15,7 +17,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,328 +25,74 @@ #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) -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)); +// The cuBLASLt forwarding of the shared BlasLt implementation in +// backends/gpu/detail +struct CublasLtApi { + using Queue = alpaka::QueueCudaRtNonBlocking; + using Handle = cublasLtHandle_t; + using BlasHandle = cublasHandle_t; + using Preference = cublasLtMatmulPreference_t; + using Stream = cudaStream_t; + using Layout = cublasLtMatrixLayout_t; + using MatmulDesc = cublasLtMatmulDesc_t; + using HeuristicResult = cublasLtMatmulHeuristicResult_t; + using Operation = cublasOperation_t; + using Epilogue = cublasLtEpilogue_t; + + static constexpr auto OpN = CUBLAS_OP_N; + static constexpr auto OpT = CUBLAS_OP_T; + static constexpr auto OpC = CUBLAS_OP_C; + static constexpr auto EpilogueDefault = CUBLASLT_EPILOGUE_DEFAULT; + static constexpr auto EpilogueBias = CUBLASLT_EPILOGUE_BIAS; + static constexpr auto EpilogueReluBias = CUBLASLT_EPILOGUE_RELU_BIAS; + static constexpr auto EpilogueGeluBias = CUBLASLT_EPILOGUE_GELU_BIAS; + static constexpr auto ComputeF32 = CUBLAS_COMPUTE_32F; + static constexpr auto RealF32 = CUDA_R_32F; + static constexpr auto DescTransA = CUBLASLT_MATMUL_DESC_TRANSA; + static constexpr auto DescTransB = CUBLASLT_MATMUL_DESC_TRANSB; + static constexpr auto DescEpilogue = CUBLASLT_MATMUL_DESC_EPILOGUE; + static constexpr auto DescBiasPointer = CUBLASLT_MATMUL_DESC_BIAS_POINTER; + static constexpr auto PrefMaxWorkspace = + CUBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES; + static constexpr const char *name = "cuBLASLt"; + + static constexpr auto ltCreate = cublasLtCreate; + static constexpr auto ltDestroy = cublasLtDestroy; + static constexpr auto blasCreate = cublasCreate; + static constexpr auto blasDestroy = cublasDestroy; + static constexpr auto blasSetStream = cublasSetStream; + static constexpr auto prefCreate = cublasLtMatmulPreferenceCreate; + static constexpr auto prefDestroy = cublasLtMatmulPreferenceDestroy; + static constexpr auto prefSetAttribute = cublasLtMatmulPreferenceSetAttribute; + static constexpr auto layoutCreate = cublasLtMatrixLayoutCreate; + static constexpr auto layoutDestroy = cublasLtMatrixLayoutDestroy; + static constexpr auto descCreate = cublasLtMatmulDescCreate; + static constexpr auto descDestroy = cublasLtMatmulDescDestroy; + static constexpr auto descSetAttribute = cublasLtMatmulDescSetAttribute; + static constexpr auto getHeuristic = cublasLtMatmulAlgoGetHeuristic; + static constexpr auto matmul = cublasLtMatmul; + static constexpr auto sgemmStridedBatched = cublasSgemmStridedBatched; + + // cudaMalloc has a templated C++ overload, so a pointer to it is ambiguous + static cudaError_t rtMalloc(void **ptr, std::size_t size) { + return cudaMalloc(ptr, size); } + static cudaError_t rtFree(void *ptr) { return cudaFree(ptr); } }; -struct PairEq { - bool operator()(const std::pair &a, - const std::pair &b) const noexcept { - return a.first == b.first && a.second == b.second; - } -}; - -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 - 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; - -public: - 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(cublasLtMatmulPreferenceCreate(&preference)); - CHECK_CUDA(cudaMalloc(&d_workspace, workspaceSize)); - CHECK_CUBLAS(cublasLtMatmulPreferenceSetAttribute( - preference, CUBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &workspaceSize, - sizeof(workspaceSize))); - } - - ~BlasCuda() { - for (auto& [key, layout] : LayoutStore) { - if (layout) { - cublasLtMatrixLayoutDestroy(layout); - } - } - LayoutStore.clear(); - - if (preference) - cublasLtMatmulPreferenceDestroy(preference); - if (operationDesc) - cublasLtMatmulDescDestroy(operationDesc); - if (ltHandle) - cublasLtDestroy(ltHandle); - if (d_workspace) - cudaFree(d_workspace); - - } - - inline cublasOperation_t charToCuBlasTranspose(char trans) { - switch (trans) { - case 'N': - case 'n': - return CUBLAS_OP_N; - case 'T': - case 't': - return CUBLAS_OP_T; - case 'C': - case 'c': - return CUBLAS_OP_C; - default: - throw std::invalid_argument("Invalid transpose character for cuBLAS."); - } - } - - void AddLayoutConfig(std::size_t m, std::size_t n, std::size_t k) { - CheckAndAddLayout(k, m); - CheckAndAddLayout(k, n); - CheckAndAddLayout(m, n); - } - -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)); -} +#define SOFIEBLAS_CHECK_LT(status) CHECK_CUBLAS(status) +#define SOFIEBLAS_CHECK_RT(err) CHECK_CUDA(err) -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)); +#include "../gpu/detail/sofieBLAS_blaslt_common.tpp" - 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_RELU_BIAS; - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_EPILOGUE, &ep, sizeof(ep))); - - 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); - } - - 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 gemmgelu(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)); - - void *bias_ptr = reinterpret_cast(alpaka::getPtrNative(bias)); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_BIAS_POINTER, &bias_ptr, - sizeof(bias_ptr))); - - cublasOperation_t transB = charToCuBlasTranspose(transb); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_TRANSB, &transB, sizeof(transB))); - - cublasOperation_t transA = charToCuBlasTranspose(transa); - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute( - localDesc, CUBLASLT_MATMUL_DESC_TRANSA, &transA, sizeof(transA))); - SetGeluActivation(); - - 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); - } - - 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)); - } - -private: - alpaka::QueueCudaRtNonBlocking m_queue; - - void CheckAndAddLayout(size_t rows, size_t cols) { - auto key = std::make_pair(rows, cols); - if (LayoutStore.find(key) == LayoutStore.end()) { - cublasLtMatrixLayout_t temp = nullptr; - size_t ld = rows; - CHECK_CUBLAS( - cublasLtMatrixLayoutCreate(&temp, CUDA_R_32F, rows, cols, ld)); - LayoutStore.emplace(key, temp); - } - } - - void ResetActivation() { - epilogue = CUBLASLT_EPILOGUE_BIAS; - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute(operationDesc, - CUBLASLT_MATMUL_DESC_EPILOGUE, - &epilogue, sizeof(epilogue))); - } - - void SetReluActivation() { - epilogue = CUBLASLT_EPILOGUE_RELU_BIAS; - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute(operationDesc, - CUBLASLT_MATMUL_DESC_EPILOGUE, - &epilogue, sizeof(epilogue))); - } - - void SetGeluActivation() { - epilogue = CUBLASLT_EPILOGUE_GELU; - CHECK_CUBLAS(cublasLtMatmulDescSetAttribute(operationDesc, - CUBLASLT_MATMUL_DESC_EPILOGUE, - &epilogue, sizeof(epilogue))); - } -}; +using BlasCuda = BlasLt; namespace traits { diff --git a/include/sofieBLAS/backends/gpu/detail/sofieBLAS_blaslt_common.tpp b/include/sofieBLAS/backends/gpu/detail/sofieBLAS_blaslt_common.tpp new file mode 100644 index 0000000..b688b76 --- /dev/null +++ b/include/sofieBLAS/backends/gpu/detail/sofieBLAS_blaslt_common.tpp @@ -0,0 +1,418 @@ +// Shared implementation of the cuBLASLt and hipBLASLt backends. The two +// vendor APIs have the same shape under different names, so the backend is +// written once against an Api table. A vendor header defines that table +// (the types, constants and functions of its library), defines the check +// macros SOFIEBLAS_CHECK_LT and SOFIEBLAS_CHECK_RT, includes the vendor and +// standard headers (, , , , +// , , , , alpaka), and then +// includes this file. + +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; // backend transpose enum encoded as int + int transB; + int epilogue; // backend epilogue enum 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; + } +}; + +template class BlasLt { + typename Api::Handle ltHandle = nullptr; + typename Api::BlasHandle handle = nullptr; + typename Api::Preference preference = nullptr; + void *d_workspace = nullptr; + size_t workspaceSize = 1u << 25; // 32 MB + typename Api::Stream stream = nullptr; + + std::unordered_map, typename Api::Layout, + PairHash, PairEq> + layoutStore; + + std::unordered_map descStore; + + // One cache entry per exact GEMM configuration: the heuristic result to + // reuse, plus this entry's position in the recency list so a hit can mark + // itself most-recently-used in O(1). The position is only maintained when a + // cache limit is set; with no limit the list stays empty. + struct CacheEntry { + typename Api::HeuristicResult h{}; + std::list::iterator lru{}; + }; + std::unordered_map algoCache; + // entries ordered most- to least-recently used; drives eviction + std::list lruOrder; + // 0 = unbounded + std::size_t algoCacheLimit = 0; + +public: + std::size_t algoCacheSize() const { return algoCache.size(); } + + BlasLt(const BlasLt &) = delete; + BlasLt &operator=(const BlasLt &) = delete; + BlasLt(BlasLt &&) = delete; + BlasLt &operator=(BlasLt &&) = delete; + + BlasLt(typename Api::Queue &queue, std::size_t cacheLimit = 0) + : algoCacheLimit{cacheLimit}, m_queue{queue} { + stream = static_cast(m_queue.getNativeHandle()); + + SOFIEBLAS_CHECK_LT(Api::ltCreate(<Handle)); + + SOFIEBLAS_CHECK_LT(Api::blasCreate(&handle)); + SOFIEBLAS_CHECK_LT(Api::blasSetStream(handle, stream)); + + SOFIEBLAS_CHECK_LT(Api::prefCreate(&preference)); + SOFIEBLAS_CHECK_RT(Api::rtMalloc(&d_workspace, workspaceSize)); + SOFIEBLAS_CHECK_LT(Api::prefSetAttribute(preference, Api::PrefMaxWorkspace, + &workspaceSize, + sizeof(workspaceSize))); + } + + ~BlasLt() { + for (auto &[key, layout] : layoutStore) + if (layout) + Api::layoutDestroy(layout); + for (auto &[key, desc] : descStore) + if (desc) + Api::descDestroy(desc); + if (preference) + Api::prefDestroy(preference); + if (ltHandle) + Api::ltDestroy(ltHandle); + if (handle) + Api::blasDestroy(handle); + if (d_workspace) + Api::rtFree(d_workspace); + } + + inline typename Api::Operation charToTranspose(char trans) { + switch (trans) { + case 'N': + case 'n': + return Api::OpN; + case 'T': + case 't': + return Api::OpT; + case 'C': + case 'c': + return Api::OpC; + default: + throw std::invalid_argument( + std::string("Invalid transpose character for ") + Api::name + "."); + } + } + + // Registers a call site's construction-time shape: creates the three matrix + // layouts and resolves the multiply algorithm for them up front, so the + // first call at this shape finds everything cached. + void addOperationConfig(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, Epilogue epilogue) { + const auto shapeA = layoutKeyA(transa, m, k); + const auto shapeB = layoutKeyB(transb, k, n); + const std::pair shapeC{m, n}; + getOrCreateLayout(shapeA, lda); + getOrCreateLayout(shapeB, ldb); + getOrCreateLayout(shapeC, ldc); + + typename Api::Epilogue apiEpilogue = Api::EpilogueDefault; + switch (epilogue) { + case Epilogue::Bias: + apiEpilogue = Api::EpilogueBias; + break; + case Epilogue::ReluBias: + apiEpilogue = Api::EpilogueReluBias; + break; + case Epilogue::GeluBias: + apiEpilogue = Api::EpilogueGeluBias; + break; + case Epilogue::Default: + break; + } + getOrComputeAlgo(charToTranspose(transa), charToTranspose(transb), + apiEpilogue, shapeA, shapeB, shapeC); + } + + // Each multiply variant comes as one generic overload, where A, B, bias and + // C are any alpaka buffers or views (anything alpaka::getPtrNative + // accepts), and one raw device-pointer overload, which generated code + // calls. + template + inline void gemm(char transa, char transb, unsigned int m, unsigned int n, + unsigned int k, float alpha, TA const &A, TB const &B, + float beta, TBias &bias, TC &C) { + executeMatmul(charToTranspose(transa), charToTranspose(transb), + Api::EpilogueBias, 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(charToTranspose(transa), charToTranspose(transb), + Api::EpilogueBias, 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, TA const &A, TB const &B, + float beta, TBias &bias, TC &C) { + executeMatmul(charToTranspose(transa), charToTranspose(transb), + Api::EpilogueReluBias, 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(charToTranspose(transa), charToTranspose(transb), + Api::EpilogueReluBias, 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, TA const &A, TB const &B, + float beta, TBias &bias, TC &C) { + executeMatmul(charToTranspose(transa), charToTranspose(transb), + Api::EpilogueGeluBias, 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(charToTranspose(transa), charToTranspose(transb), + Api::EpilogueGeluBias, 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, TA const &A, TB const &B, + float beta, TC &C) { + auto *c = alpaka::getPtrNative(C); + executeMatmul(charToTranspose(transa), charToTranspose(transb), + Api::EpilogueDefault, 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(charToTranspose(transa), charToTranspose(transb), + Api::EpilogueDefault, 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) { + SOFIEBLAS_CHECK_LT(Api::sgemmStridedBatched( + handle, charToTranspose(transa), charToTranspose(transb), m, n, k, + &alpha, A, lda, strideA, B, ldb, strideB, &beta, C, ldc, strideC, + batchCount)); + } + +private: + typename Api::Queue 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); + } + + // Returns the layout describing a (rows, cols) matrix, creating and caching + // it on first use. Every caller passes ld = rows (dense column-major). + typename Api::Layout + getOrCreateLayout(const std::pair &shape, + std::size_t ld) { + auto it = layoutStore.find(shape); + if (it != layoutStore.end()) + return it->second; + typename Api::Layout layout = nullptr; + SOFIEBLAS_CHECK_LT(Api::layoutCreate(&layout, Api::RealF32, shape.first, + shape.second, ld)); + layoutStore.emplace(shape, layout); + return layout; + } + + typename Api::MatmulDesc &getOrCreateDesc(typename Api::Operation transA, + typename Api::Operation transB, + typename Api::Epilogue epilogue) { + DescKey key{(int)transA, (int)transB, (int)epilogue}; + auto it = descStore.find(key); + if (it != descStore.end()) + return it->second; + + typename Api::MatmulDesc desc = nullptr; + SOFIEBLAS_CHECK_LT(Api::descCreate(&desc, Api::ComputeF32, Api::RealF32)); + SOFIEBLAS_CHECK_LT( + Api::descSetAttribute(desc, Api::DescTransA, &transA, sizeof(transA))); + SOFIEBLAS_CHECK_LT( + Api::descSetAttribute(desc, Api::DescTransB, &transB, sizeof(transB))); + SOFIEBLAS_CHECK_LT(Api::descSetAttribute(desc, Api::DescEpilogue, &epilogue, + sizeof(epilogue))); + // For bias epilogues: set a non-null dummy pointer so the descriptor is + // valid for the heuristic query. + if (epilogue != Api::EpilogueDefault) { + const void *dummy = d_workspace; + SOFIEBLAS_CHECK_LT(Api::descSetAttribute(desc, Api::DescBiasPointer, + &dummy, sizeof(dummy))); + } + descStore.emplace(key, desc); + return descStore.at(key); + } + + typename Api::HeuristicResult & + getOrComputeAlgo(typename Api::Operation transA, + typename Api::Operation transB, + typename Api::Epilogue epilogue, + const std::pair &shapeA, + const std::pair &shapeB, + const std::pair &shapeC) { + AlgoKey key{{(int)transA, (int)transB, (int)epilogue}, + shapeA.first, + shapeA.second, + shapeB.first, + shapeB.second}; + auto it = algoCache.find(key); + if (it != algoCache.end()) { + if (algoCacheLimit) + lruOrder.splice(lruOrder.begin(), lruOrder, it->second.lru); + return it->second.h; + } + + auto &desc = getOrCreateDesc(transA, transB, epilogue); + auto lA = getOrCreateLayout(shapeA, shapeA.first); + auto lB = getOrCreateLayout(shapeB, shapeB.first); + auto lC = getOrCreateLayout(shapeC, shapeC.first); + typename Api::HeuristicResult h{}; + int returnedResults = 0; + SOFIEBLAS_CHECK_LT(Api::getHeuristic(ltHandle, desc, lA, lB, lC, lC, + preference, 1, &h, &returnedResults)); + if (returnedResults == 0) { + std::cerr << "[sofieBLAS] No suitable " << Api::name + << " algorithm found for " + << "transA=" << transA << " transB=" << transB + << " epilogue=" << epilogue << " A=[" << shapeA.first << "x" + << shapeA.second << "]" + << " B=[" << shapeB.first << "x" << shapeB.second << "]\n"; + exit(EXIT_FAILURE); + } + auto ins = algoCache.emplace(key, CacheEntry{h, {}}).first; + if (algoCacheLimit) { + lruOrder.push_front(key); + ins->second.lru = lruOrder.begin(); + while (algoCache.size() > algoCacheLimit) { + algoCache.erase(lruOrder.back()); + lruOrder.pop_back(); + } + } + return ins->second.h; + } + + void executeMatmul(typename Api::Operation transA, + typename Api::Operation transB, + typename Api::Epilogue 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 &shapeA, + const std::pair &shapeB, + const std::pair &shapeC) { + // Retrieve (or lazily compute) the cached algorithm for this shape + auto &h = + getOrComputeAlgo(transA, transB, epilogue, shapeA, shapeB, shapeC); + + // Retrieve the cached descriptor and patch the real bias pointer in-place + auto &desc = getOrCreateDesc(transA, transB, epilogue); + if (bias_ptr) { + SOFIEBLAS_CHECK_LT(Api::descSetAttribute(desc, Api::DescBiasPointer, + &bias_ptr, sizeof(bias_ptr))); + } + + auto lA = getOrCreateLayout(shapeA, shapeA.first); + auto lB = getOrCreateLayout(shapeB, shapeB.first); + auto lC = getOrCreateLayout(shapeC, shapeC.first); + SOFIEBLAS_CHECK_LT(Api::matmul(ltHandle, desc, &alpha, A, lA, B, lB, &beta, + D_in, lC, C_out, lC, &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..4e178df --- /dev/null +++ b/include/sofieBLAS/backends/hip/sofieBLAS_hipblaslt.hpp @@ -0,0 +1,107 @@ +#pragma once + +#ifdef ALPAKA_ACC_GPU_HIP_ENABLED + +#include +#include +#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) + +// The hipBLASLt forwarding of the shared BlasLt implementation in +// backends/gpu/detail +struct HipblasLtApi { + using Queue = alpaka::QueueHipRtNonBlocking; + using Handle = hipblasLtHandle_t; + using BlasHandle = hipblasHandle_t; + using Preference = hipblasLtMatmulPreference_t; + using Stream = hipStream_t; + using Layout = hipblasLtMatrixLayout_t; + using MatmulDesc = hipblasLtMatmulDesc_t; + using HeuristicResult = hipblasLtMatmulHeuristicResult_t; + using Operation = hipblasOperation_t; + using Epilogue = hipblasLtEpilogue_t; + + static constexpr auto OpN = HIPBLAS_OP_N; + static constexpr auto OpT = HIPBLAS_OP_T; + static constexpr auto OpC = HIPBLAS_OP_C; + static constexpr auto EpilogueDefault = HIPBLASLT_EPILOGUE_DEFAULT; + static constexpr auto EpilogueBias = HIPBLASLT_EPILOGUE_BIAS; + static constexpr auto EpilogueReluBias = HIPBLASLT_EPILOGUE_RELU_BIAS; + static constexpr auto EpilogueGeluBias = HIPBLASLT_EPILOGUE_GELU_BIAS; + static constexpr auto ComputeF32 = HIPBLAS_COMPUTE_32F; + static constexpr auto RealF32 = HIP_R_32F; + static constexpr auto DescTransA = HIPBLASLT_MATMUL_DESC_TRANSA; + static constexpr auto DescTransB = HIPBLASLT_MATMUL_DESC_TRANSB; + static constexpr auto DescEpilogue = HIPBLASLT_MATMUL_DESC_EPILOGUE; + static constexpr auto DescBiasPointer = HIPBLASLT_MATMUL_DESC_BIAS_POINTER; + static constexpr auto PrefMaxWorkspace = + HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES; + static constexpr const char *name = "hipBLASLt"; + + static constexpr auto ltCreate = hipblasLtCreate; + static constexpr auto ltDestroy = hipblasLtDestroy; + static constexpr auto blasCreate = hipblasCreate; + static constexpr auto blasDestroy = hipblasDestroy; + static constexpr auto blasSetStream = hipblasSetStream; + static constexpr auto prefCreate = hipblasLtMatmulPreferenceCreate; + static constexpr auto prefDestroy = hipblasLtMatmulPreferenceDestroy; + static constexpr auto prefSetAttribute = + hipblasLtMatmulPreferenceSetAttribute; + static constexpr auto layoutCreate = hipblasLtMatrixLayoutCreate; + static constexpr auto layoutDestroy = hipblasLtMatrixLayoutDestroy; + static constexpr auto descCreate = hipblasLtMatmulDescCreate; + static constexpr auto descDestroy = hipblasLtMatmulDescDestroy; + static constexpr auto descSetAttribute = hipblasLtMatmulDescSetAttribute; + static constexpr auto getHeuristic = hipblasLtMatmulAlgoGetHeuristic; + static constexpr auto matmul = hipblasLtMatmul; + static constexpr auto sgemmStridedBatched = hipblasSgemmStridedBatched; + + // hipMalloc has a templated C++ overload, so a pointer to it is ambiguous + static hipError_t rtMalloc(void **ptr, std::size_t size) { + return hipMalloc(ptr, size); + } + static hipError_t rtFree(void *ptr) { return hipFree(ptr); } +}; + +#define SOFIEBLAS_CHECK_LT(status) CHECK_HIPBLAS(status) +#define SOFIEBLAS_CHECK_RT(err) CHECK_HIP(err) + +#include "../gpu/detail/sofieBLAS_blaslt_common.tpp" + +using BlasHip = BlasLt; + +namespace traits { + +template <> class sofieBLAS { +public: + using Impl = BlasHip; +}; + +} // namespace traits + +#endif // ALPAKA_ACC_GPU_HIP_ENABLED diff --git a/include/sofieBLAS/core.hpp b/include/sofieBLAS/core.hpp index f2f3da9..1fb8890 100644 --- a/include/sofieBLAS/core.hpp +++ b/include/sofieBLAS/core.hpp @@ -6,3 +6,5 @@ template class sofieBLAS; template using sofieBLAS = typename traits::sofieBLAS::Impl; + +enum class Epilogue { Default, Bias, ReluBias, GeluBias }; 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..512f59b 100644 --- a/tests/test.cc +++ b/tests/test.cc @@ -1,121 +1,991 @@ #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]; + } + } +} + +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; +} + +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)); +} + +static int gFailures = 0; + +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; + } + } + 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"); - // Host platform and device - alpaka::PlatformCpu host_platform{}; - auto host = alpaka::getDevByIdx(host_platform, 0u); + // --- 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"); - // 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); + // --- 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"); + } - // 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); + // --- 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"); + } - for (int i = 0; i < size * size; ++i) { - A[i] = dist(gen); - B[i] = dist(gen); + // --- 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"); } - std::cout << "Matrix A:\n"; - print(A, size); - std::cout << '\n'; - std::cout << "Matrix B:\n"; - print(B, size); - std::cout << '\n'; +} + +#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.addOperationConfig(M, N, K, ldaFor('N', M, K), ldbFor('N', K, N), M, 'N', + 'N', Epilogue::Default); + 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.addOperationConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, + 'T', 'N', Epilogue::Default); + 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.addOperationConfig(M, N, K, ldaFor('N', M, K), ldbFor('T', K, N), M, + 'N', 'T', Epilogue::Default); + 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.addOperationConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, + 'T', 'N', Epilogue::Bias); + 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.addOperationConfig(M, N, K, M, K, M, 'N', 'N', Epilogue::ReluBias); + 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 ---- + { + 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)"); + } + + // ---- 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"); + + // ---- 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"); + + // ---- 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"); + + // ---- 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"); + } +} + +static void runDynamicShapeTests() { + std::cout << "\n=== CUDA Dynamic-Shape Tests ===\n"; + + alpaka::PlatformCudaRt platform{}; + auto dev = alpaka::getDevByIdx(platform, 0u); + alpaka::Queue queue{dev}; + + alpaka::PlatformCpu hostPlatform{}; + auto hostDev = alpaka::getDevByIdx(hostPlatform, 0u); + + // M0 is the construction-time size given to addOperationConfig; the buffers + // hold MCAP rows so sizes above M0 are exercised too. + constexpr int MCAP = 96, M0 = 64, N = 3, K = 5; + + auto hA = alpaka::allocBuf(hostDev, static_cast(MCAP * K)); + auto hB = alpaka::allocBuf(hostDev, static_cast(K * N)); + auto hC = alpaka::allocBuf(hostDev, static_cast(MCAP * N)); + float *A = alpaka::getPtrNative(hA); + float *B = alpaka::getPtrNative(hB); + float *C = alpaka::getPtrNative(hC); + fillSeq(A, MCAP * K, 0.5f, 0.25f); + fillSeq(B, K * N, 1.f, 0.5f); + + auto dA = + alpaka::allocAsyncBuf(queue, static_cast(MCAP * K)); + auto dB = alpaka::allocAsyncBuf(queue, static_cast(K * N)); + auto dC = + alpaka::allocAsyncBuf(queue, static_cast(MCAP * N)); + alpaka::memcpy(queue, dA, hA); + alpaka::memcpy(queue, dB, hB); + alpaka::wait(queue); + + // One instance serving sizes never passed to addOperationConfig (issue #10), + // including m=1 and a size above the construction-time one. + sofieBLAS blas(queue); + blas.addOperationConfig(M0, N, K, ldaFor('N', M0, K), ldbFor('N', K, N), M0, + 'N', 'N', Epilogue::Default); + + std::vector ref; + auto runAt = [&](int m, const std::string &name) { + ref.assign(static_cast(m) * N, 0.f); + refMatmul(ref.data(), A, B, m, N, K, 1.f, 0.f, false, false); + blas.matmul('N', 'N', static_cast(m), static_cast(N), + static_cast(K), 1.f, dA, dB, 0.f, dC); + alpaka::memcpy(queue, hC, dC); + alpaka::wait(queue); + checkClose(C, ref.data(), m * N, name); + }; + + for (int m : {M0, 37, 8, 51, 1, M0, MCAP}) + runAt(m, "cuda::dynamic m=" + std::to_string(m)); + + // Generated code calls the raw-pointer overloads; one call keeps them + // compiled and resolving to the right overload. + ref.assign(static_cast(45) * N, 0.f); + refMatmul(ref.data(), A, B, 45, N, K, 1.f, 0.f, false, false); + blas.matmul('N', 'N', 45u, static_cast(N), static_cast(K), + 1.f, alpaka::getPtrNative(dA), alpaka::getPtrNative(dB), 0.f, + alpaka::getPtrNative(dC)); + alpaka::memcpy(queue, hC, dC); + alpaka::wait(queue); + checkClose(C, ref.data(), 45 * N, "cuda::dynamic raw pointers m=45"); + + // 32 distinct sizes through a cache limited to 8 entries. + { + sofieBLAS capped(queue, 8); + capped.addOperationConfig(M0, N, K, ldaFor('N', M0, K), ldbFor('N', K, N), + M0, 'N', 'N', Epilogue::Default); + float worst = 0.f; + for (int m = M0 + 1; m <= MCAP; ++m) { + ref.assign(static_cast(m) * N, 0.f); + refMatmul(ref.data(), A, B, m, N, K, 1.f, 0.f, false, false); + capped.matmul('N', 'N', static_cast(m), + static_cast(N), static_cast(K), 1.f, dA, + dB, 0.f, dC); + alpaka::memcpy(queue, hC, dC); + alpaka::wait(queue); + for (std::size_t i = 0; i < ref.size(); ++i) + worst = std::max(worst, std::abs(C[i] - ref[i])); + } + if (capped.algoCacheSize() <= 8 && worst < 1e-3f) { + std::cout << " PASS cuda::cache limit honoured\n"; + } else { + std::cerr << " FAIL [cuda::cache limit honoured] " + << capped.algoCacheSize() << " entries, worst err " << worst + << "\n"; + ++gFailures; + } + } +} + +#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); + checkClose(C, ref.data(), M * N, name); + }; + + // ---- matmul NN ---- + blas.addOperationConfig(M, N, K, ldaFor('N', M, K), ldbFor('N', K, N), M, 'N', + 'N', Epilogue::Default); + 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.addOperationConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, + 'T', 'N', Epilogue::Default); + 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.addOperationConfig(M, N, K, ldaFor('N', M, K), ldbFor('T', K, N), M, + 'N', 'T', Epilogue::Default); + 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.addOperationConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, + 'T', 'N', Epilogue::Bias); + 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 ---- { - alpaka::PlatformCpu platform; - alpaka::DevCpu device = alpaka::getDevByIdx(platform, 0u); - alpaka::Queue queue{device}; + 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.addOperationConfig(M, N, K, M, K, M, 'N', 'N', Epilogue::ReluBias); + 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"); + } - 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: 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)"); + } - const float alpha = 1.0f; - const float beta = 0.0f; + // ---- 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"); - 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 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"); - sofieBLAS blas(queue); - blas.gemm('n', 'n', m, n, k, alpha, A, lda, B, ldb, beta, C, ldc); + // ---- 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::cout << "CPU Matrix C = A × B:\n"; - print(C, size); - std::cout << '\n'; + 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"); } +} + +static void runHipDynamicShapeTests() { + std::cout << "\n=== HIP Dynamic-Shape Tests ===\n"; + + alpaka::PlatformHipRt platform{}; + auto dev = alpaka::getDevByIdx(platform, 0u); + alpaka::Queue queue{dev}; + + alpaka::PlatformCpu hostPlatform{}; + auto hostDev = alpaka::getDevByIdx(hostPlatform, 0u); + + // M0 is the construction-time size given to addOperationConfig; the buffers + // hold MCAP rows so sizes above M0 are exercised too. + constexpr int MCAP = 96, M0 = 64, N = 3, K = 5; + + auto hA = alpaka::allocBuf(hostDev, static_cast(MCAP * K)); + auto hB = alpaka::allocBuf(hostDev, static_cast(K * N)); + auto hC = alpaka::allocBuf(hostDev, static_cast(MCAP * N)); + float *A = alpaka::getPtrNative(hA); + float *B = alpaka::getPtrNative(hB); + float *C = alpaka::getPtrNative(hC); + fillSeq(A, MCAP * K, 0.5f, 0.25f); + fillSeq(B, K * N, 1.f, 0.5f); + + auto dA = + alpaka::allocAsyncBuf(queue, static_cast(MCAP * K)); + auto dB = alpaka::allocAsyncBuf(queue, static_cast(K * N)); + auto dC = + alpaka::allocAsyncBuf(queue, static_cast(MCAP * N)); + alpaka::memcpy(queue, dA, hA); + alpaka::memcpy(queue, dB, hB); + alpaka::wait(queue); + + // One instance serving sizes never passed to addOperationConfig (issue #10), + // including m=1 and a size above the construction-time one. + sofieBLAS blas(queue); + blas.addOperationConfig(M0, N, K, ldaFor('N', M0, K), ldbFor('N', K, N), M0, + 'N', 'N', Epilogue::Default); + + std::vector ref; + auto runAt = [&](int m, const std::string &name) { + ref.assign(static_cast(m) * N, 0.f); + refMatmul(ref.data(), A, B, m, N, K, 1.f, 0.f, false, false); + blas.matmul('N', 'N', static_cast(m), static_cast(N), + static_cast(K), 1.f, dA, dB, 0.f, dC); + alpaka::memcpy(queue, hC, dC); + alpaka::wait(queue); + checkClose(C, ref.data(), m * N, name); + }; + + for (int m : {M0, 37, 8, 51, 1, M0, MCAP}) + runAt(m, "hip::dynamic m=" + std::to_string(m)); + + // Generated code calls the raw-pointer overloads; one call keeps them + // compiled and resolving to the right overload. + ref.assign(static_cast(45) * N, 0.f); + refMatmul(ref.data(), A, B, 45, N, K, 1.f, 0.f, false, false); + blas.matmul('N', 'N', 45u, static_cast(N), static_cast(K), + 1.f, alpaka::getPtrNative(dA), alpaka::getPtrNative(dB), 0.f, + alpaka::getPtrNative(dC)); + alpaka::memcpy(queue, hC, dC); + alpaka::wait(queue); + checkClose(C, ref.data(), 45 * N, "hip::dynamic raw pointers m=45"); + + // 32 distinct sizes through a cache limited to 8 entries. + { + sofieBLAS capped(queue, 8); + capped.addOperationConfig(M0, N, K, ldaFor('N', M0, K), ldbFor('N', K, N), + M0, 'N', 'N', Epilogue::Default); + float worst = 0.f; + for (int m = M0 + 1; m <= MCAP; ++m) { + ref.assign(static_cast(m) * N, 0.f); + refMatmul(ref.data(), A, B, m, N, K, 1.f, 0.f, false, false); + capped.matmul('N', 'N', static_cast(m), + static_cast(N), static_cast(K), 1.f, dA, + dB, 0.f, dC); + alpaka::memcpy(queue, hC, dC); + alpaka::wait(queue); + for (std::size_t i = 0; i < ref.size(); ++i) + worst = std::max(worst, std::abs(C[i] - ref[i])); + } + if (capped.algoCacheSize() <= 8 && worst < 1e-3f) { + std::cout << " PASS hip::cache limit honoured\n"; + } else { + std::cerr << " FAIL [hip::cache limit honoured] " + << capped.algoCacheSize() << " entries, worst err " << worst + << "\n"; + ++gFailures; + } + } +} + +#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(); + runDynamicShapeTests(); #endif +#ifdef ALPAKA_ACC_GPU_HIP_ENABLED + runHipTests(); + runHipDynamicShapeTests(); +#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; }