Skip to content

[tests] remove cuda/hip test duplication with a templated fixture #13

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
harz05 wants to merge 2 commits into
base: dev
Choose a base branch
from
Open

[tests] remove cuda/hip test duplication with a templated fixture #13

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
248fdf0
test: deduplicate cuda and hip tests into a tag templated fixture test
harz05 Aug 13, 2026
7d1aa90
build: detect gpu architectures and skip unavailable backends
harz05 Aug 13, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
67 changes: 42 additions & 25 deletions tests/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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=<arch> 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=<arch> if needed.
check_language(HIP)
if(CMAKE_HIP_COMPILER)
enable_language(HIP)
Expand All @@ -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)
Expand Down Expand Up @@ -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)
Expand All @@ -177,7 +194,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()
Expand Down
221 changes: 221 additions & 0 deletions tests/gpu_tests.tpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,221 @@

// shared test for CUDA/HIP

template <typename TTag>

static void runGpuTests(const std::string &backend) {
std::cout << "\n=== " << backend << " Tests ===\n";

using Acc = alpaka::TagToAcc<TTag, Dim1D, Idx>;
using DevAcc = alpaka::Dev<Acc>;
using PlatformAcc = alpaka::Platform<Acc>;

PlatformAcc platform{};
auto dev = alpaka::getDevByIdx(platform, 0u);
alpaka::Queue<DevAcc, alpaka::NonBlocking> queue{dev};
sofieBLAS<TTag> blas(queue);

alpaka::PlatformCpu hostPlatform{};
auto hostDev = alpaka::getDevByIdx(hostPlatform, 0u);

constexpr int M = 4, N = 3, K = 5;

auto hA = alpaka::allocBuf<float, Idx>(hostDev, static_cast<Idx>(M * K));
auto hB = alpaka::allocBuf<float, Idx>(hostDev, static_cast<Idx>(K * N));
auto hC = alpaka::allocBuf<float, Idx>(hostDev, static_cast<Idx>(M * N));
auto hBias = alpaka::allocBuf<float, Idx>(hostDev, static_cast<Idx>(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<float, Idx>(queue, static_cast<Idx>(M * K));
auto dB = alpaka::allocAsyncBuf<float, Idx>(queue, static_cast<Idx>(K * N));
auto dC = alpaka::allocAsyncBuf<float, Idx>(queue, static_cast<Idx>(M * N));
auto dBias =
alpaka::allocAsyncBuf<float, Idx>(queue, static_cast<Idx>(M * N));

alpaka::memcpy(queue, dA, hA);
alpaka::memcpy(queue, dB, hB);
alpaka::memcpy(queue, dBias, hBias);
alpaka::wait(queue);

std::vector<float> 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<float, Idx>(hostDev, static_cast<Idx>(K * M));
float *At = alpaka::getPtrNative(hAt);
fillSeq(At, K * M);
auto dAt =
alpaka::allocAsyncBuf<float, Idx>(queue, static_cast<Idx>(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<float, Idx>(hostDev, static_cast<Idx>(N * K));
float *Bt = alpaka::getPtrNative(hBt);
fillSeq(Bt, N * K, 1.f, 0.5f);
auto dBt =
alpaka::allocAsyncBuf<float, Idx>(queue, static_cast<Idx>(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<float, Idx>(hostDev, static_cast<Idx>(K * M));
float *At = alpaka::getPtrNative(hAt);
fillSeq(At, K * M);
auto dAt =
alpaka::allocAsyncBuf<float, Idx>(queue, static_cast<Idx>(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<float, Idx>(hostDev, static_cast<Idx>(M * K));
auto hBp = alpaka::allocBuf<float, Idx>(hostDev, static_cast<Idx>(K * N));
auto hBiasz =
alpaka::allocBuf<float, Idx>(hostDev, static_cast<Idx>(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<float, Idx>(queue, static_cast<Idx>(M * K));
auto dBp =
alpaka::allocAsyncBuf<float, Idx>(queue, static_cast<Idx>(K * N));
auto dBiasz =
alpaka::allocAsyncBuf<float, Idx>(queue, static_cast<Idx>(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<float, Idx>(hostDev, static_cast<Idx>(M * N));
fillVal(alpaka::getPtrNative(hBiasz), M * N, 0.f);
auto dBiasz =
alpaka::allocAsyncBuf<float, Idx>(queue, static_cast<Idx>(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<float, Idx>(hostDev, static_cast<Idx>(M * K));
fillVal(alpaka::getPtrNative(hZero), M * K, 0.f);
auto dZero =
alpaka::allocAsyncBuf<float, Idx>(queue, static_cast<Idx>(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");
}
}
Loading