From 248fdf028cd57188f891c064c801442fd900e2e1 Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Thu, 13 Aug 2026 23:43:36 +0530 Subject: [PATCH 1/2] 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/2] 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";