This document describes the Catch-based test framework for HIP functionality in the LLVM Test Suite.
The HIP Catch Tests framework provides comprehensive HIP tests using the Catch2 testing framework. The framework includes:
Currently included test categories:
unit/compiler: Tests for HIP compiler features (kernels, device code compilation)ROCm installation: A valid ROCm installation with HIP support
LLVM Test Suite: This repository contains the Catch test infrastructure in External/HIP/catch/
clang++: ROCm's clang++ compiler for building HIP code
Catch2 v2.13.10+: Obtained automatically via:
Configure the test suite with Catch tests enabled:
cmake -G Ninja \ -DENABLE_HIP_CATCH_TESTS=ON \ -DTEST_SUITE_EXTERNALS_DIR=/path/to/externals \ -DCMAKE_CXX_COMPILER=/path/to/clang++ \ -DCMAKE_C_COMPILER=/path/to/clang \ -DAMDGPU_ARCHS=gfx90a \ /path/to/llvm-test-suite
Note: Catch tests are disabled by default. Set -DENABLE_HIP_CATCH_TESTS=ON to enable. Catch2 is automatically obtained via find_package (system) or FetchContent (download).
The Catch tests are platform-agnostic and support both AMD and NVIDIA GPUs through HIP:
Use AMD's clang++ compiler directly:
cmake -G Ninja \ -DENABLE_HIP_CATCH_TESTS=ON \ -DTEST_SUITE_EXTERNALS_DIR=/path/to/externals \ -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ \ -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang \ -DAMDGPU_ARCHS=gfx90a \ /path/to/llvm-test-suite
Use HIP's hipcc wrapper (built with CUDA backend):
cmake -G Ninja \ -DENABLE_HIP_CATCH_TESTS=ON \ -DTEST_SUITE_EXTERNALS_DIR=/path/to/externals \ -DCMAKE_CXX_COMPILER=/path/to/hip-cuda/bin/hipcc \ -DCMAKE_C_COMPILER=/path/to/hip-cuda/bin/hipcc \ -DCUDA_ARCH=sm_75 \ /path/to/llvm-test-suite
Note: The hipcc wrapper automatically handles platform-specific compile and link flags. The same test source code works on both platforms without modification!
After configuration, hierarchical targets are available at three levels: top-level (all catch tests), category-level (e.g., unit tests), and subdirectory-level (e.g., unit/compiler tests). Each level has both aggregated (all ROCm variants) and per-variant targets.
Build or run all Catch tests across all categories:
| Target | Description |
|---|---|
ninja hip-tests-catch | Build all Catch tests, all variants |
ninja hip-tests-catch-hip-7.2.0 | Build all Catch tests for specific variant |
ninja check-hip-catch | Run all Catch tests, all variants |
ninja check-hip-catch-hip-7.2.0 | Run all Catch tests for specific variant |
Build or run tests for a specific category (e.g., unit, stress):
| Target | Description |
|---|---|
ninja hip-tests-catch-unit | Build all unit tests, all variants |
ninja hip-tests-catch-unit-hip-7.2.0 | Build all unit tests for specific variant |
ninja check-hip-catch-unit | Run all unit tests, all variants |
ninja check-hip-catch-unit-hip-7.2.0 | Run all unit tests for specific variant |
Build or run tests for a specific subdirectory within a category (e.g., unit/compiler):
| Target | Description |
|---|---|
ninja hip-tests-catch-unit-compiler | Build unit/compiler tests, all variants |
ninja hip-tests-catch-unit-compiler-hip-7.2.0 | Build unit/compiler tests for specific variant |
ninja check-hip-catch-unit-compiler | Run unit/compiler tests, all variants |
ninja check-hip-catch-unit-compiler-hip-7.2.0 | Run unit/compiler tests for specific variant |
| Target | Description |
|---|---|
ninja hip-tests-all | Build all HIP tests (including Catch tests if enabled) |
Type: Boolean Default: OFF Description: Master switch to enable or disable the entire HIP Catch test framework. When disabled, no Catch test targets are created, and the build only includes simple HIP tests.
When ENABLED (default):
hip-tests-catch, hip-tests-catch-unit, hip-tests-catch-unit-compilerhip-tests-all includes both simple and catch testsWhen DISABLED:
hip-tests-all includes only simple testsExamples:
# Enable Catch tests (default) -DENABLE_HIP_CATCH_TESTS=ON # Disable Catch tests (only simple tests) -DENABLE_HIP_CATCH_TESTS=OFF
Use cases for disabling:
Type: Boolean (ON/OFF) Default: OFF Description: Controls output verbosity level for Catch2 test runs
Both modes use unified LIT-based execution (metrics always collected). The difference is only in output detail:
.test.out files-a flag)Example:
# Enable verbose output (shows full Catch2 output) -DHIP_CATCH_TEST_VERBOSE=ON # Quiet mode (default - only shows LIT summary) -DHIP_CATCH_TEST_VERBOSE=OFF
Output Comparison:
With HIP_CATCH_TEST_VERBOSE=OFF (default):
Testing Time: 0.06s Total Discovered Tests: 3 Passed: 3 (100.00%) ======================================== Detailed Test Summary: Test Suites: 3 Total Tests: 10 Passed: 8 Failed: 0 Skipped: 2 Skipped Tests: - Unit_test_generic_target_in_regular_fatbin [catch_unit_compiler_hipSquareGenericTarget-hip-7.2.0] - Unit_test_generic_target_only_in_regular_fatbin [catch_unit_compiler_hipSquareGenericTarget-hip-7.2.0] ========================================
With HIP_CATCH_TEST_VERBOSE=ON:
PASS: test-suite :: External/HIP/catch_unit_compiler_hipSquare-hip-7.2.0.test (1 of 3) ... full command and output shown ... PASS: test-suite :: External/HIP/catch_unit_compiler_hipClassKernel-hip-7.2.0.test (2 of 3) ... full command and output shown ... PASS: test-suite :: External/HIP/catch_unit_compiler_hipSquareGenericTarget-hip-7.2.0.test (3 of 3) ... full command and output shown ... Testing Time: 0.06s Total Discovered Tests: 3 Passed: 3 (100.00%) ======================================== Detailed Test Summary: Test Suites: 3 Total Tests: 10 Passed: 8 Failed: 0 Skipped: 2 Skipped Tests: - Unit_test_generic_target_in_regular_fatbin [catch_unit_compiler_hipSquareGenericTarget-hip-7.2.0] - Unit_test_generic_target_only_in_regular_fatbin [catch_unit_compiler_hipSquareGenericTarget-hip-7.2.0] ========================================
Summary Categories:
Triage Lists: The summary includes lists of failed, skipped, and crashed tests by name for easier debugging:
======================================== Detailed Test Summary: Test Suites: 3 Total Tests: 10 Passed: 0 Failed: 2 Skipped: 2 Crashed/Error: 6 Failed Tests: - Unit_hipClassKernel_Overload_Override [catch_unit_compiler_hipClassKernel-hip-7.2.0] - Unit_test_compressed_codeobject [catch_unit_compiler_hipSquare-hip-7.2.0] Skipped Tests: - Unit_test_generic_target_in_regular_fatbin [catch_unit_compiler_hipSquareGenericTarget-hip-7.2.0] - Unit_test_generic_target_only_in_regular_fatbin [catch_unit_compiler_hipSquareGenericTarget-hip-7.2.0] Crashed/Error Tests: - Unit_hipClassKernel_Friend [catch_unit_compiler_hipClassKernel-hip-7.2.0] - Unit_hipClassKernel_Empty [catch_unit_compiler_hipClassKernel-hip-7.2.0] - Unit_hipClassKernel_BSize [catch_unit_compiler_hipClassKernel-hip-7.2.0] - Unit_hipClassKernel_Size [catch_unit_compiler_hipClassKernel-hip-7.2.0] - Unit_hipClassKernel_Virtual [catch_unit_compiler_hipClassKernel-hip-7.2.0] - Unit_hipClassKernel_Value [catch_unit_compiler_hipClassKernel-hip-7.2.0] ========================================
Note: The verbose TEST_CASE output appears in the terminal during the test run. It's also saved to .out files in build/External/HIP/Output/ for later analysis.
Type: Semicolon-separated list Default: “unit” Description: Test categories to include in the build
Available categories:
unit: Unit tests for core HIP APIs (currently: compiler tests)Example:
# Include unit tests (default and currently only available category) -DCATCH_TEST_CATEGORIES="unit"
Type: Semicolon-separated list Default: "" (empty - includes all subdirectories) Description: Specific subdirectories to include within enabled categories. When empty, all subdirectories are automatically discovered and included.
This option provides fine-grained control over which tests to build. For example, within the unit category, you can selectively include only compiler, memory, and stream tests while excluding others.
Examples:
# Include only compiler tests from enabled categories -DCATCH_TEST_SUBDIRS="compiler" # Include multiple specific subdirectories -DCATCH_TEST_SUBDIRS="compiler;memory;stream;kernel" # Include all subdirectories (default behavior) -DCATCH_TEST_SUBDIRS="" # or simply omit the option
Note: The subdirectory filter applies to all enabled categories. For example, if you set -DCATCH_TEST_CATEGORIES="unit;stress" and -DCATCH_TEST_SUBDIRS="compiler", the framework will include:
catch/unit/compiler (if it exists)catch/stress/compiler (if it exists)Available subdirectories in the unit category include:
compiler - Compiler-specific testsmemory - Memory management testsstream - Stream management testsevent - Event handling testskernel - Kernel execution testsdevice - Device management testsmath - Math operations testsType: Integer (seconds) Default: 60 Description: Timeout for individual Catch tests
-DHIP_CATCH_TEST_TIMEOUT=120
#!/bin/bash export CLANG_DIR=/opt/rocm-7.2.0/llvm export ROCM_PATH=/opt/rocm-7.2.0 export AMDGPU_ARCHS=gfx90a cmake -G Ninja \ -DENABLE_HIP_CATCH_TESTS=ON \ -DTEST_SUITE_EXTERNALS_DIR=${ROCM_PATH} \ -DCMAKE_CXX_COMPILER="${CLANG_DIR}/bin/clang++" \ -DCMAKE_C_COMPILER="${CLANG_DIR}/bin/clang" \ -DAMDGPU_ARCHS=${AMDGPU_ARCHS} \ /path/to/llvm-test-suite
Build only specific subdirectories (e.g., compiler tests):
#!/bin/bash export CLANG_DIR=/opt/rocm-7.2.0/llvm export ROCM_PATH=/opt/rocm-7.2.0 export AMDGPU_ARCHS=gfx90a cmake -G Ninja \ -DENABLE_HIP_CATCH_TESTS=ON \ -DTEST_SUITE_EXTERNALS_DIR=${ROCM_PATH} \ -DCMAKE_CXX_COMPILER="${CLANG_DIR}/bin/clang++" \ -DCMAKE_C_COMPILER="${CLANG_DIR}/bin/clang" \ -DAMDGPU_ARCHS=${AMDGPU_ARCHS} \ -DCATCH_TEST_CATEGORIES="unit" \ -DCATCH_TEST_SUBDIRS="compiler" \ /path/to/llvm-test-suite # Build compiler tests ninja hip-tests-catch-unit-compiler # Run compiler tests ninja check-hip-catch-unit-compiler
#!/bin/bash export CLANG_DIR=/opt/rocm-7.2.0/llvm export EXTERNAL_DIR=/opt/rocm-7.2.0 export AMDGPU_ARCHS=gfx90a export TEST_SUITE_DIR=/path/to/llvm-test-suite cd /path/to/build-llvm-test-suite PATH=${CLANG_DIR}/bin:$PATH \ CXX=clang++ \ CC=clang \ cmake -G Ninja \ -DENABLE_HIP_CATCH_TESTS=ON \ -DTEST_SUITE_EXTERNALS_DIR=${EXTERNAL_DIR} \ -DCMAKE_CXX_COMPILER="${CLANG_DIR}/bin/clang++" \ -DCMAKE_C_COMPILER="${CLANG_DIR}/bin/clang" \ -DAMDGPU_ARCHS=${AMDGPU_ARCHS} \ ${TEST_SUITE_DIR} # Build all tests ninja hip-tests-all # Run simple tests ninja check-hip-simple # Run Catch tests at different levels ninja check-hip-catch # All catch tests ninja check-hip-catch-unit # All unit tests ninja check-hip-catch-unit-compiler # Only unit/compiler tests
For quick builds or when you only need simple HIP tests:
#!/bin/bash export CLANG_DIR=/opt/rocm-7.2.0/llvm export ROCM_PATH=/opt/rocm-7.2.0 export AMDGPU_ARCHS=gfx90a cmake -G Ninja \ -DENABLE_HIP_CATCH_TESTS=OFF \ -DTEST_SUITE_EXTERNALS_DIR=${ROCM_PATH} \ -DCMAKE_CXX_COMPILER="${CLANG_DIR}/bin/clang++" \ -DCMAKE_C_COMPILER="${CLANG_DIR}/bin/clang" \ -DAMDGPU_ARCHS=${AMDGPU_ARCHS} \ /path/to/llvm-test-suite # With Catch tests disabled, only simple test targets available ninja hip-tests-all # Build all simple tests ninja check-hip-simple # Run all simple tests ninja hip-tests-simple-hip-7.2.0 # Build simple tests for specific variant ninja check-hip-simple-hip-7.2.0 # Run simple tests for specific variant # Catch test targets are NOT created (will error if you try to use them) # ninja hip-tests-catch # ERROR: target not found # ninja check-hip-catch-unit # ERROR: target not found
The framework creates a three-level hierarchical structure for organizing and running tests (when ENABLE_HIP_CATCH_TESTS=ON):
Note: When ENABLE_HIP_CATCH_TESTS=OFF, NO Catch test targets are created at all. Only simple HIP test targets are available.
Level 1: Top-Level (ALL catch tests) ├── hip-tests-catch (all categories, all variants) ├── hip-tests-catch-hip-7.2.0 (all categories, one variant) │ Level 2: Category (e.g., unit tests) ├── hip-tests-catch-unit (all unit subdirs, all variants) ├── hip-tests-catch-unit-hip-7.2.0 (all unit subdirs, one variant) │ Level 3: Subdirectory (e.g., unit/compiler tests) ├── hip-tests-catch-unit-compiler (compiler tests, all variants) └── hip-tests-catch-unit-compiler-hip-7.2.0 (compiler tests, one variant)
Scenario 1: Compiler developer working on code generation
# Quick iteration: test only compiler tests for one ROCm version ninja hip-tests-catch-unit-compiler-hip-7.2.0 ninja check-hip-catch-unit-compiler-hip-7.2.0
Scenario 2: Testing a change across all ROCm versions
# Build compiler tests for all installed ROCm versions ninja hip-tests-catch-unit-compiler ninja check-hip-catch-unit-compiler
Scenario 3: Pre-commit validation
# Run all unit tests across all variants ninja check-hip-catch-unit
Scenario 4: Nightly CI testing
# Run everything ninja check-hip-catch
The framework organizes tests by:
ROCm Version: Each ROCm installation gets its own test variant
hip-7.2.0, hip-6.2.0Category: Tests are grouped by category
unit: Core API functionalitystress: High-load scenariosperformance: Benchmarkingperftests: Detailed metricsSubdirectory: Each subdirectory within a category
unit/compiler, unit/memory, unit/stream, etc.Test Executable: Each subdirectory produces a test executable
catch_unit_compiler-hip-7.2.0# Configure to include only unit tests cmake -DCATCH_TEST_CATEGORIES="unit" ... # Build all unit tests (all variants) ninja hip-tests-catch-unit # Build all unit tests for specific variant ninja hip-tests-catch-unit-hip-7.2.0
# Configure to include specific subdirectories cmake -DCATCH_TEST_SUBDIRS="compiler;memory" ... # Build unit/compiler tests (all variants) ninja hip-tests-catch-unit-compiler # Build unit/compiler tests for specific variant ninja hip-tests-catch-unit-compiler-hip-7.2.0
# Build all catch tests (all categories, all variants) ninja hip-tests-catch # Build all catch tests for specific variant ninja hip-tests-catch-hip-7.2.0
# Run all catch tests (all categories, all variants) ninja check-hip-catch # Run all catch tests for specific variant ninja check-hip-catch-hip-7.2.0
# Run all unit tests (all variants) ninja check-hip-catch-unit # Run all unit tests for specific variant ninja check-hip-catch-unit-hip-7.2.0
# Run unit/compiler tests (all variants) ninja check-hip-catch-unit-compiler # Run unit/compiler tests for specific variant ninja check-hip-catch-unit-compiler-hip-7.2.0
Each .cc source file is compiled into a separate test executable, allowing for individual test execution:
# Run a specific test executable directly ./External/HIP/catch_tests/catch_unit_compiler_hipSquare-hip-7.2.0 ./External/HIP/catch_tests/catch_unit_compiler_hipClassKernel-hip-7.2.0 # Run with Catch2 filtering to select specific test cases ./External/HIP/catch_tests/catch_unit_compiler_hipSquare-hip-7.2.0 "[tag]" ./External/HIP/catch_tests/catch_unit_compiler_hipClassKernel-hip-7.2.0 "Unit_hipClassKernel_*"
The tests are integrated with the LIT test runner. Each .cc source file creates a separate LIT test:
# Run all HIP tests lit External/HIP # Run only Catch tests (filter by name pattern) lit -a External/HIP | grep catch_ # Run a specific test file llvm-lit catch_unit_compiler_hipSquare-hip-7.2.0.test
Note: LIT reports one test per source file. For example, unit/compiler with 3 .cc files will show “Total Discovered Tests: 3” in LIT output. Each test executable may contain multiple Catch2 TEST_CASE definitions internally.
Viewing Verbose Output: When HIP_CATCH_TEST_VERBOSE=ON, LIT's -a flag is enabled, showing the full test output in the terminal after each test completes. The output is also saved to .test.out files in the Output/ directory for later analysis. Both verbose and non-verbose modes use the same LIT-based execution and always collect metrics.
Error: FetchContent failed to download Catch2
Cause: No internet connection during first configure, and Catch2 is not installed system-wide.
Solutions:
Install Catch2 system-wide (recommended for offline builds):
# Ubuntu/Debian sudo apt install catch2 # Or build from source (exact version) git clone -b v2.13.10 https://github.com/catchorg/Catch2.git cd Catch2 cmake -B build -DCMAKE_INSTALL_PREFIX=/usr/local cmake --build build sudo cmake --install build
Ensure internet connectivity during first CMake configure. Subsequent builds use the cached download.
Error: No test sources found in category/subdir
Solution: The CMakeLists.txt parser may have failed. This can happen if:
Check available test categories:
ls /path/to/hip-tests/catch/
Common issues:
AMDGPU_ARCHS matches your GPUThe framework consists of:
HipCatchTests.cmake: Main CMake module with functions for:
.cc source file)Modified CMakeLists.txt: Integration points in the main HIP test CMakeLists.txt:
create_hip_tests()create_hip_test()hip-tests-allTest Granularity: Each .cc source file is compiled into a separate test executable, allowing:
Test Wrappers: Shell scripts for LIT integration
To add support for additional test categories:
Add the category to CATCH_TEST_CATEGORIES:
-DCATCH_TEST_CATEGORIES="unit;stress;mycategory"
The framework will automatically discover subdirectories in:
llvm-test-suite/External/HIP/catch/mycategory/
.cc file is compiled into a separate executable, which can increase build time. Use specific categories or subdirectories to reduce build scope..cc source file creates a separate executable per variant. For example, 3 source files with 2 variants = 6 executables.-j flag to parallelize builds (particularly beneficial with multiple test executables):ninja -j16 hip-tests-catch
To contribute improvements to the framework:
Potential improvements:
The Catch test infrastructure is located in External/HIP/catch/:
External/HIP/catch/
├── unit/ # Unit test category
│ └── compiler/ # Compiler test subdirectory
│ ├── hipClassKernel.cc
│ ├── hipClassKernel.h
│ ├── hipSquare.cc
│ └── hipSquareGenericTarget.cc
├── include/ # Common test headers
│ ├── hip_test_common.hh
│ ├── hip_test_context.hh
│ ├── hip_test_features.hh
│ ├── hip_test_filesystem.hh
│ └── cmd_options.hh
├── hipTestMain/ # Test framework main files
│ ├── main.cc
│ ├── hip_test_context.cc
│ └── hip_test_features.cc
└── external/ # Third-party libraries
└── picojson/
└── picojson.h # JSON parser (vendored)
# Note: Catch2 is obtained via find_package or FetchContent, not vendored
Note: The kernels/ directory is not needed for unit/compiler tests. These tests define kernels inline using __global__ functions.