From 248fdf028cd57188f891c064c801442fd900e2e1 Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Thu, 13 Aug 2026 23:43:36 +0530 Subject: [PATCH 1/6] test: deduplicate cuda and hip tests into a tag templated fixture test --- tests/CMakeLists.txt | 2 +- tests/gpu_tests.tpp | 221 +++++++++++++++++++++ tests/test.cc | 448 +------------------------------------------ 3 files changed, 227 insertions(+), 444 deletions(-) create mode 100644 tests/gpu_tests.tpp diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 1247ec8..d51b49a 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -177,7 +177,7 @@ if(CMAKE_HIP_COMPILER) ) target_compile_definitions(test_hip PRIVATE ALPAKA_ACC_GPU_HIP_ENABLED) - target_include_directories(test_hip PRIVATE ${ALPAKA_BASE}/include "../sofieBLAS/include" ${ROCM_BASE}/include) + target_include_directories(test_hip PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${ALPAKA_BASE}/include "../sofieBLAS/include" ${ROCM_BASE}/include) target_link_directories(test_hip PRIVATE ${ROCM_BASE}/lib) target_link_libraries(test_hip PRIVATE hipblaslt hipblas amdhip64) endif() diff --git a/tests/gpu_tests.tpp b/tests/gpu_tests.tpp new file mode 100644 index 0000000..eb764ed --- /dev/null +++ b/tests/gpu_tests.tpp @@ -0,0 +1,221 @@ + +// shared test for CUDA/HIP + +template + +static void runGpuTests(const std::string &backend) { + std::cout << "\n=== " << backend << " Tests ===\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.addLayoutConfig(M, N, K, ldaFor('N', M, K), ldbFor('N', K, N), M, 'N', + 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.f, false, false); + blas.matmul('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dC); + verify(backend + "::matmul NN"); + + // ---- matmul TN ---- + { + auto hAt = alpaka::allocBuf(hostDev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + auto dAt = + alpaka::allocAsyncBuf(queue, static_cast(K * M)); + alpaka::memcpy(queue, dAt, hAt); + alpaka::wait(queue); + blas.addLayoutConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, 'T', + 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), At, B, M, N, K, 1.f, 0.f, true, false); + blas.matmul('T', 'N', M, N, K, 1.f, dAt, dB, 0.f, dC); + verify(backend + "::matmul TN"); + } + + // ---- matmul NT ---- + { + auto hBt = alpaka::allocBuf(hostDev, static_cast(N * K)); + float *Bt = alpaka::getPtrNative(hBt); + fillSeq(Bt, N * K, 1.f, 0.5f); + auto dBt = + alpaka::allocAsyncBuf(queue, static_cast(N * K)); + alpaka::memcpy(queue, dBt, hBt); + alpaka::wait(queue); + blas.addLayoutConfig(M, N, K, ldaFor('N', M, K), ldbFor('T', K, N), M, 'N', + 'T'); + std::fill(ref.begin(), ref.end(), 0.f); + refMatmul(ref.data(), A, Bt, M, N, K, 1.f, 0.f, false, true); + blas.matmul('N', 'T', M, N, K, 1.f, dA, dBt, 0.f, dC); + verify(backend + "::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(backend + "::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(backend + "::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(backend + "::gemm NN beta=1"); + + // ---- gemm TN ---- + { + auto hAt = alpaka::allocBuf(hostDev, static_cast(K * M)); + float *At = alpaka::getPtrNative(hAt); + fillSeq(At, K * M); + auto dAt = + alpaka::allocAsyncBuf(queue, static_cast(K * M)); + alpaka::memcpy(queue, dAt, hAt); + alpaka::wait(queue); + blas.addLayoutConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, 'T', + 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refGemm(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); + blas.gemm('T', 'N', M, N, K, 1.f, dAt, dB, 0.f, dBias, dC); + verify(backend + "::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.addLayoutConfig(M, N, K, M, K, M, 'N', 'N'); + std::fill(ref.begin(), ref.end(), 0.f); + refGemmRelu(ref.data(), Ap, Bp, alpaka::getPtrNative(hBiasz), M, N, K, 1.f, + 0.f, false, false); + blas.gemmrelu('N', 'N', M, N, K, 1.f, dAp, dBp, 0.f, dBiasz, dC); + verify(backend + "::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(backend + "::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(backend + "::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(backend + "::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(backend + "::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(backend + "::matmul zero-A"); + } +} diff --git a/tests/test.cc b/tests/test.cc index 4b47b67..da88625 100644 --- a/tests/test.cc +++ b/tests/test.cc @@ -344,447 +344,9 @@ static int ldbFor(char trans, int k, int n) { } #endif -// --------------------------------------------------------------------------- -// CUDA tests -// --------------------------------------------------------------------------- - -#ifdef ALPAKA_ACC_GPU_CUDA_ENABLED - -static void runCudaTests() { - std::cout << "\n=== CUDA Tests ===\n"; - - alpaka::PlatformCudaRt platform{}; - auto dev = alpaka::getDevByIdx(platform, 0u); - alpaka::Queue queue{dev}; - sofieBLAS blas(queue); - - alpaka::PlatformCpu hostPlatform{}; - auto hostDev = alpaka::getDevByIdx(hostPlatform, 0u); - - constexpr int M = 4, N = 3, K = 5; - - auto hA = alpaka::allocBuf(hostDev, static_cast(M * K)); - auto hB = alpaka::allocBuf(hostDev, static_cast(K * N)); - auto hC = alpaka::allocBuf(hostDev, static_cast(M * N)); - auto hBias = alpaka::allocBuf(hostDev, static_cast(M * N)); - - float *A = alpaka::getPtrNative(hA); - float *B = alpaka::getPtrNative(hB); - float *bias = alpaka::getPtrNative(hBias); - - fillSeq(A, M * K); - fillSeq(B, K * N, 1.f, 0.5f); - fillVal(bias, M * N, 0.f); - - auto dA = alpaka::allocAsyncBuf(queue, static_cast(M * K)); - auto dB = alpaka::allocAsyncBuf(queue, static_cast(K * N)); - auto dC = alpaka::allocAsyncBuf(queue, static_cast(M * N)); - auto dBias = - alpaka::allocAsyncBuf(queue, static_cast(M * N)); - - alpaka::memcpy(queue, dA, hA); - alpaka::memcpy(queue, dB, hB); - alpaka::memcpy(queue, dBias, hBias); - alpaka::wait(queue); - - std::vector ref(M * N); - float *C = alpaka::getPtrNative(hC); - - auto verify = [&](const std::string &name) { - alpaka::memcpy(queue, hC, dC); - alpaka::wait(queue); - checkClose(C, ref.data(), M * N, name); - }; - - // ---- matmul NN ---- - blas.addLayoutConfig(M, N, K, ldaFor('N', M, K), ldbFor('N', K, N), M, 'N', - 'N'); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.f, false, false); - blas.matmul('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dC); - verify("cuda::matmul NN"); - - // ---- matmul TN ---- - { - auto hAt = alpaka::allocBuf(hostDev, static_cast(K * M)); - float *At = alpaka::getPtrNative(hAt); - fillSeq(At, K * M); - auto dAt = - alpaka::allocAsyncBuf(queue, static_cast(K * M)); - alpaka::memcpy(queue, dAt, hAt); - alpaka::wait(queue); - blas.addLayoutConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, 'T', - 'N'); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), At, B, M, N, K, 1.f, 0.f, true, false); - blas.matmul('T', 'N', M, N, K, 1.f, dAt, dB, 0.f, dC); - verify("cuda::matmul TN"); - } - - // ---- matmul NT ---- - { - auto hBt = alpaka::allocBuf(hostDev, static_cast(N * K)); - float *Bt = alpaka::getPtrNative(hBt); - fillSeq(Bt, N * K, 1.f, 0.5f); - auto dBt = - alpaka::allocAsyncBuf(queue, static_cast(N * K)); - alpaka::memcpy(queue, dBt, hBt); - alpaka::wait(queue); - blas.addLayoutConfig(M, N, K, ldaFor('N', M, K), ldbFor('T', K, N), M, 'N', - 'T'); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), A, Bt, M, N, K, 1.f, 0.f, false, true); - blas.matmul('N', 'T', M, N, K, 1.f, dA, dBt, 0.f, dC); - verify("cuda::matmul NT"); - } - - // ---- 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.addLayoutConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, 'T', - 'N'); - std::fill(ref.begin(), ref.end(), 0.f); - refGemm(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); - blas.gemm('T', 'N', M, N, K, 1.f, dAt, dB, 0.f, dBias, dC); - verify("cuda::gemm TN"); - } - - // ---- gemmrelu: all-positive (relu is identity) ---- - { - auto hAp = alpaka::allocBuf(hostDev, static_cast(M * K)); - auto hBp = alpaka::allocBuf(hostDev, static_cast(K * N)); - auto hBiasz = - alpaka::allocBuf(hostDev, static_cast(M * N)); - float *Ap = alpaka::getPtrNative(hAp); - float *Bp = alpaka::getPtrNative(hBp); - fillSeq(Ap, M * K, 0.1f, 0.1f); - fillSeq(Bp, K * N, 0.1f, 0.1f); - fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); - auto dAp = - alpaka::allocAsyncBuf(queue, static_cast(M * K)); - auto dBp = - alpaka::allocAsyncBuf(queue, static_cast(K * N)); - auto dBiasz = - alpaka::allocAsyncBuf(queue, static_cast(M * N)); - alpaka::memcpy(queue, dAp, hAp); - alpaka::memcpy(queue, dBp, hBp); - alpaka::memcpy(queue, dBiasz, hBiasz); - alpaka::wait(queue); - blas.addLayoutConfig(M, N, K, M, K, M, 'N', 'N'); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmRelu(ref.data(), Ap, Bp, alpaka::getPtrNative(hBiasz), M, N, K, 1.f, - 0.f, false, false); - blas.gemmrelu('N', 'N', M, N, K, 1.f, dAp, dBp, 0.f, dBiasz, dC); - verify("cuda::gemmrelu all-positive"); - } - - // ---- 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"); - } -} - -#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.addLayoutConfig(M, N, K, ldaFor('N', M, K), ldbFor('N', K, N), M, 'N', - 'N'); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), A, B, M, N, K, 1.f, 0.f, false, false); - blas.matmul('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dC); - verify("hip::matmul NN"); - - // ---- matmul TN ---- - { - auto hAt = alpaka::allocBuf(hostDev, static_cast(K * M)); - float *At = alpaka::getPtrNative(hAt); - fillSeq(At, K * M); - auto dAt = - alpaka::allocAsyncBuf(queue, static_cast(K * M)); - alpaka::memcpy(queue, dAt, hAt); - alpaka::wait(queue); - blas.addLayoutConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, 'T', - 'N'); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), At, B, M, N, K, 1.f, 0.f, true, false); - blas.matmul('T', 'N', M, N, K, 1.f, dAt, dB, 0.f, dC); - verify("hip::matmul TN"); - } - - // ---- matmul NT ---- - { - auto hBt = alpaka::allocBuf(hostDev, static_cast(N * K)); - float *Bt = alpaka::getPtrNative(hBt); - fillSeq(Bt, N * K, 1.f, 0.5f); - auto dBt = - alpaka::allocAsyncBuf(queue, static_cast(N * K)); - alpaka::memcpy(queue, dBt, hBt); - alpaka::wait(queue); - blas.addLayoutConfig(M, N, K, ldaFor('N', M, K), ldbFor('T', K, N), M, 'N', - 'T'); - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), A, Bt, M, N, K, 1.f, 0.f, false, true); - blas.matmul('N', 'T', M, N, K, 1.f, dA, dBt, 0.f, dC); - verify("hip::matmul NT"); - } - - // ---- matmul alpha=2.5 ---- - std::fill(ref.begin(), ref.end(), 0.f); - refMatmul(ref.data(), A, B, M, N, K, 2.5f, 0.f, false, false); - blas.matmul('N', 'N', M, N, K, 2.5f, dA, dB, 0.f, dC); - verify("hip::matmul alpha=2.5"); - - // ---- gemm NN beta=0 ---- - fillSeq(bias, M * N, 0.1f, 0.1f); - alpaka::memcpy(queue, dBias, hBias); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - blas.gemm('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); - verify("hip::gemm NN beta=0"); - - // ---- gemm NN beta=1 ---- - // D_in = bias, so result = A*B + 1*bias_matrix + bias_vec - std::fill(ref.begin(), ref.end(), 0.f); - refGemm(ref.data(), A, B, bias, M, N, K, 1.f, 1.f, false, false); - blas.gemm('N', 'N', M, N, K, 1.f, dA, dB, 1.f, dBias, dC); - verify("hip::gemm NN beta=1"); - - // ---- gemm TN ---- - { - auto hAt = alpaka::allocBuf(hostDev, static_cast(K * M)); - float *At = alpaka::getPtrNative(hAt); - fillSeq(At, K * M); - auto dAt = - alpaka::allocAsyncBuf(queue, static_cast(K * M)); - alpaka::memcpy(queue, dAt, hAt); - alpaka::wait(queue); - blas.addLayoutConfig(M, N, K, ldaFor('T', M, K), ldbFor('N', K, N), M, 'T', - 'N'); - std::fill(ref.begin(), ref.end(), 0.f); - refGemm(ref.data(), At, B, bias, M, N, K, 1.f, 0.f, true, false); - blas.gemm('T', 'N', M, N, K, 1.f, dAt, dB, 0.f, dBias, dC); - verify("hip::gemm TN"); - } - - // ---- gemmrelu: all-positive ---- - { - auto hAp = alpaka::allocBuf(hostDev, static_cast(M * K)); - auto hBp = alpaka::allocBuf(hostDev, static_cast(K * N)); - auto hBiasz = - alpaka::allocBuf(hostDev, static_cast(M * N)); - float *Ap = alpaka::getPtrNative(hAp); - float *Bp = alpaka::getPtrNative(hBp); - fillSeq(Ap, M * K, 0.1f, 0.1f); - fillSeq(Bp, K * N, 0.1f, 0.1f); - fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); - auto dAp = - alpaka::allocAsyncBuf(queue, static_cast(M * K)); - auto dBp = - alpaka::allocAsyncBuf(queue, static_cast(K * N)); - auto dBiasz = - alpaka::allocAsyncBuf(queue, static_cast(M * N)); - alpaka::memcpy(queue, dAp, hAp); - alpaka::memcpy(queue, dBp, hBp); - alpaka::memcpy(queue, dBiasz, hBiasz); - alpaka::wait(queue); - blas.addLayoutConfig(M, N, K, M, K, M, 'N', 'N'); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmRelu(ref.data(), Ap, Bp, alpaka::getPtrNative(hBiasz), M, N, K, 1.f, - 0.f, false, false); - blas.gemmrelu('N', 'N', M, N, K, 1.f, dAp, dBp, 0.f, dBiasz, dC); - verify("hip::gemmrelu all-positive"); - } - - // ---- gemmrelu: alpha=-1 forces negatives ---- - { - auto hBiasz = - alpaka::allocBuf(hostDev, static_cast(M * N)); - fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f); - auto dBiasz = - alpaka::allocAsyncBuf(queue, static_cast(M * N)); - alpaka::memcpy(queue, dBiasz, hBiasz); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmRelu(ref.data(), A, B, alpaka::getPtrNative(hBiasz), M, N, K, -1.f, - 0.f, false, false); - blas.gemmrelu('N', 'N', M, N, K, -1.f, dA, dB, 0.f, dBiasz, dC); - verify("hip::gemmrelu alpha=-1 (clamped)"); - } - - // ---- gemmrelu with mixed bias ---- - fillSeq(bias, M * N, -5.f, 2.f); - alpaka::memcpy(queue, dBias, hBias); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmRelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - blas.gemmrelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); - verify("hip::gemmrelu with mixed bias"); - - // ---- gemmgelu NN ---- - fillVal(bias, M * N, 0.f); - alpaka::memcpy(queue, dBias, hBias); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - blas.gemmgelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); - verify("hip::gemmgelu NN"); - - // ---- gemmgelu with bias ---- - fillSeq(bias, M * N, -2.f, 0.5f); - alpaka::memcpy(queue, dBias, hBias); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - refGemmGelu(ref.data(), A, B, bias, M, N, K, 1.f, 0.f, false, false); - blas.gemmgelu('N', 'N', M, N, K, 1.f, dA, dB, 0.f, dBias, dC); - verify("hip::gemmgelu with bias"); - - // ---- edge: zero A ---- - { - auto hZero = alpaka::allocBuf(hostDev, static_cast(M * K)); - fillVal(alpaka::getPtrNative(hZero), M * K, 0.f); - auto dZero = - alpaka::allocAsyncBuf(queue, static_cast(M * K)); - alpaka::memcpy(queue, dZero, hZero); - alpaka::wait(queue); - std::fill(ref.begin(), ref.end(), 0.f); - blas.matmul('N', 'N', M, N, K, 1.f, dZero, dB, 0.f, dC); - verify("hip::matmul zero-A"); - } -} - -#endif // ALPAKA_ACC_GPU_HIP_ENABLED +#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) || defined(ALPAKA_ACC_GPU_HIP_ENABLED) +#include "gpu_tests.tpp" +#endif // --------------------------------------------------------------------------- // main @@ -795,10 +357,10 @@ int main() { runCpuTests(); #endif #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED - runCudaTests(); + runGpuTests("cuda"); #endif #ifdef ALPAKA_ACC_GPU_HIP_ENABLED - runHipTests(); + runGpuTests("hip"); #endif std::cout << "\n"; From 7d1aa9051819c0f25959f47dca8cb71c1e4f74cd Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Fri, 14 Aug 2026 01:05:17 +0530 Subject: [PATCH 2/6] build: detect gpu architectures and skip unavailable backends --- tests/CMakeLists.txt | 65 ++++++++++++++++++++++++++++---------------- tests/test.cc | 4 +-- 2 files changed, 43 insertions(+), 26 deletions(-) diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index d51b49a..700fbbc 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -24,16 +24,30 @@ 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) +include(CheckLanguage) + # --- CUDA architecture: must be set before enable_language(CUDA) so all # targets inherit a valid default (CMake 3.18+ requires CUDA_ARCHITECTURES -# to be non-empty on every target once CUDA is enabled globally). -set(CMAKE_CUDA_ARCHITECTURES 86) -enable_language(CUDA) +# to be non-empty on every target once CUDA is enabled globally). Defaults +# to the architecture of the GPU in the build machine; override with +# -DCMAKE_CUDA_ARCHITECTURES= to build for a different one. +if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES) + if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.24) + set(CMAKE_CUDA_ARCHITECTURES native) + else() + set(CMAKE_CUDA_ARCHITECTURES 75) + endif() +endif() +check_language(CUDA) +if(CMAKE_CUDA_COMPILER) + enable_language(CUDA) +else() + message(STATUS "No CUDA compiler found, test_cuda target will not be built") +endif() -# --- HIP architecture: must be set before enable_language(HIP) for the same -# reason as CUDA_ARCHITECTURES above. -set(CMAKE_HIP_ARCHITECTURES gfx1100) -include(CheckLanguage) +# --- HIP architecture: CMAKE_HIP_ARCHITECTURES is left unset on purpose, so +# CMake initializes it from rocm_agent_enumerator, i.e. the GPUs actually +# installed. Override with -DCMAKE_HIP_ARCHITECTURES= if needed. check_language(HIP) if(CMAKE_HIP_COMPILER) enable_language(HIP) @@ -46,10 +60,10 @@ include_directories(${ALPAKA_BASE}/include "../include") # --- Functions to find BLAS libraries --- function(find_openblas blas_lib blas_define blas_include) - find_path(OPENBLAS_PATH NAMES libopenblas.a - PATHS /usr/lib/x86_64-linux-gnu/openblas-serial /usr/lib/x86_64-linux-gnu) + find_path(OPENBLAS_PATH NAMES libopenblas.a libopenblas.so libopenblas.so.0 + PATHS /usr/lib64 /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}) + find_library(OPENBLAS_LIB NAMES openblas libopenblas.so.0 PATHS ${OPENBLAS_PATH}) find_path(OPENBLAS_INCLUDE_DIR NAMES cblas.h PATHS /usr/include /usr/include/openblas /usr/include/cblas) set(${blas_lib} ${OPENBLAS_LIB} PARENT_SCOPE) @@ -143,22 +157,25 @@ target_include_directories(test_cpu PRIVATE "../sofieBLAS/include" ${ALPAKA_BASE target_link_libraries(test_cpu PRIVATE ${BLAS_LIBS}) -add_executable(test_cuda) -target_sources(test_cuda PRIVATE test.cc) -set_source_files_properties(test.cc PROPERTIES LANGUAGE CUDA) -set_target_properties(test_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON) -target_compile_features(test_cuda PUBLIC cxx_std_20) +# --- test_cuda target --- +if(CMAKE_CUDA_COMPILER) + add_executable(test_cuda) + target_sources(test_cuda PRIVATE test.cc) + set_source_files_properties(test.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_options(test_cuda PRIVATE + ${CXXFLAGS} + ${CXX_CUDA_FLAGS} + ${XCOMPILER_FLAGS} + ) -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 cublasLt cublas cudart) + 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 cublasLt cublas cudart) +endif() # --- test_hip target --- if(CMAKE_HIP_COMPILER) diff --git a/tests/test.cc b/tests/test.cc index da88625..fa2be9a 100644 --- a/tests/test.cc +++ b/tests/test.cc @@ -357,10 +357,10 @@ int main() { runCpuTests(); #endif #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED - runGpuTests("cuda"); + runGpuTests("CUDA"); #endif #ifdef ALPAKA_ACC_GPU_HIP_ENABLED - runGpuTests("hip"); + runGpuTests("HIP"); #endif std::cout << "\n"; From 2405df01b36d19fbae6148ef5795a232b6f451d7 Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Fri, 11 Sep 2026 02:34:41 +0530 Subject: [PATCH 3/6] tests: move cpu and gpu tests into unit_test.tpp files and share the dynamic shape test --- tests/cpu/unit_test.tpp | 220 +++++++++++ tests/{gpu_tests.tpp => gpu/unit_test.tpp} | 97 +++++ tests/test.cc | 423 +-------------------- 3 files changed, 321 insertions(+), 419 deletions(-) create mode 100644 tests/cpu/unit_test.tpp rename tests/{gpu_tests.tpp => gpu/unit_test.tpp} (67%) 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_tests.tpp b/tests/gpu/unit_test.tpp similarity index 67% rename from tests/gpu_tests.tpp rename to tests/gpu/unit_test.tpp index c49157b..9b04003 100644 --- a/tests/gpu_tests.tpp +++ b/tests/gpu/unit_test.tpp @@ -219,3 +219,100 @@ static void runGpuTests(const std::string &backend) { verify(backend + "::matmul zero-A"); } } + +template +static void runGpuDynamicShapeTests(const std::string &backend) { + std::cout << "\n=== " << backend << " Dynamic-Shape Tests ===\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, backend + "::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, backend + "::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 " << backend << "::cache limit honoured\n"; + } else { + std::cerr << " FAIL [" << backend << "::cache limit honoured] " + << capped.algoCacheSize() << " entries, worst err " << worst + << "\n"; + ++gFailures; + } + } +} + diff --git a/tests/test.cc b/tests/test.cc index def820e..d69a6ea 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,207 +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_tests.tpp" +#include "gpu/unit_test.tpp" #endif -// --------------------------------------------------------------------------- -// CUDA tests -// --------------------------------------------------------------------------- - -#ifdef ALPAKA_ACC_GPU_CUDA_ENABLED - -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 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 // --------------------------------------------------------------------------- @@ -549,11 +134,11 @@ int main() { #endif #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED runGpuTests("CUDA"); - runDynamicShapeTests(); + runGpuDynamicShapeTests("CUDA"); #endif #ifdef ALPAKA_ACC_GPU_HIP_ENABLED runGpuTests("HIP"); - runHipDynamicShapeTests(); + runGpuDynamicShapeTests("HIP"); #endif std::cout << "\n"; From 6e9728f6bfb19aa8ddc407099a14f48101a83b3f Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Tue, 15 Sep 2026 14:40:15 +0530 Subject: [PATCH 4/6] tests: identify the gpu backend from the template tag instead of a name argument --- tests/gpu/unit_test.tpp | 44 ++++++++++++++++++++--------------------- tests/test.cc | 8 ++++---- 2 files changed, 25 insertions(+), 27 deletions(-) diff --git a/tests/gpu/unit_test.tpp b/tests/gpu/unit_test.tpp index 9b04003..614b23a 100644 --- a/tests/gpu/unit_test.tpp +++ b/tests/gpu/unit_test.tpp @@ -1,10 +1,8 @@ // shared test for CUDA/HIP -template - -static void runGpuTests(const std::string &backend) { - std::cout << "\n=== " << backend << " Tests ===\n"; +template static void runGpuTests() { + std::cout << "\n=== " << __PRETTY_FUNCTION__ << " ===\n"; using Acc = alpaka::TagToAcc; using DevAcc = alpaka::Dev; @@ -59,7 +57,7 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::matmul NN"); + verify("matmul NN"); // ---- matmul TN ---- { @@ -75,7 +73,7 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::matmul TN"); + verify("matmul TN"); } // ---- matmul NT ---- @@ -92,14 +90,14 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::matmul NT"); + 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(backend + "::matmul alpha=2.5"); + verify("matmul alpha=2.5"); // ---- gemm NN beta=0 ---- fillSeq(bias, M * N, 0.1f, 0.1f); @@ -108,14 +106,14 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::gemm NN beta=0"); + 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(backend + "::gemm NN beta=1"); + verify("gemm NN beta=1"); // ---- gemm TN ---- { @@ -131,7 +129,7 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::gemm TN"); + verify("gemm TN"); } // ---- gemmrelu: all-positive (relu is identity) ---- @@ -160,7 +158,7 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::gemmrelu all-positive"); + verify("gemmrelu all-positive"); } // ---- gemmrelu: alpha=-1 forces negatives -> clamped to zero ---- @@ -176,7 +174,7 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::gemmrelu alpha=-1 (clamped)"); + verify("gemmrelu alpha=-1 (clamped)"); } // ---- gemmrelu with mixed bias ---- @@ -186,7 +184,7 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::gemmrelu with mixed bias"); + verify("gemmrelu with mixed bias"); // ---- gemmgelu NN ---- fillVal(bias, M * N, 0.f); @@ -195,7 +193,7 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::gemmgelu NN"); + verify("gemmgelu NN"); // ---- gemmgelu with bias ---- fillSeq(bias, M * N, -2.f, 0.5f); @@ -204,7 +202,7 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::gemmgelu with bias"); + verify("gemmgelu with bias"); // ---- edge: zero A ---- { @@ -216,13 +214,13 @@ static void runGpuTests(const std::string &backend) { 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(backend + "::matmul zero-A"); + verify("matmul zero-A"); } } template -static void runGpuDynamicShapeTests(const std::string &backend) { - std::cout << "\n=== " << backend << " Dynamic-Shape Tests ===\n"; +static void runGpuDynamicShapeTests() { + std::cout << "\n=== " << __PRETTY_FUNCTION__ << " ===\n"; using Acc = alpaka::TagToAcc; using DevAcc = alpaka::Dev; @@ -275,7 +273,7 @@ static void runGpuDynamicShapeTests(const std::string &backend) { }; for (int m : {M0, 37, 8, 51, 1, M0, MCAP}) - runAt(m, backend + "::dynamic m=" + std::to_string(m)); + 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. @@ -286,7 +284,7 @@ static void runGpuDynamicShapeTests(const std::string &backend) { alpaka::getPtrNative(dC)); alpaka::memcpy(queue, hC, dC); alpaka::wait(queue); - checkClose(C, ref.data(), 45 * N, backend + "::dynamic raw pointers m=45"); + checkClose(C, ref.data(), 45 * N, "dynamic raw pointers m=45"); // 32 distinct sizes through a cache limited to 8 entries. { @@ -306,9 +304,9 @@ static void runGpuDynamicShapeTests(const std::string &backend) { worst = std::max(worst, std::abs(C[i] - ref[i])); } if (capped.algoCacheSize() <= 8 && worst < 1e-3f) { - std::cout << " PASS " << backend << "::cache limit honoured\n"; + std::cout << " PASS cache limit honoured\n"; } else { - std::cerr << " FAIL [" << backend << "::cache limit honoured] " + 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 d69a6ea..9e2ce01 100644 --- a/tests/test.cc +++ b/tests/test.cc @@ -133,12 +133,12 @@ int main() { runCpuTests(); #endif #ifdef ALPAKA_ACC_GPU_CUDA_ENABLED - runGpuTests("CUDA"); - runGpuDynamicShapeTests("CUDA"); + runGpuTests(); + runGpuDynamicShapeTests(); #endif #ifdef ALPAKA_ACC_GPU_HIP_ENABLED - runGpuTests("HIP"); - runGpuDynamicShapeTests("HIP"); + runGpuTests(); + runGpuDynamicShapeTests(); #endif std::cout << "\n"; From e8dd17b75ad7a798fca961ead08e27d6231e9bce Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Tue, 15 Sep 2026 15:01:15 +0530 Subject: [PATCH 5/6] ci: run the cpu tests against both OpenBLAS and MKL --- .github/workflows/tests.yml | 32 +++++++++++++++++++++++++------- 1 file changed, 25 insertions(+), 7 deletions(-) diff --git a/.github/workflows/tests.yml b/.github/workflows/tests.yml index e41ae81..ef8a3a6 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,24 @@ 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 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 From 4730fda2cf374bc0b1ccefb32649f05fda5b7232 Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Tue, 15 Sep 2026 15:12:51 +0530 Subject: [PATCH 6/6] ci: use the sequential MKL threading layer so libmkl_rt needs no OpenMP runtime --- .github/workflows/tests.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/workflows/tests.yml b/.github/workflows/tests.yml index ef8a3a6..e14c3f2 100644 --- a/.github/workflows/tests.yml +++ b/.github/workflows/tests.yml @@ -83,6 +83,7 @@ jobs: 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