diff --git a/.github/workflows/tests.yml b/.github/workflows/tests.yml index e41ae81..e14c3f2 100644 --- a/.github/workflows/tests.yml +++ b/.github/workflows/tests.yml @@ -19,7 +19,11 @@ env: jobs: cpu-tests: - name: CPU Unit Tests + name: CPU Unit Tests (${{ matrix.blas }}) + strategy: + fail-fast: false + matrix: + blas: [OpenBLAS, MKL] if: | github.event_name != 'issue_comment' || ( @@ -28,7 +32,7 @@ jobs: !startsWith(github.event.comment.body, '/runbenchmark') ) concurrency: - group: cpu-tests-${{ github.event.issue.number || github.ref }} + group: cpu-tests-${{ matrix.blas }}-${{ github.event.issue.number || github.ref }} cancel-in-progress: true runs-on: ubuntu-latest timeout-minutes: 30 @@ -38,10 +42,19 @@ jobs: with: ref: ${{ github.event_name == 'issue_comment' && format('refs/pull/{0}/merge', github.event.issue.number) || github.ref }} - - name: Install OpenBLAS + - name: Install ${{ matrix.blas }} run: | - sudo apt-get update - sudo apt-get install -y libopenblas-dev + if [ "${{ matrix.blas }}" = MKL ]; then + wget -qO- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB \ + | gpg --dearmor | sudo tee /usr/share/keyrings/oneapi-archive-keyring.gpg > /dev/null + echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" \ + | sudo tee /etc/apt/sources.list.d/oneAPI.list + sudo apt-get update + sudo apt-get install -y intel-oneapi-mkl-devel + else + sudo apt-get update + sudo apt-get install -y libopenblas-dev + fi - name: Cache FetchContent dependencies uses: actions/cache@v4 @@ -51,6 +64,7 @@ jobs: restore-keys: cmake-deps-cpu- - name: Configure + shell: bash run: | cmake -B build -S . \ -DCMAKE_BUILD_TYPE=${{ env.BUILD_TYPE }} \ @@ -58,20 +72,25 @@ jobs: -DSOFIEBLAS_BUILD_BENCHMARKS=ON \ -DSOFIEBLAS_ENABLE_CUDA=OFF \ -DSOFIEBLAS_ENABLE_HIP=OFF \ - "-DFETCHCONTENT_BASE_DIR=${{ env.DEPS_CACHE }}" + -DCPU_BLAS_LIB=${{ matrix.blas }} \ + "-DFETCHCONTENT_BASE_DIR=${{ env.DEPS_CACHE }}" | tee configure.log + grep -q "using CPU BLAS library ${{ matrix.blas }}" configure.log - name: Build run: cmake --build build -j"$(nproc)" - name: Run tests working-directory: build + env: + LD_LIBRARY_PATH: /opt/intel/oneapi/mkl/latest/lib/intel64:/opt/intel/oneapi/mkl/latest/lib + MKL_THREADING_LAYER: SEQUENTIAL 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 }} + name: cpu-test-log-${{ matrix.blas }}-${{ github.run_id }} path: build/Testing/Temporary/LastTest.log if-no-files-found: ignore diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 2c55d06..1619bef 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -33,6 +33,7 @@ if(SOFIEBLAS_CUDA_ENABLED) ) target_compile_definitions(test_cuda PRIVATE ALPAKA_ACC_GPU_CUDA_ENABLED) + target_include_directories(test_cuda PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}) target_link_libraries(test_cuda PRIVATE sofieBLAS::sofieBLAS alpaka::alpaka CUDA::cudart CUDA::cublas CUDA::cublasLt) @@ -51,7 +52,7 @@ if(SOFIEBLAS_HIP_ENABLED) target_compile_options(test_hip PRIVATE ${CXXFLAGS} ${CXX_HOST_FLAGS}) target_compile_definitions(test_hip PRIVATE ALPAKA_ACC_GPU_HIP_ENABLED) - target_include_directories(test_hip PRIVATE ${ROCM_BASE}/include) + target_include_directories(test_hip PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${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) diff --git a/tests/cpu/unit_test.tpp b/tests/cpu/unit_test.tpp new file mode 100644 index 0000000..fb3d03f --- /dev/null +++ b/tests/cpu/unit_test.tpp @@ -0,0 +1,220 @@ +// CPU backend tests, included from test.cc + +static void runCpuTests() { + std::cout << "\n=== CPU Tests ===\n"; + + alpaka::PlatformCpu platform{}; + auto dev = alpaka::getDevByIdx(platform, 0u); + alpaka::Queue queue{dev}; + sofieBLAS blas(queue); + + constexpr int M = 4, N = 3, K = 5; + + // Allocate host buffers + auto hA = alpaka::allocBuf(dev, static_cast(M * K)); + auto hB = alpaka::allocBuf(dev, static_cast(K * N)); + auto hC = alpaka::allocBuf(dev, static_cast(M * N)); + auto hBias = alpaka::allocBuf(dev, static_cast(M * N)); + + float *A = alpaka::getPtrNative(hA); + float *B = alpaka::getPtrNative(hB); + float *C = alpaka::getPtrNative(hC); + float *bias = alpaka::getPtrNative(hBias); + + fillSeq(A, M * K); + fillSeq(B, K * N, 1.f, 0.5f); + fillSeq(bias, M * N, 0.1f, 0.1f); + + std::vector ref(M * N); + + // --- matmul NN --- + fillVal(C, M * N, 0.f); + blas.matmul('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hC); + std::copy(C, C + M * N, ref.data()); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::matmul NN"); + + // --- matmul TN (A^T: K×M physical → M×K logical) --- + { + auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + fillVal(C, M * N, 0.f); + blas.matmul('T', 'N', M, N, K, 1.f, hAt, hB, 0.f, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), At, B, M, N, K, 1.f, 0.f, true, false); + checkClose(C, ref.data(), M * N, "cpu::matmul TN"); + } + + // --- matmul NT --- + { + auto hBt = alpaka::allocBuf(dev, static_cast(N * K)); + float *Bt = alpaka::getPtrNative(hBt); + fillSeq(Bt, N * K, 1.f, 0.5f); + fillVal(C, M * N, 0.f); + blas.matmul('N', 'T', M, N, K, 1.f, hA, hBt, 0.f, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, Bt, M, N, K, 1.f, 0.f, false, true); + checkClose(C, ref.data(), M * N, "cpu::matmul NT"); + } + + // --- matmul TT --- + { + auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); + auto hBt = alpaka::allocBuf(dev, static_cast(N * K)); + float *At = alpaka::getPtrNative(hAt); + float *Bt = alpaka::getPtrNative(hBt); + fillSeq(At, K * M); + fillSeq(Bt, N * K, 1.f, 0.5f); + fillVal(C, M * N, 0.f); + blas.matmul('T', 'T', M, N, K, 1.f, hAt, hBt, 0.f, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), At, Bt, M, N, K, 1.f, 0.f, true, true); + checkClose(C, ref.data(), M * N, "cpu::matmul TT"); + } + + // --- matmul: alpha scaling --- + fillVal(C, M * N, 0.f); + blas.matmul('N', 'N', M, N, K, 2.5f, hA, hB, 0.f, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, B, M, N, K, 2.5f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::matmul alpha=2.5"); + + // --- matmul: beta accumulation --- + fillSeq(C, M * N, 10.f); // pre-fill C + blas.matmul('N', 'N', M, N, K, 1.f, hA, hB, 0.5f, hC); + { + std::vector C0(M * N); + fillSeq(C0.data(), M * N, 10.f); + refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.5f, false, false); + std::copy(C0.begin(), C0.end(), ref.data()); + refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.5f, false, false); + } + checkClose(C, ref.data(), M * N, "cpu::matmul beta=0.5"); + + // --- gemm NN (beta=0, no prior accumulation) --- + fillVal(C, M * N, 0.f); + fillSeq(bias, M * N, 0.1f, 0.1f); + blas.gemm('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemm NN beta=0"); + + // --- gemm NN (beta=1 accumulation) --- + fillVal(C, M * N, 0.f); + blas.gemm('N', 'N', M, N, K, 1.f, hA, hB, 1.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 1.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemm NN beta=1"); + + // --- gemm TN --- + { + auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + fillVal(C, M * N, 0.f); + blas.gemm('T', 'N', M, N, K, 1.f, hAt, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); + checkClose(C, ref.data(), M * N, "cpu::gemm TN"); + } + + // --- gemmrelu: all-positive matmul result stays unchanged --- + { + // A and B with positive values ensure result is positive before bias + auto hAp = alpaka::allocBuf(dev, static_cast(M * K)); + auto hBp = alpaka::allocBuf(dev, static_cast(K * N)); + auto hBiasp = alpaka::allocBuf(dev, static_cast(M * N)); + float *Ap = alpaka::getPtrNative(hAp); + float *Bp = alpaka::getPtrNative(hBp); + float *biasp = alpaka::getPtrNative(hBiasp); + fillSeq(Ap, M * K, 0.1f, 0.1f); + fillSeq(Bp, K * N, 0.1f, 0.1f); + fillVal(biasp, M * N, 0.f); + fillVal(C, M * N, 0.f); + blas.gemmrelu('N', 'N', M, N, K, 1.f, hAp, hBp, 0.f, hBiasp, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), Ap, Bp, biasp, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemmrelu all-positive"); + } + + // --- gemmrelu: negative values clamped to zero --- + { + // Use alpha=-1 to force negative results + auto hBiasz = alpaka::allocBuf(dev, static_cast(M * N)); + fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); + fillVal(C, M * N, 0.f); + blas.gemmrelu('N', 'N', M, N, K, -1.f, hA, hB, 0.f, hBiasz, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), A, B, alpaka::getPtrNative(hBiasz), M, N, K, -1.f, + 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemmrelu alpha=-1 (clamped)"); + } + + // --- gemmrelu with bias --- + fillVal(C, M * N, 0.f); + fillSeq(bias, M * N, -5.f, 2.f); + blas.gemmrelu('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemmrelu with mixed bias"); + + // --- gemmgelu NN --- + fillVal(C, M * N, 0.f); + fillVal(bias, M * N, 0.f); + blas.gemmgelu('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemmgelu NN"); + + // --- gemmgelu with bias --- + fillVal(C, M * N, 0.f); + fillSeq(bias, M * N, -2.f, 0.5f); + blas.gemmgelu('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); + checkClose(C, ref.data(), M * N, "cpu::gemmgelu with bias"); + + // --- gemmgelu TN --- + { + auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + fillVal(C, M * N, 0.f); + fillVal(bias, M * N, 0.f); + blas.gemmgelu('T', 'N', M, N, K, 1.f, hAt, hB, 0.f, hBias, hC); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmGelu(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); + checkClose(C, ref.data(), M * N, "cpu::gemmgelu TN"); + } + + // --- edge: zero matrix --- + { + auto hZ = alpaka::allocBuf(dev, static_cast(M * K)); + fillVal(alpaka::getPtrNative(hZ), M * K, 0.f); + fillVal(C, M * N, 99.f); + fillVal(bias, M * N, 0.f); + blas.matmul('N', 'N', M, N, K, 1.f, hZ, hB, 0.f, hC); + std::fill(ref.begin(), ref.end(), 0.f); + checkClose(C, ref.data(), M * N, "cpu::matmul zero-A"); + } + + // --- edge: identity-like (square, known result) --- + { + constexpr int S = 3; + auto hI = alpaka::allocBuf(dev, static_cast(S * S)); + auto hX = alpaka::allocBuf(dev, static_cast(S * S)); + auto hY = alpaka::allocBuf(dev, static_cast(S * S)); + float *I = alpaka::getPtrNative(hI); + float *X = alpaka::getPtrNative(hX); + float *Y = alpaka::getPtrNative(hY); + fillVal(I, S * S, 0.f); + for (int i = 0; i < S; ++i) + I[i * S + i] = 1.f; + fillSeq(X, S * S); + fillVal(Y, S * S, 0.f); + blas.matmul('N', 'N', S, S, S, 1.f, hI, hX, 0.f, hY); + checkClose(Y, X, S * S, "cpu::matmul identity×X=X"); + } +} diff --git a/tests/gpu/unit_test.tpp b/tests/gpu/unit_test.tpp new file mode 100644 index 0000000..614b23a --- /dev/null +++ b/tests/gpu/unit_test.tpp @@ -0,0 +1,316 @@ + +// shared test for CUDA/HIP + +template static void runGpuTests() { + std::cout << "\n=== " << __PRETTY_FUNCTION__ << " ===\n"; + + using Acc = alpaka::TagToAcc; + using DevAcc = alpaka::Dev; + using PlatformAcc = alpaka::Platform; + + PlatformAcc 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("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("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("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("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("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("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("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); + 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("gemmrelu all-positive"); + } + + // ---- 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("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("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("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("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("matmul zero-A"); + } +} + +template +static void runGpuDynamicShapeTests() { + std::cout << "\n=== " << __PRETTY_FUNCTION__ << " ===\n"; + + using Acc = alpaka::TagToAcc; + using DevAcc = alpaka::Dev; + using PlatformAcc = alpaka::Platform; + + PlatformAcc 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, "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, "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 cache limit honoured\n"; + } else { + std::cerr << " FAIL [cache limit honoured] " + << capped.algoCacheSize() << " entries, worst err " << worst + << "\n"; + ++gFailures; + } + } +} + diff --git a/tests/test.cc b/tests/test.cc index 512f59b..9e2ce01 100644 --- a/tests/test.cc +++ b/tests/test.cc @@ -110,224 +110,7 @@ static void fillVal(float *M, int n, float v) { #ifdef ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED -static void runCpuTests() { - std::cout << "\n=== CPU Tests ===\n"; - - alpaka::PlatformCpu platform{}; - auto dev = alpaka::getDevByIdx(platform, 0u); - alpaka::Queue queue{dev}; - sofieBLAS blas(queue); - - constexpr int M = 4, N = 3, K = 5; - - // Allocate host buffers - auto hA = alpaka::allocBuf(dev, static_cast(M * K)); - auto hB = alpaka::allocBuf(dev, static_cast(K * N)); - auto hC = alpaka::allocBuf(dev, static_cast(M * N)); - auto hBias = alpaka::allocBuf(dev, static_cast(M * N)); - - float *A = alpaka::getPtrNative(hA); - float *B = alpaka::getPtrNative(hB); - float *C = alpaka::getPtrNative(hC); - float *bias = alpaka::getPtrNative(hBias); - - fillSeq(A, M * K); - fillSeq(B, K * N, 1.f, 0.5f); - fillSeq(bias, M * N, 0.1f, 0.1f); - - std::vector ref(M * N); - - // --- matmul NN --- - fillVal(C, M * N, 0.f); - blas.matmul('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hC); - std::copy(C, C + M * N, ref.data()); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.f, false, false); - checkClose(C, ref.data(), M * N, "cpu::matmul NN"); - - // --- matmul TN (A^T: K×M physical → M×K logical) --- - { - auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); - float *At = alpaka::getPtrNative(hAt); - fillSeq(At, K * M); - fillVal(C, M * N, 0.f); - blas.matmul('T', 'N', M, N, K, 1.f, hAt, hB, 0.f, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), At, B, M, N, K, 1.f, 0.f, true, false); - checkClose(C, ref.data(), M * N, "cpu::matmul TN"); - } - - // --- matmul NT --- - { - auto hBt = alpaka::allocBuf(dev, static_cast(N * K)); - float *Bt = alpaka::getPtrNative(hBt); - fillSeq(Bt, N * K, 1.f, 0.5f); - fillVal(C, M * N, 0.f); - blas.matmul('N', 'T', M, N, K, 1.f, hA, hBt, 0.f, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), A, Bt, M, N, K, 1.f, 0.f, false, true); - checkClose(C, ref.data(), M * N, "cpu::matmul NT"); - } - - // --- matmul TT --- - { - auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); - auto hBt = alpaka::allocBuf(dev, static_cast(N * K)); - float *At = alpaka::getPtrNative(hAt); - float *Bt = alpaka::getPtrNative(hBt); - fillSeq(At, K * M); - fillSeq(Bt, N * K, 1.f, 0.5f); - fillVal(C, M * N, 0.f); - blas.matmul('T', 'T', M, N, K, 1.f, hAt, hBt, 0.f, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), At, Bt, M, N, K, 1.f, 0.f, true, true); - checkClose(C, ref.data(), M * N, "cpu::matmul TT"); - } - - // --- matmul: alpha scaling --- - fillVal(C, M * N, 0.f); - blas.matmul('N', 'N', M, N, K, 2.5f, hA, hB, 0.f, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), A, B, M, N, K, 2.5f, 0.f, false, false); - checkClose(C, ref.data(), M * N, "cpu::matmul alpha=2.5"); - - // --- matmul: beta accumulation --- - fillSeq(C, M * N, 10.f); // pre-fill C - blas.matmul('N', 'N', M, N, K, 1.f, hA, hB, 0.5f, hC); - { - std::vector C0(M * N); - fillSeq(C0.data(), M * N, 10.f); - refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.5f, false, false); - std::copy(C0.begin(), C0.end(), ref.data()); - refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.5f, false, false); - } - checkClose(C, ref.data(), M * N, "cpu::matmul beta=0.5"); - - // --- gemm NN (beta=0, no prior accumulation) --- - fillVal(C, M * N, 0.f); - fillSeq(bias, M * N, 0.1f, 0.1f); - blas.gemm('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - checkClose(C, ref.data(), M * N, "cpu::gemm NN beta=0"); - - // --- gemm NN (beta=1 accumulation) --- - fillVal(C, M * N, 0.f); - blas.gemm('N', 'N', M, N, K, 1.f, hA, hB, 1.f, hBias, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 1.f, false, false); - checkClose(C, ref.data(), M * N, "cpu::gemm NN beta=1"); - - // --- gemm TN --- - { - auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); - float *At = alpaka::getPtrNative(hAt); - fillSeq(At, K * M); - fillVal(C, M * N, 0.f); - blas.gemm('T', 'N', M, N, K, 1.f, hAt, hB, 0.f, hBias, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refGemm(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); - checkClose(C, ref.data(), M * N, "cpu::gemm TN"); - } - - // --- gemmrelu: all-positive matmul result stays unchanged --- - { - // A and B with positive values ensure result is positive before bias - auto hAp = alpaka::allocBuf(dev, static_cast(M * K)); - auto hBp = alpaka::allocBuf(dev, static_cast(K * N)); - auto hBiasp = alpaka::allocBuf(dev, static_cast(M * N)); - float *Ap = alpaka::getPtrNative(hAp); - float *Bp = alpaka::getPtrNative(hBp); - float *biasp = alpaka::getPtrNative(hBiasp); - fillSeq(Ap, M * K, 0.1f, 0.1f); - fillSeq(Bp, K * N, 0.1f, 0.1f); - fillVal(biasp, M * N, 0.f); - fillVal(C, M * N, 0.f); - blas.gemmrelu('N', 'N', M, N, K, 1.f, hAp, hBp, 0.f, hBiasp, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmRelu(ref.data(), Ap, Bp, biasp, M, N, K, 1.f, 0.f, false, false); - checkClose(C, ref.data(), M * N, "cpu::gemmrelu all-positive"); - } - - // --- gemmrelu: negative values clamped to zero --- - { - // Use alpha=-1 to force negative results - auto hBiasz = alpaka::allocBuf(dev, static_cast(M * N)); - fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); - fillVal(C, M * N, 0.f); - blas.gemmrelu('N', 'N', M, N, K, -1.f, hA, hB, 0.f, hBiasz, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmRelu(ref.data(), A, B, alpaka::getPtrNative(hBiasz), M, N, K, -1.f, - 0.f, false, false); - checkClose(C, ref.data(), M * N, "cpu::gemmrelu alpha=-1 (clamped)"); - } - - // --- gemmrelu with bias --- - fillVal(C, M * N, 0.f); - fillSeq(bias, M * N, -5.f, 2.f); - blas.gemmrelu('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmRelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - checkClose(C, ref.data(), M * N, "cpu::gemmrelu with mixed bias"); - - // --- gemmgelu NN --- - fillVal(C, M * N, 0.f); - fillVal(bias, M * N, 0.f); - blas.gemmgelu('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - checkClose(C, ref.data(), M * N, "cpu::gemmgelu NN"); - - // --- gemmgelu with bias --- - fillVal(C, M * N, 0.f); - fillSeq(bias, M * N, -2.f, 0.5f); - blas.gemmgelu('N', 'N', M, N, K, 1.f, hA, hB, 0.f, hBias, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - checkClose(C, ref.data(), M * N, "cpu::gemmgelu with bias"); - - // --- gemmgelu TN --- - { - auto hAt = alpaka::allocBuf(dev, static_cast(K * M)); - float *At = alpaka::getPtrNative(hAt); - fillSeq(At, K * M); - fillVal(C, M * N, 0.f); - fillVal(bias, M * N, 0.f); - blas.gemmgelu('T', 'N', M, N, K, 1.f, hAt, hB, 0.f, hBias, hC); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmGelu(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); - checkClose(C, ref.data(), M * N, "cpu::gemmgelu TN"); - } - - // --- edge: zero matrix --- - { - auto hZ = alpaka::allocBuf(dev, static_cast(M * K)); - fillVal(alpaka::getPtrNative(hZ), M * K, 0.f); - fillVal(C, M * N, 99.f); - fillVal(bias, M * N, 0.f); - blas.matmul('N', 'N', M, N, K, 1.f, hZ, hB, 0.f, hC); - std::fill(ref.begin(), ref.end(), 0.f); - checkClose(C, ref.data(), M * N, "cpu::matmul zero-A"); - } - - // --- edge: identity-like (square, known result) --- - { - constexpr int S = 3; - auto hI = alpaka::allocBuf(dev, static_cast(S * S)); - auto hX = alpaka::allocBuf(dev, static_cast(S * S)); - auto hY = alpaka::allocBuf(dev, static_cast(S * S)); - float *I = alpaka::getPtrNative(hI); - float *X = alpaka::getPtrNative(hX); - float *Y = alpaka::getPtrNative(hY); - fillVal(I, S * S, 0.f); - for (int i = 0; i < S; ++i) - I[i * S + i] = 1.f; - fillSeq(X, S * S); - fillVal(Y, S * S, 0.f); - blas.matmul('N', 'N', S, S, S, 1.f, hI, hX, 0.f, hY); - checkClose(Y, X, S * S, "cpu::matmul identity×X=X"); - } -} +#include "cpu/unit_test.tpp" #endif // ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED @@ -338,632 +121,9 @@ static int ldaFor(char trans, int m, int k) { static int ldbFor(char trans, int k, int n) { return (trans == 'N' || trans == 'n') ? k : n; } +#include "gpu/unit_test.tpp" #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 ---- - { - 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"); - } - - // ---- 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"); - } - - // ---- 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"); - - // ---- 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"); - - // ---- 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"); - - // ---- 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); - 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"); - } - - // ---- 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 ---- - { - 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"); - } - - // ---- gemmrelu: alpha=-1 forces negatives ---- - { - auto hBiasz = - alpaka::allocBuf(hostDev, static_cast(M * N)); - fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); - auto dBiasz = - alpaka::allocAsyncBuf(queue, static_cast(M * N)); - alpaka::memcpy(queue, dBiasz, hBiasz); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmRelu(ref.data(), A, B, alpaka::getPtrNative(hBiasz), M, N, K, -1.f, - 0.f, false, false); - blas.gemmrelu('N', 'N', M, N, K, -1.f, dA, dB, 0.f, dBiasz, dC); - verify("hip::gemmrelu alpha=-1 (clamped)"); - } - - // ---- gemmrelu with mixed bias ---- - fillSeq(bias, M * N, -5.f, 2.f); - alpaka::memcpy(queue, dBias, hBias); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmRelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - blas.gemmrelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); - verify("hip::gemmrelu with mixed bias"); - - // ---- gemmgelu NN ---- - fillVal(bias, M * N, 0.f); - alpaka::memcpy(queue, dBias, hBias); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - blas.gemmgelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); - verify("hip::gemmgelu NN"); - - // ---- gemmgelu with bias ---- - fillSeq(bias, M * N, -2.f, 0.5f); - alpaka::memcpy(queue, dBias, hBias); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - blas.gemmgelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); - verify("hip::gemmgelu with bias"); - - // ---- edge: zero A ---- - { - auto hZero = alpaka::allocBuf(hostDev, static_cast(M * K)); - fillVal(alpaka::getPtrNative(hZero), M * K, 0.f); - auto dZero = - alpaka::allocAsyncBuf(queue, static_cast(M * K)); - alpaka::memcpy(queue, dZero, hZero); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - blas.matmul('N', 'N', M, N, K, 1.f, dZero, dB, 0.f, dC); - verify("hip::matmul zero-A"); - } -} - -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 // --------------------------------------------------------------------------- @@ -973,12 +133,12 @@ int main() { runCpuTests(); #endif #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED - runCudaTests(); - runDynamicShapeTests(); + runGpuTests(); + runGpuDynamicShapeTests(); #endif #ifdef ALPAKA_ACC_GPU_HIP_ENABLED - runHipTests(); - runHipDynamicShapeTests(); + runGpuTests(); + runGpuDynamicShapeTests(); #endif std::cout << "\n";