diff --git a/External/HIP/CATCH_TESTS_README.md b/External/HIP/CATCH_TESTS_README.md new file mode 100644 index 0000000000..c2c804499b --- /dev/null +++ b/External/HIP/CATCH_TESTS_README.md @@ -0,0 +1,766 @@ +# HIP Catch Tests Integration + +This document describes the Catch-based test framework for HIP functionality in the LLVM Test Suite. + +## Overview + +The HIP Catch Tests framework provides comprehensive HIP tests using the Catch2 testing framework. The framework includes: + +- **Unit tests**: Core HIP API functionality (memory, streams, events, kernels, compiler features, etc.) +- Support for multiple ROCm versions (per-variant testing) +- Hierarchical test organization (category and subdirectory levels) +- Integration with LLVM's LIT test runner + +Currently included test categories: +- `unit/compiler`: Tests for HIP compiler features (kernels, device code compilation) + +## Requirements + +1. **ROCm installation**: A valid ROCm installation with HIP support + +2. **LLVM Test Suite**: This repository contains the Catch test infrastructure in `External/HIP/catch/` + +3. **clang++**: ROCm's clang++ compiler for building HIP code + +4. **Catch2 v2.13.10+**: Obtained automatically via: + - **System installation**: If Catch2 >= 2.13.10 is installed, it will be used + - **FetchContent**: If not found, CMake will download Catch2 v2.13.10 from GitHub at configure time (requires internet connection for first build) + +## Quick Start + +### Basic Configuration + +Configure the test suite with Catch tests enabled: + +```bash +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). + +### Platform Support + +The Catch tests are **platform-agnostic** and support both AMD and NVIDIA GPUs through HIP: + +#### AMD ROCm Backend (Default) + +Use AMD's clang++ compiler directly: + +```bash +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 +``` + +#### NVIDIA CUDA Backend + +Use HIP's `hipcc` wrapper (built with CUDA backend): + +```bash +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! + +### Build Targets + +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. + +#### Top-Level 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 | + +#### Category-Level Targets + +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 | + +#### Subdirectory-Level Targets + +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 | + +#### Other Targets + +| Target | Description | +|--------|-------------| +| `ninja hip-tests-all` | Build all HIP tests (including Catch tests if enabled) | + +## Configuration Options + +### ENABLE_HIP_CATCH_TESTS + +**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)**: +- All Catch test targets are created (top-level, category-level, subdirectory-level, per-variant) +- Examples: `hip-tests-catch`, `hip-tests-catch-unit`, `hip-tests-catch-unit-compiler` +- `hip-tests-all` includes both simple and catch tests + +**When DISABLED**: +- No Catch test targets are created +- Only simple HIP tests are available +- `hip-tests-all` includes only simple tests +- Faster CMake configuration (no Catch test infrastructure processed) + +Examples: +```bash +# Enable Catch tests (default) +-DENABLE_HIP_CATCH_TESTS=ON + +# Disable Catch tests (only simple tests) +-DENABLE_HIP_CATCH_TESTS=OFF +``` + +**Use cases for disabling**: +- Quick builds when you only need simple HIP tests +- CI environments that don't require comprehensive testing +- Debugging simple test failures without the overhead of Catch tests +- Systems with limited resources where minimal test coverage is acceptable + +### HIP_CATCH_TEST_VERBOSE + +**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: +- **OFF**: Shows LIT summary only, full details available in `.test.out` files +- **ON**: Shows full test output in terminal (LIT `-a` flag) + +Example: +```bash +# 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:** +- **Passed**: Tests that completed successfully with passing assertions +- **Failed**: Tests that failed assertions +- **Skipped**: Tests that intentionally didn't run (detected via "is skipped" messages) - exit code 0 +- **Crashed/Error**: Tests that terminated abnormally (e.g., segfault, abort) - non-zero exit code + +**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. + +### CATCH_TEST_CATEGORIES + +**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: +```bash +# Include unit tests (default and currently only available category) +-DCATCH_TEST_CATEGORIES="unit" +``` + +### CATCH_TEST_SUBDIRS + +**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: +```bash +# 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 tests +- `memory` - Memory management tests +- `stream` - Stream management tests +- `event` - Event handling tests +- `kernel` - Kernel execution tests +- `device` - Device management tests +- `math` - Math operations tests +- And 40+ more subdirectories... + +### HIP_CATCH_TEST_TIMEOUT + +**Type**: Integer (seconds) +**Default**: 60 +**Description**: Timeout for individual Catch tests + +```bash +-DHIP_CATCH_TEST_TIMEOUT=120 +``` + +## Example Configurations + +### Example 1: Basic Unit Tests Only + +```bash +#!/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 +``` + +### Example 2: Selective Subdirectory Testing + +Build only specific subdirectories (e.g., compiler tests): + +```bash +#!/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 +``` + +### Example 3: Comprehensive Testing + +```bash +#!/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 +``` + +### Example 4: Simple Tests Only (Catch Tests Disabled) + +For quick builds or when you only need simple HIP tests: + +```bash +#!/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 +``` + +## Target Hierarchy + +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. + +### Hierarchy Levels + +``` +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) +``` + +### Benefits of Hierarchical Targets + +1. **Flexibility**: Build/run at any granularity level +2. **Efficiency**: Test only what you need during development +3. **Multi-Variant Support**: Test across multiple ROCm versions or focus on one +4. **Incremental Testing**: Quick feedback by testing changed components + +### Example Workflows + +**Scenario 1: Compiler developer working on code generation** +```bash +# 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** +```bash +# 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** +```bash +# Run all unit tests across all variants +ninja check-hip-catch-unit +``` + +**Scenario 4: Nightly CI testing** +```bash +# Run everything +ninja check-hip-catch +``` + +## Test Organization + +The framework organizes tests by: + +1. **ROCm Version**: Each ROCm installation gets its own test variant + - Example: `hip-7.2.0`, `hip-6.2.0` + +2. **Category**: Tests are grouped by category + - `unit`: Core API functionality + - `stress`: High-load scenarios + - `performance`: Benchmarking + - `perftests`: Detailed metrics + +3. **Subdirectory**: Each subdirectory within a category + - `unit/compiler`, `unit/memory`, `unit/stream`, etc. + +4. **Test Executable**: Each subdirectory produces a test executable + - Example: `catch_unit_compiler-hip-7.2.0` + +## Building Tests at Different Levels + +### Build by Category + +```bash +# 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 +``` + +### Build by Subdirectory + +```bash +# 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 Everything + +```bash +# 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 +``` + +## Running Tests at Different Levels + +### Run All Catch Tests + +```bash +# 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 Category-Level Tests + +```bash +# 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 Subdirectory-Level Tests + +```bash +# 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 +``` + +### Run Individual Test Executable + +Each `.cc` source file is compiled into a separate test executable, allowing for individual test execution: + +```bash +# 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_*" +``` + +### Using LIT + +The tests are integrated with the LIT test runner. Each `.cc` source file creates a separate LIT test: + +```bash +# 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. + +## Troubleshooting + +### Catch2 Download Failed + +**Error**: `FetchContent failed to download Catch2` + +**Cause**: No internet connection during first configure, and Catch2 is not installed system-wide. + +**Solutions**: +1. **Install Catch2 system-wide** (recommended for offline builds): + ```bash + # 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 + ``` + +2. **Ensure internet connectivity** during first CMake configure. Subsequent builds use the cached download. + +### No Test Sources Found + +**Error**: `No test sources found in category/subdir` + +**Solution**: The CMakeLists.txt parser may have failed. This can happen if: +- The test category doesn't exist in your hip-tests version +- The CMakeLists.txt format is not recognized + +Check available test categories: +```bash +ls /path/to/hip-tests/catch/ +``` + +### Compilation Errors + +**Common issues**: + +1. **Missing includes**: Ensure ROCm headers are accessible +2. **Architecture mismatch**: Verify `AMDGPU_ARCHS` matches your GPU +3. **ROCm version mismatch**: Ensure HIP runtime matches compiler version + +## Architecture + +The framework consists of: + +1. **HipCatchTests.cmake**: Main CMake module with functions for: + - Repository discovery and validation + - Test category discovery + - Test executable creation (one per `.cc` source file) + - LIT integration + +2. **Modified CMakeLists.txt**: Integration points in the main HIP test CMakeLists.txt: + - Initialization in `create_hip_tests()` + - Per-variant integration in `create_hip_test()` + - Target dependencies in `hip-tests-all` + +3. **Test Granularity**: Each `.cc` source file is compiled into a separate test executable, allowing: + - Individual test tracking in LIT + - Parallel test execution + - Fine-grained failure reporting + - Independent test filtering + +4. **Test Wrappers**: Shell scripts for LIT integration + +## Extending the Framework + +### Adding Custom Test Categories + +To add support for additional test categories: + +1. Add the category to `CATCH_TEST_CATEGORIES`: + ```bash + -DCATCH_TEST_CATEGORIES="unit;stress;mycategory" + ``` + +2. The framework will automatically discover subdirectories in: + ``` + llvm-test-suite/External/HIP/catch/mycategory/ + ``` + +## Performance Considerations + +- **Build Time**: Each `.cc` file is compiled into a separate executable, which can increase build time. Use specific categories or subdirectories to reduce build scope. +- **Disk Space**: Each `.cc` source file creates a separate executable per variant. For example, 3 source files with 2 variants = 6 executables. +- **Parallelism**: Use ninja's `-j` flag to parallelize builds (particularly beneficial with multiple test executables): + ```bash + ninja -j16 hip-tests-catch + ``` + +## Contributing + +To contribute improvements to the framework: + +1. Test changes with multiple ROCm versions +2. Verify tests build and run correctly +3. Update this documentation +4. Submit changes to the LLVM Test Suite + +## Future Enhancements + +Potential improvements: +- [ ] Add more test categories (stress, performance, etc.) +- [ ] Add more unit test subdirectories (memory, streams, etc.) +- [ ] Test selection by regex pattern +- [ ] Integration with CTest +- [ ] Test result aggregation and reporting +- [ ] Enhanced test filtering options + +## Directory Structure + +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. + +## References + +- [hip-tests in ROCm Systems](https://github.com/ROCm/rocm-systems/tree/develop/projects/hip-tests) (original source of tests) +- [ROCm Documentation](https://rocm.docs.amd.com/) +- [LLVM Test Suite](https://llvm.org/docs/TestSuiteGuide.html) +- [Catch2 Framework](https://github.com/catchorg/Catch2) diff --git a/External/HIP/CMakeLists.txt b/External/HIP/CMakeLists.txt index 37d1cf6c24..dabb71da28 100644 --- a/External/HIP/CMakeLists.txt +++ b/External/HIP/CMakeLists.txt @@ -7,6 +7,9 @@ message(STATUS "TEST_SUITE_HIP_ROOT: ${TEST_SUITE_HIP_ROOT}") get_filename_component(HIP_CLANG_PATH ${CMAKE_CXX_COMPILER} DIRECTORY) message(STATUS "HIP_CLANG_PATH: ${HIP_CLANG_PATH}") +# Include Catch tests integration module +include(${CMAKE_CURRENT_LIST_DIR}/HipCatchTests.cmake) + # Inspired from create_one_local_test. Runs hipify on the TestSource and then compiles it. # Search for the reference files next to TestSource. macro(create_one_hipify_cuda_test TestName TestSource VairantOffload VariantSuffix VariantCPPFlags VariantLibs) @@ -151,6 +154,11 @@ function(create_hip_test VariantSuffix) DEPENDS hip-tests-simple-${VariantSuffix} USES_TERMINAL) add_dependencies(check-hip-simple check-hip-simple-${VariantSuffix}) + + # Integrate Catch tests for this variant (if enabled) + if(ENABLE_HIP_CATCH_TESTS) + integrate_catch_tests(${VariantSuffix} ${_RocmPath}) + endif() endfunction(create_hip_test) macro(create_hip_tests) @@ -176,6 +184,9 @@ macro(create_hip_tests) add_custom_target(check-hip-simple COMMENT "Run all simple HIP tests") + # Initialize Catch tests framework + initialize_catch_tests() + if(NOT AMDGPU_ARCHS) list(APPEND AMDGPU_ARCHS "gfx906;gfx90a;gfx1030;gfx1100;native") endif() @@ -260,8 +271,14 @@ macro(create_hip_tests) endif() - add_custom_target(hip-tests-all DEPENDS hip-tests-simple - COMMENT "Build all HIP tests.") + # Build all HIP tests (simple + catch if enabled) + if(ENABLE_HIP_CATCH_TESTS) + add_custom_target(hip-tests-all DEPENDS hip-tests-simple hip-tests-catch + COMMENT "Build all HIP tests (simple + catch).") + else() + add_custom_target(hip-tests-all DEPENDS hip-tests-simple + COMMENT "Build all HIP tests (simple only, catch disabled).") + endif() file(COPY lit.local.cfg DESTINATION "${CMAKE_CURRENT_BINARY_DIR}") endmacro(create_hip_tests) diff --git a/External/HIP/HipCatchTests.cmake b/External/HIP/HipCatchTests.cmake new file mode 100644 index 0000000000..7ded76ce3b --- /dev/null +++ b/External/HIP/HipCatchTests.cmake @@ -0,0 +1,922 @@ +# CMake module for integrating hip-tests Catch tests into llvm-test-suite +# This module provides functions to discover and build Catch-based HIP tests + +include(CMakeParseArguments) + +# Global variables for Catch test configuration +set(ENABLE_HIP_CATCH_TESTS OFF CACHE BOOL "Enable HIP Catch test framework and all catch test targets") +set(CATCH_TEST_CATEGORIES "unit" CACHE STRING "Semicolon-separated list of test categories to include (unit;stress;performance;perftests)") +set(CATCH_TEST_SUBDIRS "" CACHE STRING "Semicolon-separated list of test subdirectories to include (e.g., compiler;memory;stream). Empty means all subdirectories within enabled categories.") +set(HIP_CATCH_TEST_TIMEOUT 60 CACHE STRING "Timeout for individual Catch tests in seconds") +set(HIP_CATCH_TEST_VERBOSE OFF CACHE BOOL "Show verbose output with individual TEST_CASE results from Catch2") + +# Local paths for Catch test infrastructure +set(HIP_CATCH_TESTS_DIR "${CMAKE_CURRENT_LIST_DIR}/catch") + +# Try to find system-installed Catch2 v2.13.10+ +# Note: v2.13.10 is used because v2.13.4 has glibc 2.34+ incompatibility (MINSIGSTKSZ issue) +find_package(Catch2 2.13.10 QUIET) + +if(Catch2_FOUND) + message(STATUS "Using system Catch2: ${Catch2_DIR}") + get_target_property(CATCH2_INCLUDE_PATH Catch2::Catch2 INTERFACE_INCLUDE_DIRECTORIES) +else() + message(STATUS "Catch2 >= 2.13.10 not found on system, fetching v2.13.10...") + include(FetchContent) + FetchContent_Declare( + Catch2 + GIT_REPOSITORY https://github.com/catchorg/Catch2.git + GIT_TAG v2.13.10 + GIT_SHALLOW TRUE + ) + FetchContent_MakeAvailable(Catch2) + set(CATCH2_INCLUDE_PATH "${catch2_SOURCE_DIR}/single_include/catch2") +endif() + +set(CATCH2_FOUND TRUE) + +# Global tracking for hierarchical targets +# These will be set as GLOBAL properties to track across function calls +define_property(GLOBAL PROPERTY CATCH_ALL_VARIANTS + BRIEF_DOCS "List of all discovered HIP variants" + FULL_DOCS "Tracks all variant suffixes (e.g., hip-7.2.0) discovered during configuration") + +define_property(GLOBAL PROPERTY CATCH_CATEGORY_TARGETS_CREATED + BRIEF_DOCS "List of categories for which aggregated targets have been created" + FULL_DOCS "Prevents duplicate target creation for category-level aggregated targets") + +define_property(GLOBAL PROPERTY CATCH_SUBDIR_TARGETS_CREATED + BRIEF_DOCS "List of category-subdir pairs for which aggregated targets have been created" + FULL_DOCS "Prevents duplicate target creation for subdirectory-level aggregated targets") + +# Function to validate Catch test infrastructure +function(validate_catch_tests_infrastructure) + set(_required_paths + "${HIP_CATCH_TESTS_DIR}/unit/compiler" + "${HIP_CATCH_TESTS_DIR}/external/picojson/picojson.h" + "${HIP_CATCH_TESTS_DIR}/hipTestMain" + "${HIP_CATCH_TESTS_DIR}/include" + ) + # Note: Catch2 is now obtained via find_package or FetchContent, not vendored + # Note: kernels/ directory not required for unit/compiler tests + + foreach(_path ${_required_paths}) + if(NOT EXISTS "${_path}") + message(FATAL_ERROR "Required Catch test path not found: ${_path}") + endif() + endforeach() + + message(STATUS "Using local Catch test infrastructure: ${HIP_CATCH_TESTS_DIR}") + message(STATUS "Catch2 include path: ${CATCH2_INCLUDE_PATH}") +endfunction() + +# Function to discover test sources from hip-tests +# Arguments: +# CATEGORY - Test category (unit, stress, performance, perftests) +# SUBDIRS - Specific subdirectories to include (optional, default: all) +# Returns: +# Sets ${CATEGORY}_TEST_DIRS in parent scope +function(discover_catch_test_category CATEGORY) + set(options) + set(oneValueArgs) + set(multiValueArgs SUBDIRS) + cmake_parse_arguments(ARG "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN}) + + set(_catch_base "${HIP_CATCH_TESTS_DIR}/${CATEGORY}") + + if(NOT IS_DIRECTORY "${_catch_base}") + message(STATUS "Category ${CATEGORY} not found in Catch tests, skipping") + return() + endif() + + # If specific subdirectories requested via argument, use them; + # otherwise check CATCH_TEST_SUBDIRS cache variable; + # otherwise discover all + if(ARG_SUBDIRS) + set(_test_subdirs ${ARG_SUBDIRS}) + elseif(CATCH_TEST_SUBDIRS) + # Use subdirectories from CATCH_TEST_SUBDIRS, filtering to only those that exist + set(_test_subdirs "") + foreach(_subdir ${CATCH_TEST_SUBDIRS}) + if(IS_DIRECTORY "${_catch_base}/${_subdir}") + list(APPEND _test_subdirs "${_subdir}") + message(STATUS "Including subdirectory from CATCH_TEST_SUBDIRS: ${CATEGORY}/${_subdir}") + else() + message(STATUS "Subdirectory ${CATEGORY}/${_subdir} not found, skipping") + endif() + endforeach() + else() + # Discover all subdirectories automatically. + file(GLOB _potential_subdirs RELATIVE "${_catch_base}" "${_catch_base}/*") + set(_test_subdirs "") + foreach(_subdir ${_potential_subdirs}) + if(IS_DIRECTORY "${_catch_base}/${_subdir}") + # Check if directory has any .cc test files + file(GLOB _test_files "${_catch_base}/${_subdir}/*.cc") + if(_test_files) + list(APPEND _test_subdirs "${_subdir}") + endif() + endif() + endforeach() + endif() + + if(_test_subdirs) + message(STATUS "Discovered ${CATEGORY} test subdirectories: ${_test_subdirs}") + else() + message(STATUS "No test subdirectories found for ${CATEGORY} category") + endif() + set(${CATEGORY}_TEST_DIRS "${_test_subdirs}" PARENT_SCOPE) +endfunction() + +# Helper function to track test targets at multiple levels +# Arguments: +# TEST_TARGET - The test target name (e.g., catch_unit_compiler-hip-7.2.0.test) +# CATEGORY - Category name (e.g., unit) +# SUBDIR - Subdirectory name (e.g., compiler) +# VARIANT - Variant suffix (e.g., hip-7.2.0) +function(track_test_target_multi_level TEST_TARGET CATEGORY SUBDIR VARIANT) + # Track at global level (all catch tests) + set_property(GLOBAL APPEND PROPERTY CATCH_ALL_TEST_TARGETS "${TEST_TARGET}") + + # Track at per-variant global level + set_property(GLOBAL APPEND PROPERTY "CATCH_TEST_TARGETS_${VARIANT}" "${TEST_TARGET}") + + # Track at per-category level (all variants) + set_property(GLOBAL APPEND PROPERTY "CATCH_${CATEGORY}_TEST_TARGETS" "${TEST_TARGET}") + + # Track at per-category-variant level + set_property(GLOBAL APPEND PROPERTY "CATCH_${CATEGORY}_TEST_TARGETS_${VARIANT}" "${TEST_TARGET}") + + # Track at per-subdirectory level (all variants) + set_property(GLOBAL APPEND PROPERTY "CATCH_${CATEGORY}_${SUBDIR}_TEST_TARGETS" "${TEST_TARGET}") + + # Track at per-subdirectory-variant level + set_property(GLOBAL APPEND PROPERTY "CATCH_${CATEGORY}_${SUBDIR}_TEST_TARGETS_${VARIANT}" "${TEST_TARGET}") +endfunction() + +# Function to create special generic target executables for hipSquareGenericTarget test +# This test requires additional executables built with generic-only offload architectures +# Arguments: +# TEST_BASENAME - Base name of the test (e.g., hipSquareGenericTarget) +# TEST_DIR - Directory containing the test sources +# VARIANT_SUFFIX - Variant suffix (e.g., hip-7.2.0) +# ROCM_PATH - Path to ROCm installation +function(create_generic_target_executables TEST_BASENAME TEST_DIR VARIANT_SUFFIX ROCM_PATH) + # Check if this is AMD platform (generic targets are AMD-specific) + get_filename_component(_compiler_name "${CMAKE_CXX_COMPILER}" NAME) + if(NOT (_compiler_name MATCHES "hipcc" OR _compiler_name MATCHES "clang")) + message(STATUS "Skipping generic target executables (not AMD platform)") + return() + endif() + + message(STATUS "Creating generic target executables for ${TEST_BASENAME}-${VARIANT_SUFFIX}") + + # Generic target architecture flags + set(_generic_archs + "--offload-arch=gfx9-generic" + "--offload-arch=gfx9-4-generic:sramecc+:xnack-" + "--offload-arch=gfx9-4-generic:sramecc-:xnack-" + "--offload-arch=gfx9-4-generic:xnack+" + "--offload-arch=gfx10-1-generic" + "--offload-arch=gfx10-3-generic" + "--offload-arch=gfx11-generic" + "--offload-arch=gfx12-generic" + ) + + set(_source_file "${TEST_DIR}/${TEST_BASENAME}.cc") + set(_output_dir "${CMAKE_CURRENT_BINARY_DIR}/catch_tests") + + # Common source files + set(_common_sources + "${_source_file}" + "${HIP_CATCH_TESTS_DIR}/hipTestMain/hip_test_context.cc" + "${HIP_CATCH_TESTS_DIR}/hipTestMain/hip_test_features.cc" + "${HIP_CATCH_TESTS_DIR}/hipTestMain/main.cc" + ) + + # Common include directories + set(_include_flags + "-I${ROCM_PATH}/include" + "-I${HIP_CATCH_TESTS_DIR}/include" + "-I${CATCH2_INCLUDE_PATH}" + "-I${HIP_CATCH_TESTS_DIR}/external/picojson" + ) + + # Determine library linking flags + if(WIN32) + set(_libfs_flag "") + set(_exe_suffix ".exe") + else() + set(_libfs_flag "-lstdc++fs") + set(_exe_suffix "") + endif() + + # 1. Build hipSquareGenericTargetOnly (regular fatbin with generic targets only) + set(_exe_name_regular "hipSquareGenericTargetOnly${_exe_suffix}") + set(_output_path_regular "${_output_dir}/${_exe_name_regular}") + + add_custom_command( + OUTPUT "${_output_path_regular}" + COMMAND ${CMAKE_COMMAND} -E make_directory "${_output_dir}" + COMMAND ${CMAKE_CXX_COMPILER} + -DNO_GENERIC_TARGET_ONLY_TEST + --std=c++17 + -x hip + -mcode-object-version=6 + -w + ${_generic_archs} + ${_common_sources} + -o "${_output_path_regular}" + --hip-path=${ROCM_PATH} + --rocm-path=${ROCM_PATH} + --hip-link + -rtlib=compiler-rt + -unwindlib=libgcc + -frtlib-add-rpath + ${_include_flags} + ${_libfs_flag} + DEPENDS ${_common_sources} + WORKING_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}" + COMMENT "Building ${_exe_name_regular} for ${VARIANT_SUFFIX}" + VERBATIM + ) + + # 2. Build hipSquareGenericTargetOnlyCompressed (compressed fatbin with generic targets only) + set(_exe_name_compressed "hipSquareGenericTargetOnlyCompressed${_exe_suffix}") + set(_output_path_compressed "${_output_dir}/${_exe_name_compressed}") + + add_custom_command( + OUTPUT "${_output_path_compressed}" + COMMAND ${CMAKE_COMMAND} -E make_directory "${_output_dir}" + COMMAND ${CMAKE_CXX_COMPILER} + -DNO_GENERIC_TARGET_ONLY_TEST + -DGENERIC_COMPRESSED + --std=c++17 + -x hip + -mcode-object-version=6 + --offload-compress + -w + ${_generic_archs} + ${_common_sources} + -o "${_output_path_compressed}" + --hip-path=${ROCM_PATH} + --rocm-path=${ROCM_PATH} + --hip-link + -rtlib=compiler-rt + -unwindlib=libgcc + -frtlib-add-rpath + ${_include_flags} + ${_libfs_flag} + DEPENDS ${_common_sources} + WORKING_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}" + COMMENT "Building ${_exe_name_compressed} for ${VARIANT_SUFFIX}" + VERBATIM + ) + + # Create custom targets for these executables + add_custom_target(hipSquareGenericTargetOnly-${VARIANT_SUFFIX} + DEPENDS "${_output_path_regular}" + ) + + add_custom_target(hipSquareGenericTargetOnlyCompressed-${VARIANT_SUFFIX} + DEPENDS "${_output_path_compressed}" + ) + + # Make the main test executable depend on these + set(_main_test_exe "catch_unit_compiler_${TEST_BASENAME}-${VARIANT_SUFFIX}") + if(TARGET ${_main_test_exe}) + add_dependencies(${_main_test_exe} + hipSquareGenericTargetOnly-${VARIANT_SUFFIX} + hipSquareGenericTargetOnlyCompressed-${VARIANT_SUFFIX} + ) + message(STATUS "Added generic target executable dependencies to ${_main_test_exe}") + endif() + + # Also add to the build target hierarchy + if(TARGET hip-tests-catch-unit-compiler-${VARIANT_SUFFIX}) + add_dependencies(hip-tests-catch-unit-compiler-${VARIANT_SUFFIX} + hipSquareGenericTargetOnly-${VARIANT_SUFFIX} + hipSquareGenericTargetOnlyCompressed-${VARIANT_SUFFIX} + ) + endif() +endfunction() + +# Function to create a Catch test executable +# Arguments: +# TEST_NAME - Name of the test +# TEST_SOURCES - Source file(s) (can be a single file or list of files) +# TEST_DIR - Directory containing the test sources +# CATEGORY - Test category (unit, stress, etc.) +# SUBDIR - Subdirectory name +# VARIANT_SUFFIX - Variant suffix (e.g., hip-7.2.0) +# ROCM_PATH - Path to ROCm installation +macro(create_catch_test_executable TEST_NAME TEST_SOURCES TEST_DIR CATEGORY SUBDIR VARIANT_SUFFIX ROCM_PATH) + set(_test_exe "${TEST_NAME}-${VARIANT_SUFFIX}") + set(_test_sources "") + + # Build full paths to source files + foreach(_src ${TEST_SOURCES}) + list(APPEND _test_sources "${TEST_DIR}/${_src}") + endforeach() + + # Add hipTestMain sources (required for catch2 integration) + list(APPEND _test_sources + "${HIP_CATCH_TESTS_DIR}/hipTestMain/main.cc" + "${HIP_CATCH_TESTS_DIR}/hipTestMain/hip_test_context.cc" + "${HIP_CATCH_TESTS_DIR}/hipTestMain/hip_test_features.cc" + ) + + # Check if sources exist + set(_valid_sources "") + foreach(_src ${_test_sources}) + if(EXISTS "${_src}") + list(APPEND _valid_sources "${_src}") + else() + message(STATUS "Source file not found: ${_src}") + endif() + endforeach() + + if(NOT _valid_sources) + message(STATUS "No valid sources found for ${TEST_NAME}, skipping") + return() + endif() + + # Create the executable + add_executable(${_test_exe} EXCLUDE_FROM_ALL ${_valid_sources}) + + # Ensure timeit tool is built first (needed for compilation timing) + if(TARGET build-timeit) + add_dependencies(${_test_exe} build-timeit) + endif() + + # Set properties + set_target_properties(${_test_exe} PROPERTIES + CXX_STANDARD 17 + CXX_STANDARD_REQUIRED ON + RUNTIME_OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/catch_tests" + ) + + # Include directories + target_include_directories(${_test_exe} PRIVATE + ${CATCH2_INCLUDE_PATH} + "${HIP_CATCH_TESTS_DIR}/include" + "${HIP_CATCH_TESTS_DIR}/external/picojson" + ) + + # Add HIP runtime includes + # hipcc wrapper provides these automatically, direct compilers need explicit paths + if(NOT _compiler_name MATCHES "hipcc") + target_include_directories(${_test_exe} PRIVATE + "${ROCM_PATH}/include" + ) + endif() + + # Compile definitions + # Note: __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__ is automatically + # defined by HIP headers based on the compiler backend, so no manual + # platform definition is needed + # target_compile_definitions(${_test_exe} PRIVATE ...) + + # Compile options + target_compile_options(${_test_exe} PRIVATE + ${VariantCPPFLAGS} + -x hip + -Wall + -Wno-deprecated + -Wno-unused-command-line-argument + ) + + # Special handling for hipSquareGenericTarget test - add generic target architectures + if("${TEST_NAME}" MATCHES "hipSquareGenericTarget" AND "${CATEGORY}" STREQUAL "unit" AND "${SUBDIR}" STREQUAL "compiler") + target_compile_options(${_test_exe} PRIVATE + -mcode-object-version=6 + -w + --offload-arch=gfx9-generic + --offload-arch=gfx9-4-generic:sramecc+:xnack- + --offload-arch=gfx9-4-generic:sramecc-:xnack- + --offload-arch=gfx9-4-generic:xnack+ + --offload-arch=gfx10-1-generic + --offload-arch=gfx10-3-generic + --offload-arch=gfx11-generic + --offload-arch=gfx12-generic + ) + message(STATUS "Added generic target compile options to ${_test_exe}") + endif() + + # Link options - platform-specific handling + # If using hipcc wrapper (AMD or NVIDIA backend), it handles flags automatically + # Otherwise, add explicit flags for AMD clang + get_filename_component(_compiler_name "${CMAKE_CXX_COMPILER}" NAME) + if(_compiler_name MATCHES "hipcc") + # hipcc wrapper handles platform-specific flags automatically + message(VERBOSE "Using hipcc wrapper: ${CMAKE_CXX_COMPILER}") + else() + # Direct compiler (AMD clang) - add explicit HIP link flags + target_link_options(${_test_exe} PRIVATE + --rocm-path=${ROCM_PATH} + --hip-link + -rtlib=compiler-rt + -unwindlib=libgcc + -frtlib-add-rpath + ) + endif() + + # Link libraries + target_link_libraries(${_test_exe} PRIVATE + ${VariantLibs} + stdc++fs + dl + pthread + rt + ) + + # Add to subdirectory variant target (lowest level) + add_dependencies(hip-tests-catch-${CATEGORY}-${SUBDIR}-${VARIANT_SUFFIX} ${_test_exe}) + + # Create a test wrapper script + set(_test_wrapper "${CMAKE_CURRENT_BINARY_DIR}/catch_tests/${_test_exe}_wrapper.sh") + file(WRITE "${_test_wrapper}" "#!/bin/bash\n") + file(APPEND "${_test_wrapper}" "# Auto-generated wrapper for ${_test_exe}\n") + + # Special handling for hipSquareGenericTarget - needs to run from catch_tests dir + # so that helper executables (hipSquareGenericTargetOnly, etc.) can be found + if("${TEST_NAME}" MATCHES "hipSquareGenericTarget") + file(APPEND "${_test_wrapper}" "cd \"${CMAKE_CURRENT_BINARY_DIR}/catch_tests\" || exit 1\n") + endif() + + # Add verbose reporting if enabled + # Use console reporter for consistent output parsing in summary scripts + # Console reporter outputs: "test cases: X | Y passed | Z failed" + # Also capture exit code for accurate skip vs crash detection + file(APPEND "${_test_wrapper}" "echo \"=== Running: ${_test_exe} ===\"\n") + if("${TEST_NAME}" MATCHES "hipSquareGenericTarget") + # Use relative path since we're in catch_tests dir + file(APPEND "${_test_wrapper}" "\"./${_test_exe}\" --reporter console \"$@\"\n") + else() + file(APPEND "${_test_wrapper}" "\"${CMAKE_CURRENT_BINARY_DIR}/catch_tests/${_test_exe}\" --reporter console \"$@\"\n") + endif() + file(APPEND "${_test_wrapper}" "TEST_EXIT_CODE=\$?\n") + file(APPEND "${_test_wrapper}" "echo \"EXIT_CODE: \$TEST_EXIT_CODE\"\n") + file(APPEND "${_test_wrapper}" "exit \$TEST_EXIT_CODE\n") + + execute_process(COMMAND chmod +x "${_test_wrapper}") + + # Register with LIT + # Use relative path from build directory for the wrapper script + llvm_test_run(EXECUTABLE "/bin/bash" "catch_tests/${_test_exe}_wrapper.sh") + + # Add verification to check if test passed + # Catch2 prints "test cases: X | Y failed" when tests fail + # Check for failure indicators in output (%o expands to Output/.test.out) + llvm_test_verify("! grep -q failed %o") + + llvm_add_test(${_test_exe}.test catch_tests/${_test_exe}_wrapper.sh) + + # Track this test target at all hierarchy levels + track_test_target_multi_level(${_test_exe}.test ${CATEGORY} ${SUBDIR} ${VARIANT_SUFFIX}) + + # Also add to VARIANT_CATCH_TEST_TARGETS for backward compatibility + list(APPEND VARIANT_CATCH_TEST_TARGETS ${_test_exe}.test) + set(VARIANT_CATCH_TEST_TARGETS ${VARIANT_CATCH_TEST_TARGETS} PARENT_SCOPE) + + message(STATUS "Created Catch test executable: ${_test_exe}") +endmacro() + +# Function to create Catch tests for a specific category and subdirectory +# Arguments: +# CATEGORY - Test category (unit, stress, performance) +# SUBDIR - Subdirectory name +# VARIANT_SUFFIX - Variant suffix (e.g., hip-7.2.0) +# ROCM_PATH - Path to ROCm installation +function(create_catch_tests_for_subdir CATEGORY SUBDIR VARIANT_SUFFIX ROCM_PATH) + set(_test_dir "${HIP_CATCH_TESTS_DIR}/${CATEGORY}/${SUBDIR}") + + # Create unique identifier for this category-subdir pair + set(_subdir_id "${CATEGORY}-${SUBDIR}") + + # Create aggregated subdirectory-level targets (once per category-subdir pair) + get_property(_created_subdirs GLOBAL PROPERTY CATCH_SUBDIR_TARGETS_CREATED) + if(NOT "${_subdir_id}" IN_LIST _created_subdirs) + message(STATUS "Creating aggregated subdirectory targets for: ${CATEGORY}/${SUBDIR}") + + # Create aggregated build target + add_custom_target(hip-tests-catch-${CATEGORY}-${SUBDIR} + COMMENT "Build all HIP Catch ${CATEGORY}/${SUBDIR} tests across all variants") + add_dependencies(hip-tests-catch-${CATEGORY} hip-tests-catch-${CATEGORY}-${SUBDIR}) + + # Create aggregated check target (will be populated later) + add_custom_target(check-hip-catch-${CATEGORY}-${SUBDIR} + COMMENT "Run all HIP Catch ${CATEGORY}/${SUBDIR} tests across all variants") + add_dependencies(check-hip-catch-${CATEGORY} check-hip-catch-${CATEGORY}-${SUBDIR}) + + # Mark as created + set_property(GLOBAL APPEND PROPERTY CATCH_SUBDIR_TARGETS_CREATED "${_subdir_id}") + endif() + + # Create per-variant subdirectory-level targets + message(STATUS "Creating per-variant subdirectory targets for: ${CATEGORY}/${SUBDIR}-${VARIANT_SUFFIX}") + add_custom_target(hip-tests-catch-${CATEGORY}-${SUBDIR}-${VARIANT_SUFFIX} + COMMENT "Build HIP Catch ${CATEGORY}/${SUBDIR} tests for variant ${VARIANT_SUFFIX}") + + # Wire dependencies: variant-specific target to aggregated subdirectory target + add_dependencies(hip-tests-catch-${CATEGORY}-${SUBDIR} hip-tests-catch-${CATEGORY}-${SUBDIR}-${VARIANT_SUFFIX}) + + # Wire dependencies: category variant target depends on subdirectory variant target + add_dependencies(hip-tests-catch-${CATEGORY}-${VARIANT_SUFFIX} hip-tests-catch-${CATEGORY}-${SUBDIR}-${VARIANT_SUFFIX}) + + # Discover test sources directly from .cc files + file(GLOB _test_sources RELATIVE "${_test_dir}" "${_test_dir}/*.cc") + + if(NOT _test_sources) + message(STATUS "No test sources found in ${CATEGORY}/${SUBDIR}, skipping") + return() + endif() + + message(STATUS "Discovered test sources in ${CATEGORY}/${SUBDIR}: ${_test_sources}") + + # Create a separate test executable for each source file + # This allows LIT to report statistics for each individual test + foreach(_src ${_test_sources}) + # Get the test name from the source filename (without extension) + get_filename_component(_test_basename "${_src}" NAME_WE) + set(_test_name "catch_${CATEGORY}_${SUBDIR}_${_test_basename}") + + message(STATUS "Creating Catch test: ${_test_name} from ${_src}") + + # Create the test executable with just this one source file + create_catch_test_executable("${_test_name}" "${_src}" "${_test_dir}" "${CATEGORY}" "${SUBDIR}" "${VARIANT_SUFFIX}" "${ROCM_PATH}") + + # Special handling for hipSquareGenericTarget test (unit/compiler only) + # This test requires additional executables with generic-only targets + if("${_test_basename}" STREQUAL "hipSquareGenericTarget" AND + "${CATEGORY}" STREQUAL "unit" AND "${SUBDIR}" STREQUAL "compiler") + create_generic_target_executables("${_test_basename}" "${_test_dir}" "${VARIANT_SUFFIX}" "${ROCM_PATH}") + endif() + endforeach() + + # Create per-variant check target for this subdirectory + # Get test targets for this subdirectory-variant combination + get_property(_subdir_variant_tests GLOBAL PROPERTY "CATCH_${CATEGORY}_${SUBDIR}_TEST_TARGETS_${VARIANT_SUFFIX}") + + if(_subdir_variant_tests) + # Unified LIT-based execution with configurable verbosity + # HIP_CATCH_TEST_VERBOSE controls output detail level (-a flag), not execution mode + # Metrics are always collected via LIT + + # Configure LIT flags based on verbosity setting + if(HIP_CATCH_TEST_VERBOSE) + set(_lit_verbosity_flags "-a") # Show all test output + else() + set(_lit_verbosity_flags "") # Quiet mode - only show pass/fail + endif() + + # Generate enhanced summary script with individual test tracking + set(_summary_script "${CMAKE_CURRENT_BINARY_DIR}/catch_tests/summary_${CATEGORY}_${SUBDIR}_${VARIANT_SUFFIX}.sh") + file(WRITE "${_summary_script}" "#!/bin/bash\n") + file(APPEND "${_summary_script}" "# Enhanced summary script for Catch2 TEST_CASE statistics\n") + file(APPEND "${_summary_script}" "cd ${CMAKE_CURRENT_BINARY_DIR}\n") + file(APPEND "${_summary_script}" "echo \"\"\n") + file(APPEND "${_summary_script}" "echo \"========================================\"\n") + file(APPEND "${_summary_script}" "echo \"Detailed Test Summary:\"\n") + file(APPEND "${_summary_script}" "TOTAL_FILES=0\n") + file(APPEND "${_summary_script}" "TOTAL_TESTS=0\n") + file(APPEND "${_summary_script}" "PASSED_TESTS=0\n") + file(APPEND "${_summary_script}" "FAILED_TESTS=0\n") + file(APPEND "${_summary_script}" "SKIPPED_TESTS=0\n") + file(APPEND "${_summary_script}" "CRASHED_TESTS=0\n") + file(APPEND "${_summary_script}" "# Arrays to track test names for triage\n") + file(APPEND "${_summary_script}" "FAILED_LIST=\"\"\n") + file(APPEND "${_summary_script}" "SKIPPED_LIST=\"\"\n") + file(APPEND "${_summary_script}" "CRASHED_LIST=\"\"\n") + file(APPEND "${_summary_script}" "for test in catch_tests/catch_${CATEGORY}_${SUBDIR}_*-${VARIANT_SUFFIX}; do\n") + file(APPEND "${_summary_script}" " if [ -x \"\$test\" ]; then\n") + file(APPEND "${_summary_script}" " TOTAL_FILES=\$((TOTAL_FILES + 1))\n") + file(APPEND "${_summary_script}" " TEST_BASENAME=\$(basename \"\$test\")\n") + file(APPEND "${_summary_script}" " # Get test names from --list-tests\n") + file(APPEND "${_summary_script}" " LIST_OUTPUT=\$(\"\$test\" --list-tests 2>&1)\n") + file(APPEND "${_summary_script}" " FILE_TOTAL=\$(echo \"\$LIST_OUTPUT\" | tail -1 | grep -o '^[0-9]*' || echo 0)\n") + file(APPEND "${_summary_script}" " TOTAL_TESTS=\$((TOTAL_TESTS + FILE_TOTAL))\n") + file(APPEND "${_summary_script}" " # Extract individual test names (lines starting with spaces after 'All available test cases:')\n") + file(APPEND "${_summary_script}" " TEST_NAMES=\$(echo \"\$LIST_OUTPUT\" | grep '^ ' | sed 's/^ //')\n") + file(APPEND "${_summary_script}" " # Parse the corresponding .test.out file for results\n") + file(APPEND "${_summary_script}" " OUT_FILE=\"Output/\${TEST_BASENAME}.test.out\"\n") + file(APPEND "${_summary_script}" " if [ -f \"\$OUT_FILE\" ]; then\n") + file(APPEND "${_summary_script}" " CASES_PASSED=0\n") + file(APPEND "${_summary_script}" " CASES_FAILED=0\n") + file(APPEND "${_summary_script}" " PARSED=0\n") + file(APPEND "${_summary_script}" " FILE_FAILED_NAMES=\"\"\n") + file(APPEND "${_summary_script}" " FILE_SKIPPED_NAMES=\"\"\n") + file(APPEND "${_summary_script}" " FILE_CRASHED_NAMES=\"\"\n") + file(APPEND "${_summary_script}" " # Parse Catch2 output - try multiple formats (prefer final summary over intermediate)\n") + file(APPEND "${_summary_script}" " # Format 1: 'All tests passed (N assertion in M test cases)' - final summary, most accurate\n") + file(APPEND "${_summary_script}" " ALL_PASSED_LINE=\$(grep 'All tests passed' \"\$OUT_FILE\" 2>/dev/null | tail -1 || echo \"\")\n") + file(APPEND "${_summary_script}" " if [ -n \"\$ALL_PASSED_LINE\" ]; then\n") + file(APPEND "${_summary_script}" " PARSED=1\n") + file(APPEND "${_summary_script}" " CASES_PASSED=\$(echo \"\$ALL_PASSED_LINE\" | grep -o 'in [0-9]* test case' | grep -o '[0-9]*' || echo 0)\n") + file(APPEND "${_summary_script}" " CASES_FAILED=0\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " # Format 2: 'test cases: X | Y passed | Z failed' - use LAST occurrence (final summary)\n") + file(APPEND "${_summary_script}" " if [ \$PARSED -eq 0 ]; then\n") + file(APPEND "${_summary_script}" " SUMMARY_LINE=\$(grep '^test cases:' \"\$OUT_FILE\" 2>/dev/null | tail -1 || echo \"\")\n") + file(APPEND "${_summary_script}" " if [ -n \"\$SUMMARY_LINE\" ]; then\n") + file(APPEND "${_summary_script}" " PARSED=1\n") + file(APPEND "${_summary_script}" " CASES_PASSED=\$(echo \"\$SUMMARY_LINE\" | grep -o '[0-9]* passed' | grep -o '[0-9]*' || echo 0)\n") + file(APPEND "${_summary_script}" " CASES_FAILED=\$(echo \"\$SUMMARY_LINE\" | grep -o '[0-9]* failed' | grep -o '[0-9]*' || echo 0)\n") + file(APPEND "${_summary_script}" " CASES_TOTAL=\$(echo \"\$SUMMARY_LINE\" | sed 's/test cases: \\([0-9]*\\).*/\\1/')\n") + file(APPEND "${_summary_script}" " if [ \$CASES_PASSED -eq 0 ] && [ \$CASES_FAILED -gt 0 ]; then\n") + file(APPEND "${_summary_script}" " CASES_PASSED=\$((CASES_TOTAL - CASES_FAILED))\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " # Update totals if we parsed something\n") + file(APPEND "${_summary_script}" " if [ \$PARSED -eq 1 ]; then\n") + file(APPEND "${_summary_script}" " # Detect runtime skips by parsing 'is skipped' messages\n") + file(APPEND "${_summary_script}" " RUNTIME_SKIPPED=\$(grep -i 'is skipped' \"\$OUT_FILE\" 2>/dev/null | wc -l)\n") + file(APPEND "${_summary_script}" " if [ \$RUNTIME_SKIPPED -gt 0 ] && [ \$CASES_PASSED -gt 0 ]; then\n") + file(APPEND "${_summary_script}" " if [ \$RUNTIME_SKIPPED -gt \$CASES_PASSED ]; then\n") + file(APPEND "${_summary_script}" " RUNTIME_SKIPPED=\$CASES_PASSED\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " CASES_PASSED=\$((CASES_PASSED - RUNTIME_SKIPPED))\n") + file(APPEND "${_summary_script}" " SKIPPED_TESTS=\$((SKIPPED_TESTS + RUNTIME_SKIPPED))\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " # Find failed and skipped test names by cross-referencing --list-tests with output\n") + file(APPEND "${_summary_script}" " # Catch2 console reporter shows 'TestName passed' for tests that passed\n") + file(APPEND "${_summary_script}" " FAILED_COUNT=\$CASES_FAILED\n") + file(APPEND "${_summary_script}" " SKIPPED_COUNT=\$RUNTIME_SKIPPED\n") + file(APPEND "${_summary_script}" " if [ -n \"\$TEST_NAMES\" ]; then\n") + file(APPEND "${_summary_script}" " while IFS= read -r tname; do\n") + file(APPEND "${_summary_script}" " [ -z \"\$tname\" ] && continue\n") + file(APPEND "${_summary_script}" " # Trim trailing whitespace from test name\n") + file(APPEND "${_summary_script}" " tname=\$(echo \"\$tname\" | sed 's/[[:space:]]*\$//')\n") + file(APPEND "${_summary_script}" " [ -z \"\$tname\" ] && continue\n") + file(APPEND "${_summary_script}" " # Check if this test passed\n") + file(APPEND "${_summary_script}" " if grep -q \"^\$tname passed\" \"\$OUT_FILE\" 2>/dev/null; then\n") + file(APPEND "${_summary_script}" " : # Test passed, nothing to track\n") + file(APPEND "${_summary_script}" " elif [ \$SKIPPED_COUNT -gt 0 ]; then\n") + file(APPEND "${_summary_script}" " # Test didn't pass - count as skipped first\n") + file(APPEND "${_summary_script}" " FILE_SKIPPED_NAMES=\"\${FILE_SKIPPED_NAMES}\${tname}|\"\n") + file(APPEND "${_summary_script}" " SKIPPED_COUNT=\$((SKIPPED_COUNT - 1))\n") + file(APPEND "${_summary_script}" " elif [ \$FAILED_COUNT -gt 0 ]; then\n") + file(APPEND "${_summary_script}" " # Test didn't pass and no more skips - count as failed\n") + file(APPEND "${_summary_script}" " FILE_FAILED_NAMES=\"\${FILE_FAILED_NAMES}\${tname}|\"\n") + file(APPEND "${_summary_script}" " FAILED_COUNT=\$((FAILED_COUNT - 1))\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " done <<< \"\$TEST_NAMES\"\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " PASSED_TESTS=\$((PASSED_TESTS + CASES_PASSED))\n") + file(APPEND "${_summary_script}" " FAILED_TESTS=\$((FAILED_TESTS + CASES_FAILED))\n") + file(APPEND "${_summary_script}" " # Categorize incomplete tests using exit code\n") + file(APPEND "${_summary_script}" " INCOMPLETE=\$((FILE_TOTAL - CASES_PASSED - CASES_FAILED - RUNTIME_SKIPPED))\n") + file(APPEND "${_summary_script}" " FILE_EXIT=\$(grep '^EXIT_CODE:' \"\$OUT_FILE\" 2>/dev/null | tail -1 | grep -o '[0-9]*' || echo 1)\n") + file(APPEND "${_summary_script}" " if [ \$INCOMPLETE -gt 0 ]; then\n") + file(APPEND "${_summary_script}" " if [ \"\$FILE_EXIT\" -eq 0 ]; then\n") + file(APPEND "${_summary_script}" " SKIPPED_TESTS=\$((SKIPPED_TESTS + INCOMPLETE))\n") + file(APPEND "${_summary_script}" " else\n") + file(APPEND "${_summary_script}" " CRASHED_TESTS=\$((CRASHED_TESTS + INCOMPLETE))\n") + file(APPEND "${_summary_script}" " # Track crashed test names - tests that didn't pass and weren't explicitly categorized\n") + file(APPEND "${_summary_script}" " if [ -n \"\$TEST_NAMES\" ]; then\n") + file(APPEND "${_summary_script}" " while IFS= read -r tname; do\n") + file(APPEND "${_summary_script}" " [ -z \"\$tname\" ] && continue\n") + file(APPEND "${_summary_script}" " tname=\$(echo \"\$tname\" | sed 's/[[:space:]]*\$//')\n") + file(APPEND "${_summary_script}" " [ -z \"\$tname\" ] && continue\n") + file(APPEND "${_summary_script}" " # Check if this test passed, was skipped, or was failed\n") + file(APPEND "${_summary_script}" " if ! grep -q \"^\$tname passed\" \"\$OUT_FILE\" 2>/dev/null; then\n") + file(APPEND "${_summary_script}" " # Not passed - check if already in skipped or failed\n") + file(APPEND "${_summary_script}" " if [[ \"\$FILE_SKIPPED_NAMES\" != *\"\$tname|\"* ]] && [[ \"\$FILE_FAILED_NAMES\" != *\"\$tname|\"* ]]; then\n") + file(APPEND "${_summary_script}" " FILE_CRASHED_NAMES=\"\${FILE_CRASHED_NAMES}\${tname}|\"\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " done <<< \"\$TEST_NAMES\"\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " # Accumulate test names for triage lists (names are | delimited)\n") + file(APPEND "${_summary_script}" " IFS='|' read -ra FAILED_ARR <<< \"\$FILE_FAILED_NAMES\"\n") + file(APPEND "${_summary_script}" " for name in \"\${FAILED_ARR[@]}\"; do\n") + file(APPEND "${_summary_script}" " [ -n \"\$name\" ] && FAILED_LIST=\"\${FAILED_LIST}\${name} [\${TEST_BASENAME}]\\n\"\n") + file(APPEND "${_summary_script}" " done\n") + file(APPEND "${_summary_script}" " IFS='|' read -ra SKIPPED_ARR <<< \"\$FILE_SKIPPED_NAMES\"\n") + file(APPEND "${_summary_script}" " for name in \"\${SKIPPED_ARR[@]}\"; do\n") + file(APPEND "${_summary_script}" " [ -n \"\$name\" ] && SKIPPED_LIST=\"\${SKIPPED_LIST}\${name} [\${TEST_BASENAME}]\\n\"\n") + file(APPEND "${_summary_script}" " done\n") + file(APPEND "${_summary_script}" " IFS='|' read -ra CRASHED_ARR <<< \"\$FILE_CRASHED_NAMES\"\n") + file(APPEND "${_summary_script}" " for name in \"\${CRASHED_ARR[@]}\"; do\n") + file(APPEND "${_summary_script}" " [ -n \"\$name\" ] && CRASHED_LIST=\"\${CRASHED_LIST}\${name} [\${TEST_BASENAME}]\\n\"\n") + file(APPEND "${_summary_script}" " done\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " else\n") + file(APPEND "${_summary_script}" " # No output file - test crashed before producing output\n") + file(APPEND "${_summary_script}" " if [ \$FILE_TOTAL -gt 0 ]; then\n") + file(APPEND "${_summary_script}" " CRASHED_TESTS=\$((CRASHED_TESTS + FILE_TOTAL))\n") + file(APPEND "${_summary_script}" " # Add all test names to crashed list\n") + file(APPEND "${_summary_script}" " if [ -n \"\$TEST_NAMES\" ]; then\n") + file(APPEND "${_summary_script}" " while IFS= read -r tname; do\n") + file(APPEND "${_summary_script}" " [ -z \"\$tname\" ] && continue\n") + file(APPEND "${_summary_script}" " tname=\$(echo \"\$tname\" | sed 's/[[:space:]]*\$//')\n") + file(APPEND "${_summary_script}" " [ -n \"\$tname\" ] && CRASHED_LIST=\"\${CRASHED_LIST}\${tname} [\${TEST_BASENAME}]\\n\"\n") + file(APPEND "${_summary_script}" " done <<< \"\$TEST_NAMES\"\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" " fi\n") + file(APPEND "${_summary_script}" "done\n") + file(APPEND "${_summary_script}" "echo \" Test Suites: \$TOTAL_FILES\"\n") + file(APPEND "${_summary_script}" "echo \" Total Tests: \$TOTAL_TESTS\"\n") + file(APPEND "${_summary_script}" "echo \" Passed: \$PASSED_TESTS\"\n") + file(APPEND "${_summary_script}" "echo \" Failed: \$FAILED_TESTS\"\n") + file(APPEND "${_summary_script}" "if [ \$SKIPPED_TESTS -gt 0 ]; then\n") + file(APPEND "${_summary_script}" " echo \" Skipped: \$SKIPPED_TESTS\"\n") + file(APPEND "${_summary_script}" "fi\n") + file(APPEND "${_summary_script}" "if [ \$CRASHED_TESTS -gt 0 ]; then\n") + file(APPEND "${_summary_script}" " echo \" Crashed/Error: \$CRASHED_TESTS\"\n") + file(APPEND "${_summary_script}" "fi\n") + file(APPEND "${_summary_script}" "# Print triage lists\n") + file(APPEND "${_summary_script}" "if [ -n \"\$FAILED_LIST\" ]; then\n") + file(APPEND "${_summary_script}" " echo \"\"\n") + file(APPEND "${_summary_script}" " echo \"Failed Tests:\"\n") + file(APPEND "${_summary_script}" " echo -e \"\$FAILED_LIST\" | while IFS= read -r line; do\n") + file(APPEND "${_summary_script}" " [ -n \"\$line\" ] && echo \" - \$line\"\n") + file(APPEND "${_summary_script}" " done\n") + file(APPEND "${_summary_script}" "fi\n") + file(APPEND "${_summary_script}" "if [ -n \"\$SKIPPED_LIST\" ]; then\n") + file(APPEND "${_summary_script}" " echo \"\"\n") + file(APPEND "${_summary_script}" " echo \"Skipped Tests:\"\n") + file(APPEND "${_summary_script}" " echo -e \"\$SKIPPED_LIST\" | while IFS= read -r line; do\n") + file(APPEND "${_summary_script}" " [ -n \"\$line\" ] && echo \" - \$line\"\n") + file(APPEND "${_summary_script}" " done\n") + file(APPEND "${_summary_script}" "fi\n") + file(APPEND "${_summary_script}" "if [ -n \"\$CRASHED_LIST\" ]; then\n") + file(APPEND "${_summary_script}" " echo \"\"\n") + file(APPEND "${_summary_script}" " echo \"Crashed/Error Tests:\"\n") + file(APPEND "${_summary_script}" " echo -e \"\$CRASHED_LIST\" | while IFS= read -r line; do\n") + file(APPEND "${_summary_script}" " [ -n \"\$line\" ] && echo \" - \$line\"\n") + file(APPEND "${_summary_script}" " done\n") + file(APPEND "${_summary_script}" "fi\n") + file(APPEND "${_summary_script}" "echo \"========================================\"\n") + execute_process(COMMAND chmod +x "${_summary_script}") + + # Create wrapper script that runs LIT then shows summary + set(_lit_wrapper "${CMAKE_CURRENT_BINARY_DIR}/catch_tests/lit_wrapper_${CATEGORY}_${SUBDIR}_${VARIANT_SUFFIX}.sh") + file(WRITE "${_lit_wrapper}" "#!/bin/bash\n") + file(APPEND "${_lit_wrapper}" "cd ${CMAKE_CURRENT_BINARY_DIR}\n") + string(REPLACE ";" " " _test_list_str "${_subdir_variant_tests}") + file(APPEND "${_lit_wrapper}" "${TEST_SUITE_LIT} ${TEST_SUITE_LIT_FLAGS} ${_lit_verbosity_flags} ${_test_list_str}\n") + file(APPEND "${_lit_wrapper}" "LIT_EXIT=\$?\n") + file(APPEND "${_lit_wrapper}" "${_summary_script}\n") + file(APPEND "${_lit_wrapper}" "exit \$LIT_EXIT\n") + execute_process(COMMAND chmod +x "${_lit_wrapper}") + + add_custom_target(check-hip-catch-${CATEGORY}-${SUBDIR}-${VARIANT_SUFFIX} + COMMAND ${_lit_wrapper} + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} + DEPENDS hip-tests-catch-${CATEGORY}-${SUBDIR}-${VARIANT_SUFFIX} + USES_TERMINAL + COMMENT "Run HIP Catch ${CATEGORY}/${SUBDIR} tests for variant ${VARIANT_SUFFIX}") + + # Ensure litsupport files (including lit.cfg) and timeit-target are available + if(TARGET build-litsupport) + add_dependencies(check-hip-catch-${CATEGORY}-${SUBDIR}-${VARIANT_SUFFIX} build-litsupport) + endif() + if(TARGET timeit-target) + add_dependencies(check-hip-catch-${CATEGORY}-${SUBDIR}-${VARIANT_SUFFIX} timeit-target) + endif() + + # Wire to aggregated check target (subdirectory aggregator depends on subdirectory-variant) + add_dependencies(check-hip-catch-${CATEGORY}-${SUBDIR} check-hip-catch-${CATEGORY}-${SUBDIR}-${VARIANT_SUFFIX}) + endif() +endfunction() + +# Main function to integrate Catch tests for a specific HIP variant +# Arguments: +# VARIANT_SUFFIX - Variant suffix (e.g., hip-7.2.0) +# ROCM_PATH - Path to ROCm installation +function(integrate_catch_tests VARIANT_SUFFIX ROCM_PATH) + message(STATUS "Integrating Catch tests for variant ${VARIANT_SUFFIX}") + + # Register this variant globally + get_property(_variants GLOBAL PROPERTY CATCH_ALL_VARIANTS) + if(NOT "${VARIANT_SUFFIX}" IN_LIST _variants) + set_property(GLOBAL APPEND PROPERTY CATCH_ALL_VARIANTS "${VARIANT_SUFFIX}") + endif() + + # Create catch test target for this variant + add_custom_target(hip-tests-catch-${VARIANT_SUFFIX} + COMMENT "Build Catch tests for HIP variant ${VARIANT_SUFFIX}") + + # Process each enabled category + foreach(_category ${CATCH_TEST_CATEGORIES}) + message(STATUS "Processing Catch test category: ${_category}") + + # Create aggregated category-level targets (once per category) + get_property(_created_categories GLOBAL PROPERTY CATCH_CATEGORY_TARGETS_CREATED) + if(NOT "${_category}" IN_LIST _created_categories) + message(STATUS "Creating aggregated category targets for: ${_category}") + + # Create aggregated build target + add_custom_target(hip-tests-catch-${_category} + COMMENT "Build all HIP Catch ${_category} tests across all variants") + add_dependencies(hip-tests-catch hip-tests-catch-${_category}) + + # Create aggregated check target (will be populated later) + add_custom_target(check-hip-catch-${_category} + COMMENT "Run all HIP Catch ${_category} tests across all variants") + add_dependencies(check-hip-catch check-hip-catch-${_category}) + + # Mark as created + set_property(GLOBAL APPEND PROPERTY CATCH_CATEGORY_TARGETS_CREATED "${_category}") + endif() + + # Create per-variant category-level targets + message(STATUS "Creating per-variant category targets for: ${_category}-${VARIANT_SUFFIX}") + add_custom_target(hip-tests-catch-${_category}-${VARIANT_SUFFIX} + COMMENT "Build HIP Catch ${_category} tests for variant ${VARIANT_SUFFIX}") + + # Wire dependencies: variant-specific target to aggregated category target + add_dependencies(hip-tests-catch-${_category} hip-tests-catch-${_category}-${VARIANT_SUFFIX}) + + # Wire dependencies: top-level variant target depends on category variant target + add_dependencies(hip-tests-catch-${VARIANT_SUFFIX} hip-tests-catch-${_category}-${VARIANT_SUFFIX}) + + # Discover test subdirectories + discover_catch_test_category(${_category}) + + # Get the discovered subdirectories + set(_subdir_var "${_category}_TEST_DIRS") + if(DEFINED ${_subdir_var}) + foreach(_subdir ${${_subdir_var}}) + create_catch_tests_for_subdir("${_category}" "${_subdir}" "${VARIANT_SUFFIX}" "${ROCM_PATH}") + endforeach() + endif() + endforeach() + + # Create per-variant check targets for each category + foreach(_category ${CATCH_TEST_CATEGORIES}) + # Get test targets for this category-variant combination + get_property(_category_variant_tests GLOBAL PROPERTY "CATCH_${_category}_TEST_TARGETS_${VARIANT_SUFFIX}") + + if(_category_variant_tests) + # Category-variant target is just an aggregator - subdirectory-variant targets do the actual work + add_custom_target(check-hip-catch-${_category}-${VARIANT_SUFFIX} + DEPENDS hip-tests-catch-${_category}-${VARIANT_SUFFIX} + COMMENT "Run HIP Catch ${_category} tests for variant ${VARIANT_SUFFIX}") + + # Wire category-variant to its subdirectory-variant targets + set(_subdir_var "${_category}_TEST_DIRS") + if(DEFINED ${_subdir_var}) + foreach(_subdir ${${_subdir_var}}) + if(TARGET check-hip-catch-${_category}-${_subdir}-${VARIANT_SUFFIX}) + add_dependencies(check-hip-catch-${_category}-${VARIANT_SUFFIX} check-hip-catch-${_category}-${_subdir}-${VARIANT_SUFFIX}) + endif() + endforeach() + endif() + + # Wire to aggregated check target (will run tests from all variants) + add_dependencies(check-hip-catch-${_category} check-hip-catch-${_category}-${VARIANT_SUFFIX}) + endif() + endforeach() + + # Add variant target to main catch target + add_dependencies(hip-tests-catch hip-tests-catch-${VARIANT_SUFFIX}) + + # Create variant-level check target as aggregator (category-variant targets do the actual work) + add_custom_target(check-hip-catch-${VARIANT_SUFFIX} + DEPENDS hip-tests-catch-${VARIANT_SUFFIX} + COMMENT "Run all HIP Catch tests for variant ${VARIANT_SUFFIX}") + + # Wire variant target to category-variant targets + foreach(_category ${CATCH_TEST_CATEGORIES}) + if(TARGET check-hip-catch-${_category}-${VARIANT_SUFFIX}) + add_dependencies(check-hip-catch-${VARIANT_SUFFIX} check-hip-catch-${_category}-${VARIANT_SUFFIX}) + endif() + endforeach() + + add_dependencies(check-hip-catch check-hip-catch-${VARIANT_SUFFIX}) + + message(STATUS "Integrated ${CMAKE_CURRENT_LIST_LENGTH} Catch test targets for ${VARIANT_SUFFIX}") +endfunction() + +# Initialize the Catch test framework integration +macro(initialize_catch_tests) + if(ENABLE_HIP_CATCH_TESTS) + message(STATUS "=== Initializing HIP Catch Tests Integration ===") + + # Validate local Catch test infrastructure + validate_catch_tests_infrastructure() + + # Create main targets + add_custom_target(hip-tests-catch + COMMENT "Build all HIP Catch tests") + add_custom_target(check-hip-catch + COMMENT "Run all HIP Catch tests") + + message(STATUS "Catch test categories enabled: ${CATCH_TEST_CATEGORIES}") + if(CATCH_TEST_SUBDIRS) + message(STATUS "Catch test subdirectories filter: ${CATCH_TEST_SUBDIRS}") + else() + message(STATUS "Catch test subdirectories filter: ALL (no filter)") + endif() + message(STATUS "=== Catch Tests Integration Initialized ===") + else() + message(STATUS "HIP Catch tests are DISABLED (set ENABLE_HIP_CATCH_TESTS=ON to enable)") + endif() +endmacro() diff --git a/External/HIP/catch/external/picojson/picojson.h b/External/HIP/catch/external/picojson/picojson.h new file mode 100644 index 0000000000..7048f7caf9 --- /dev/null +++ b/External/HIP/catch/external/picojson/picojson.h @@ -0,0 +1,1141 @@ +/* + * Copyright 2009-2010 Cybozu Labs, Inc. + * Copyright 2011-2014 Kazuho Oku + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are met: + * + * 1. Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimer. + * + * 2. Redistributions in binary form must reproduce the above copyright notice, + * this list of conditions and the following disclaimer in the documentation + * and/or other materials provided with the distribution. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" + * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE + * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE + * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR + * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF + * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS + * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN + * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) + * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE + * POSSIBILITY OF SUCH DAMAGE. + */ +#ifndef picojson_h +#define picojson_h + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +// for isnan/isinf +#if __cplusplus >= 201103L +#include +#else +extern "C" { +#ifdef _MSC_VER +#include +#elif defined(__INTEL_COMPILER) +#include +#else +#include +#endif +} +#endif + +#ifndef PICOJSON_USE_RVALUE_REFERENCE +#if (defined(__cpp_rvalue_references) && __cpp_rvalue_references >= 200610) || \ + (defined(_MSC_VER) && _MSC_VER >= 1600) +#define PICOJSON_USE_RVALUE_REFERENCE 1 +#else +#define PICOJSON_USE_RVALUE_REFERENCE 0 +#endif +#endif // PICOJSON_USE_RVALUE_REFERENCE + +#ifndef PICOJSON_NOEXCEPT +#if PICOJSON_USE_RVALUE_REFERENCE +#define PICOJSON_NOEXCEPT noexcept +#else +#define PICOJSON_NOEXCEPT throw() +#endif +#endif + +// experimental support for int64_t (see README.mkdn for detail) +#ifdef PICOJSON_USE_INT64 +#define __STDC_FORMAT_MACROS +#include +#if __cplusplus >= 201103L +#include +#else +extern "C" { +#include +} +#endif +#endif + +// to disable the use of localeconv(3), set PICOJSON_USE_LOCALE to 0 +#ifndef PICOJSON_USE_LOCALE +#define PICOJSON_USE_LOCALE 1 +#endif +#if PICOJSON_USE_LOCALE +extern "C" { +#include +} +#endif + +#ifndef PICOJSON_ASSERT +#define PICOJSON_ASSERT(e) \ + do { \ + if (!(e)) throw std::runtime_error(#e); \ + } while (0) +#endif + +#ifdef _MSC_VER +#define SNPRINTF _snprintf_s +#pragma warning(push) +#pragma warning(disable : 4244) // conversion from int to char +#pragma warning(disable : 4127) // conditional expression is constant +#pragma warning(disable : 4702) // unreachable code +#pragma warning(disable : 4706) // assignment within conditional expression +#else +#define SNPRINTF snprintf +#endif + +namespace picojson { + +enum { + null_type, + boolean_type, + number_type, + string_type, + array_type, + object_type +#ifdef PICOJSON_USE_INT64 + , + int64_type +#endif +}; + +enum { INDENT_WIDTH = 2, DEFAULT_MAX_DEPTHS = 100 }; + +struct null {}; + +class value { + public: + typedef std::vector array; + typedef std::map object; + union _storage { + bool boolean_; + double number_; +#ifdef PICOJSON_USE_INT64 + int64_t int64_; +#endif + std::string* string_; + array* array_; + object* object_; + }; + + protected: + int type_; + _storage u_; + + public: + value(); + value(int type, bool); + explicit value(bool b); +#ifdef PICOJSON_USE_INT64 + explicit value(int64_t i); +#endif + explicit value(double n); + explicit value(const std::string& s); + explicit value(const array& a); + explicit value(const object& o); +#if PICOJSON_USE_RVALUE_REFERENCE + explicit value(std::string&& s); + explicit value(array&& a); + explicit value(object&& o); +#endif + explicit value(const char* s); + value(const char* s, size_t len); + ~value(); + value(const value& x); + value& operator=(const value& x); +#if PICOJSON_USE_RVALUE_REFERENCE + value(value&& x) PICOJSON_NOEXCEPT; + value& operator=(value&& x) PICOJSON_NOEXCEPT; +#endif + void swap(value& x) PICOJSON_NOEXCEPT; + template bool is() const; + template const T& get() const; + template T& get(); + template void set(const T&); +#if PICOJSON_USE_RVALUE_REFERENCE + template void set(T&&); +#endif + bool evaluate_as_boolean() const; + const value& get(const size_t idx) const; + const value& get(const std::string& key) const; + value& get(const size_t idx); + value& get(const std::string& key); + + bool contains(const size_t idx) const; + bool contains(const std::string& key) const; + std::string to_str() const; + template void serialize(Iter os, bool prettify = false) const; + std::string serialize(bool prettify = false) const; + + private: + template + value(const T*); // intentionally defined to block implicit conversion of pointer to bool + template static void _indent(Iter os, int indent); + template void _serialize(Iter os, int indent) const; + std::string _serialize(int indent) const; + void clear(); +}; + +typedef value::array array; +typedef value::object object; + +inline value::value() : type_(null_type), u_() {} + +inline value::value(int type, bool) : type_(type), u_() { + switch (type) { +#define INIT(p, v) \ + case p##type: \ + u_.p = v; \ + break + INIT(boolean_, false); + INIT(number_, 0.0); +#ifdef PICOJSON_USE_INT64 + INIT(int64_, 0); +#endif + INIT(string_, new std::string()); + INIT(array_, new array()); + INIT(object_, new object()); +#undef INIT + default: + break; + } +} + +inline value::value(bool b) : type_(boolean_type), u_() { u_.boolean_ = b; } + +#ifdef PICOJSON_USE_INT64 +inline value::value(int64_t i) : type_(int64_type), u_() { u_.int64_ = i; } +#endif + +inline value::value(double n) : type_(number_type), u_() { + if ( +#ifdef _MSC_VER + !_finite(n) +#elif __cplusplus >= 201103L + std::isnan(n) || std::isinf(n) +#else + isnan(n) || isinf(n) +#endif + ) { + throw std::overflow_error(""); + } + u_.number_ = n; +} + +inline value::value(const std::string& s) : type_(string_type), u_() { + u_.string_ = new std::string(s); +} + +inline value::value(const array& a) : type_(array_type), u_() { u_.array_ = new array(a); } + +inline value::value(const object& o) : type_(object_type), u_() { u_.object_ = new object(o); } + +#if PICOJSON_USE_RVALUE_REFERENCE +inline value::value(std::string&& s) : type_(string_type), u_() { + u_.string_ = new std::string(std::move(s)); +} + +inline value::value(array&& a) : type_(array_type), u_() { u_.array_ = new array(std::move(a)); } + +inline value::value(object&& o) : type_(object_type), u_() { + u_.object_ = new object(std::move(o)); +} +#endif + +inline value::value(const char* s) : type_(string_type), u_() { u_.string_ = new std::string(s); } + +inline value::value(const char* s, size_t len) : type_(string_type), u_() { + u_.string_ = new std::string(s, len); +} + +inline void value::clear() { + switch (type_) { +#define DEINIT(p) \ + case p##type: \ + delete u_.p; \ + break + DEINIT(string_); + DEINIT(array_); + DEINIT(object_); +#undef DEINIT + default: + break; + } +} + +inline value::~value() { clear(); } + +inline value::value(const value& x) : type_(x.type_), u_() { + switch (type_) { +#define INIT(p, v) \ + case p##type: \ + u_.p = v; \ + break + INIT(string_, new std::string(*x.u_.string_)); + INIT(array_, new array(*x.u_.array_)); + INIT(object_, new object(*x.u_.object_)); +#undef INIT + default: + u_ = x.u_; + break; + } +} + +inline value& value::operator=(const value& x) { + if (this != &x) { + value t(x); + swap(t); + } + return *this; +} + +#if PICOJSON_USE_RVALUE_REFERENCE +inline value::value(value&& x) PICOJSON_NOEXCEPT : type_(null_type), u_() { swap(x); } +inline value& value::operator=(value&& x) PICOJSON_NOEXCEPT { + swap(x); + return *this; +} +#endif +inline void value::swap(value& x) PICOJSON_NOEXCEPT { + std::swap(type_, x.type_); + std::swap(u_, x.u_); +} + +#define IS(ctype, jtype) \ + template <> inline bool value::is() const { return type_ == jtype##_type; } +IS(null, null) +IS(bool, boolean) +#ifdef PICOJSON_USE_INT64 +IS(int64_t, int64) +#endif +IS(std::string, string) +IS(array, array) +IS(object, object) +#undef IS +template <> inline bool value::is() const { + return type_ == number_type +#ifdef PICOJSON_USE_INT64 + || type_ == int64_type +#endif + ; +} + +#define GET(ctype, var) \ + template <> inline const ctype& value::get() const { \ + PICOJSON_ASSERT("type mismatch! call is() before get()" && is()); \ + return var; \ + } \ + template <> inline ctype& value::get() { \ + PICOJSON_ASSERT("type mismatch! call is() before get()" && is()); \ + return var; \ + } +GET(bool, u_.boolean_) +GET(std::string, *u_.string_) +GET(array, *u_.array_) +GET(object, *u_.object_) +#ifdef PICOJSON_USE_INT64 +GET(double, (type_ == int64_type && (const_cast(this)->type_ = number_type, + (const_cast(this)->u_.number_ = u_.int64_)), + u_.number_)) +GET(int64_t, u_.int64_) +#else +GET(double, u_.number_) +#endif +#undef GET + +#define SET(ctype, jtype, setter) \ + template <> inline void value::set(const ctype& _val) { \ + clear(); \ + type_ = jtype##_type; \ + setter \ + } +SET(bool, boolean, u_.boolean_ = _val;) +SET(std::string, string, u_.string_ = new std::string(_val);) +SET(array, array, u_.array_ = new array(_val);) +SET(object, object, u_.object_ = new object(_val);) +SET(double, number, u_.number_ = _val;) +#ifdef PICOJSON_USE_INT64 +SET(int64_t, int64, u_.int64_ = _val;) +#endif +#undef SET + +#if PICOJSON_USE_RVALUE_REFERENCE +#define MOVESET(ctype, jtype, setter) \ + template <> inline void value::set(ctype && _val) { \ + clear(); \ + type_ = jtype##_type; \ + setter \ + } +MOVESET(std::string, string, u_.string_ = new std::string(std::move(_val));) +MOVESET(array, array, u_.array_ = new array(std::move(_val));) +MOVESET(object, object, u_.object_ = new object(std::move(_val));) +#undef MOVESET +#endif + +inline bool value::evaluate_as_boolean() const { + switch (type_) { + case null_type: + return false; + case boolean_type: + return u_.boolean_; + case number_type: + return u_.number_ != 0; +#ifdef PICOJSON_USE_INT64 + case int64_type: + return u_.int64_ != 0; +#endif + case string_type: + return !u_.string_->empty(); + default: + return true; + } +} + +inline const value& value::get(const size_t idx) const { + static value s_null; + PICOJSON_ASSERT(is()); + return idx < u_.array_->size() ? (*u_.array_)[idx] : s_null; +} + +inline value& value::get(const size_t idx) { + static value s_null; + PICOJSON_ASSERT(is()); + return idx < u_.array_->size() ? (*u_.array_)[idx] : s_null; +} + +inline const value& value::get(const std::string& key) const { + static value s_null; + PICOJSON_ASSERT(is()); + object::const_iterator i = u_.object_->find(key); + return i != u_.object_->end() ? i->second : s_null; +} + +inline value& value::get(const std::string& key) { + static value s_null; + PICOJSON_ASSERT(is()); + object::iterator i = u_.object_->find(key); + return i != u_.object_->end() ? i->second : s_null; +} + +inline bool value::contains(const size_t idx) const { + PICOJSON_ASSERT(is()); + return idx < u_.array_->size(); +} + +inline bool value::contains(const std::string& key) const { + PICOJSON_ASSERT(is()); + object::const_iterator i = u_.object_->find(key); + return i != u_.object_->end(); +} + +inline std::string value::to_str() const { + switch (type_) { + case null_type: + return "null"; + case boolean_type: + return u_.boolean_ ? "true" : "false"; +#ifdef PICOJSON_USE_INT64 + case int64_type: { + char buf[sizeof("-9223372036854775808")]; + SNPRINTF(buf, sizeof(buf), "%" PRId64, u_.int64_); + return buf; + } +#endif + case number_type: { + char buf[256]; + double tmp; + SNPRINTF(buf, sizeof(buf), + fabs(u_.number_) < (1ULL << 53) && modf(u_.number_, &tmp) == 0 ? "%.f" : "%.17g", + u_.number_); +#if PICOJSON_USE_LOCALE + char* decimal_point = localeconv()->decimal_point; + if (strcmp(decimal_point, ".") != 0) { + size_t decimal_point_len = strlen(decimal_point); + for (char* p = buf; *p != '\0'; ++p) { + if (strncmp(p, decimal_point, decimal_point_len) == 0) { + return std::string(buf, p) + "." + (p + decimal_point_len); + } + } + } +#endif + return buf; + } + case string_type: + return *u_.string_; + case array_type: + return "array"; + case object_type: + return "object"; + default: + PICOJSON_ASSERT(0); +#ifdef _MSC_VER + __assume(0); +#endif + } + return std::string(); +} + +template void copy(const std::string& s, Iter oi) { + std::copy(s.begin(), s.end(), oi); +} + +template struct serialize_str_char { + Iter oi; + void operator()(char c) { + switch (c) { +#define MAP(val, sym) \ + case val: \ + copy(sym, oi); \ + break + MAP('"', "\\\""); + MAP('\\', "\\\\"); + MAP('/', "\\/"); + MAP('\b', "\\b"); + MAP('\f', "\\f"); + MAP('\n', "\\n"); + MAP('\r', "\\r"); + MAP('\t', "\\t"); +#undef MAP + default: + if (static_cast(c) < 0x20 || c == 0x7f) { + char buf[7]; + SNPRINTF(buf, sizeof(buf), "\\u%04x", c & 0xff); + copy(buf, buf + 6, oi); + } else { + *oi++ = c; + } + break; + } + } +}; + +template void serialize_str(const std::string& s, Iter oi) { + *oi++ = '"'; + serialize_str_char process_char = {oi}; + std::for_each(s.begin(), s.end(), process_char); + *oi++ = '"'; +} + +template void value::serialize(Iter oi, bool prettify) const { + return _serialize(oi, prettify ? 0 : -1); +} + +inline std::string value::serialize(bool prettify) const { return _serialize(prettify ? 0 : -1); } + +template void value::_indent(Iter oi, int indent) { + *oi++ = '\n'; + for (int i = 0; i < indent * INDENT_WIDTH; ++i) { + *oi++ = ' '; + } +} + +template void value::_serialize(Iter oi, int indent) const { + switch (type_) { + case string_type: + serialize_str(*u_.string_, oi); + break; + case array_type: { + *oi++ = '['; + if (indent != -1) { + ++indent; + } + for (array::const_iterator i = u_.array_->begin(); i != u_.array_->end(); ++i) { + if (i != u_.array_->begin()) { + *oi++ = ','; + } + if (indent != -1) { + _indent(oi, indent); + } + i->_serialize(oi, indent); + } + if (indent != -1) { + --indent; + if (!u_.array_->empty()) { + _indent(oi, indent); + } + } + *oi++ = ']'; + break; + } + case object_type: { + *oi++ = '{'; + if (indent != -1) { + ++indent; + } + for (object::const_iterator i = u_.object_->begin(); i != u_.object_->end(); ++i) { + if (i != u_.object_->begin()) { + *oi++ = ','; + } + if (indent != -1) { + _indent(oi, indent); + } + serialize_str(i->first, oi); + *oi++ = ':'; + if (indent != -1) { + *oi++ = ' '; + } + i->second._serialize(oi, indent); + } + if (indent != -1) { + --indent; + if (!u_.object_->empty()) { + _indent(oi, indent); + } + } + *oi++ = '}'; + break; + } + default: + copy(to_str(), oi); + break; + } + if (indent == 0) { + *oi++ = '\n'; + } +} + +inline std::string value::_serialize(int indent) const { + std::string s; + _serialize(std::back_inserter(s), indent); + return s; +} + +template class input { + protected: + Iter cur_, end_; + bool consumed_; + int line_; + + public: + input(const Iter& first, const Iter& last) + : cur_(first), end_(last), consumed_(false), line_(1) {} + int getc() { + if (consumed_) { + if (*cur_ == '\n') { + ++line_; + } + ++cur_; + } + if (cur_ == end_) { + consumed_ = false; + return -1; + } + consumed_ = true; + return *cur_ & 0xff; + } + void ungetc() { consumed_ = false; } + Iter cur() const { + if (consumed_) { + input* self = const_cast*>(this); + self->consumed_ = false; + ++self->cur_; + } + return cur_; + } + int line() const { return line_; } + void skip_ws() { + while (1) { + int ch = getc(); + if (!(ch == ' ' || ch == '\t' || ch == '\n' || ch == '\r')) { + ungetc(); + break; + } + } + } + bool expect(const int expected) { + skip_ws(); + if (getc() != expected) { + ungetc(); + return false; + } + return true; + } + bool match(const std::string& pattern) { + for (std::string::const_iterator pi(pattern.begin()); pi != pattern.end(); ++pi) { + if (getc() != *pi) { + ungetc(); + return false; + } + } + return true; + } +}; + +template inline int _parse_quadhex(input& in) { + int uni_ch = 0, hex; + for (int i = 0; i < 4; i++) { + if ((hex = in.getc()) == -1) { + return -1; + } + if ('0' <= hex && hex <= '9') { + hex -= '0'; + } else if ('A' <= hex && hex <= 'F') { + hex -= 'A' - 0xa; + } else if ('a' <= hex && hex <= 'f') { + hex -= 'a' - 0xa; + } else { + in.ungetc(); + return -1; + } + uni_ch = uni_ch * 16 + hex; + } + return uni_ch; +} + +template +inline bool _parse_codepoint(String& out, input& in) { + int uni_ch; + if ((uni_ch = _parse_quadhex(in)) == -1) { + return false; + } + if (0xd800 <= uni_ch && uni_ch <= 0xdfff) { + if (0xdc00 <= uni_ch) { + // a second 16-bit of a surrogate pair appeared + return false; + } + // first 16-bit of surrogate pair, get the next one + if (in.getc() != '\\' || in.getc() != 'u') { + in.ungetc(); + return false; + } + int second = _parse_quadhex(in); + if (!(0xdc00 <= second && second <= 0xdfff)) { + return false; + } + uni_ch = ((uni_ch - 0xd800) << 10) | ((second - 0xdc00) & 0x3ff); + uni_ch += 0x10000; + } + if (uni_ch < 0x80) { + out.push_back(static_cast(uni_ch)); + } else { + if (uni_ch < 0x800) { + out.push_back(static_cast(0xc0 | (uni_ch >> 6))); + } else { + if (uni_ch < 0x10000) { + out.push_back(static_cast(0xe0 | (uni_ch >> 12))); + } else { + out.push_back(static_cast(0xf0 | (uni_ch >> 18))); + out.push_back(static_cast(0x80 | ((uni_ch >> 12) & 0x3f))); + } + out.push_back(static_cast(0x80 | ((uni_ch >> 6) & 0x3f))); + } + out.push_back(static_cast(0x80 | (uni_ch & 0x3f))); + } + return true; +} + +template inline bool _parse_string(String& out, input& in) { + while (1) { + int ch = in.getc(); + if (ch < ' ') { + in.ungetc(); + return false; + } else if (ch == '"') { + return true; + } else if (ch == '\\') { + if ((ch = in.getc()) == -1) { + return false; + } + switch (ch) { +#define MAP(sym, val) \ + case sym: \ + out.push_back(val); \ + break + MAP('"', '\"'); + MAP('\\', '\\'); + MAP('/', '/'); + MAP('b', '\b'); + MAP('f', '\f'); + MAP('n', '\n'); + MAP('r', '\r'); + MAP('t', '\t'); +#undef MAP + case 'u': + if (!_parse_codepoint(out, in)) { + return false; + } + break; + default: + return false; + } + } else { + out.push_back(static_cast(ch)); + } + } + return false; +} + +template inline bool _parse_array(Context& ctx, input& in) { + if (!ctx.parse_array_start()) { + return false; + } + size_t idx = 0; + if (in.expect(']')) { + return ctx.parse_array_stop(idx); + } + do { + if (!ctx.parse_array_item(in, idx)) { + return false; + } + idx++; + } while (in.expect(',')); + return in.expect(']') && ctx.parse_array_stop(idx); +} + +template +inline bool _parse_object(Context& ctx, input& in) { + if (!ctx.parse_object_start()) { + return false; + } + if (in.expect('}')) { + return ctx.parse_object_stop(); + } + do { + std::string key; + if (!in.expect('"') || !_parse_string(key, in) || !in.expect(':')) { + return false; + } + if (!ctx.parse_object_item(in, key)) { + return false; + } + } while (in.expect(',')); + return in.expect('}') && ctx.parse_object_stop(); +} + +template inline std::string _parse_number(input& in) { + std::string num_str; + while (1) { + int ch = in.getc(); + if (('0' <= ch && ch <= '9') || ch == '+' || ch == '-' || ch == 'e' || ch == 'E') { + num_str.push_back(static_cast(ch)); + } else if (ch == '.') { +#if PICOJSON_USE_LOCALE + num_str += localeconv()->decimal_point; +#else + num_str.push_back('.'); +#endif + } else { + in.ungetc(); + break; + } + } + return num_str; +} + +template inline bool _parse(Context& ctx, input& in) { + in.skip_ws(); + int ch = in.getc(); + switch (ch) { +#define IS(ch, text, op) \ + case ch: \ + if (in.match(text) && op) { \ + return true; \ + } else { \ + return false; \ + } + IS('n', "ull", ctx.set_null()); + IS('f', "alse", ctx.set_bool(false)); + IS('t', "rue", ctx.set_bool(true)); +#undef IS + case '"': + return ctx.parse_string(in); + case '[': + return _parse_array(ctx, in); + case '{': + return _parse_object(ctx, in); + default: + if (('0' <= ch && ch <= '9') || ch == '-') { + double f; + char* endp; + in.ungetc(); + std::string num_str(_parse_number(in)); + if (num_str.empty()) { + return false; + } +#ifdef PICOJSON_USE_INT64 + { + errno = 0; + intmax_t ival = strtoimax(num_str.c_str(), &endp, 10); + if (errno == 0 && std::numeric_limits::min() <= ival && + ival <= std::numeric_limits::max() && + endp == num_str.c_str() + num_str.size()) { + ctx.set_int64(ival); + return true; + } + } +#endif + f = strtod(num_str.c_str(), &endp); + if (endp == num_str.c_str() + num_str.size()) { + ctx.set_number(f); + return true; + } + return false; + } + break; + } + in.ungetc(); + return false; +} + +class deny_parse_context { + public: + bool set_null() { return false; } + bool set_bool(bool) { return false; } +#ifdef PICOJSON_USE_INT64 + bool set_int64(int64_t) { return false; } +#endif + bool set_number(double) { return false; } + template bool parse_string(input&) { return false; } + bool parse_array_start() { return false; } + template bool parse_array_item(input&, size_t) { return false; } + bool parse_array_stop(size_t) { return false; } + bool parse_object_start() { return false; } + template bool parse_object_item(input&, const std::string&) { + return false; + } +}; + +class default_parse_context { + protected: + value* out_; + size_t depths_; + + public: + default_parse_context(value* out, size_t depths = DEFAULT_MAX_DEPTHS) + : out_(out), depths_(depths) {} + bool set_null() { + *out_ = value(); + return true; + } + bool set_bool(bool b) { + *out_ = value(b); + return true; + } +#ifdef PICOJSON_USE_INT64 + bool set_int64(int64_t i) { + *out_ = value(i); + return true; + } +#endif + bool set_number(double f) { + *out_ = value(f); + return true; + } + template bool parse_string(input& in) { + *out_ = value(string_type, false); + return _parse_string(out_->get(), in); + } + bool parse_array_start() { + if (depths_ == 0) return false; + --depths_; + *out_ = value(array_type, false); + return true; + } + template bool parse_array_item(input& in, size_t) { + array& a = out_->get(); + a.push_back(value()); + default_parse_context ctx(&a.back(), depths_); + return _parse(ctx, in); + } + bool parse_array_stop(size_t) { + ++depths_; + return true; + } + bool parse_object_start() { + if (depths_ == 0) return false; + *out_ = value(object_type, false); + return true; + } + template bool parse_object_item(input& in, const std::string& key) { + object& o = out_->get(); + default_parse_context ctx(&o[key], depths_); + return _parse(ctx, in); + } + bool parse_object_stop() { + ++depths_; + return true; + } + + private: + default_parse_context(const default_parse_context&); + default_parse_context& operator=(const default_parse_context&); +}; + +class null_parse_context { + protected: + size_t depths_; + + public: + struct dummy_str { + void push_back(int) {} + }; + + public: + null_parse_context(size_t depths = DEFAULT_MAX_DEPTHS) : depths_(depths) {} + bool set_null() { return true; } + bool set_bool(bool) { return true; } +#ifdef PICOJSON_USE_INT64 + bool set_int64(int64_t) { return true; } +#endif + bool set_number(double) { return true; } + template bool parse_string(input& in) { + dummy_str s; + return _parse_string(s, in); + } + bool parse_array_start() { + if (depths_ == 0) return false; + --depths_; + return true; + } + template bool parse_array_item(input& in, size_t) { + return _parse(*this, in); + } + bool parse_array_stop(size_t) { + ++depths_; + return true; + } + bool parse_object_start() { + if (depths_ == 0) return false; + --depths_; + return true; + } + template bool parse_object_item(input& in, const std::string&) { + ++depths_; + return _parse(*this, in); + } + bool parse_object_stop() { return true; } + + private: + null_parse_context(const null_parse_context&); + null_parse_context& operator=(const null_parse_context&); +}; + +// obsolete, use the version below +template inline std::string parse(value& out, Iter& pos, const Iter& last) { + std::string err; + pos = parse(out, pos, last, &err); + return err; +} + +template +inline Iter _parse(Context& ctx, const Iter& first, const Iter& last, std::string* err) { + input in(first, last); + if (!_parse(ctx, in) && err != NULL) { + char buf[64]; + SNPRINTF(buf, sizeof(buf), "syntax error at line %d near: ", in.line()); + *err = buf; + while (1) { + int ch = in.getc(); + if (ch == -1 || ch == '\n') { + break; + } else if (ch >= ' ') { + err->push_back(static_cast(ch)); + } + } + } + return in.cur(); +} + +template +inline Iter parse(value& out, const Iter& first, const Iter& last, std::string* err) { + default_parse_context ctx(&out); + return _parse(ctx, first, last, err); +} + +inline std::string parse(value& out, const std::string& s) { + std::string err; + parse(out, s.begin(), s.end(), &err); + return err; +} + +inline std::string parse(value& out, std::istream& is) { + std::string err; + parse(out, std::istreambuf_iterator(is.rdbuf()), std::istreambuf_iterator(), &err); + return err; +} + +template struct last_error_t { + static std::string s; +}; +template std::string last_error_t::s; + +inline void set_last_error(const std::string& s) { last_error_t::s = s; } + +inline const std::string& get_last_error() { return last_error_t::s; } + +inline bool operator==(const value& x, const value& y) { + if (x.is()) return y.is(); +#define PICOJSON_CMP(type) \ + if (x.is()) return y.is() && x.get() == y.get() + PICOJSON_CMP(bool); + PICOJSON_CMP(double); + PICOJSON_CMP(std::string); + PICOJSON_CMP(array); + PICOJSON_CMP(object); +#undef PICOJSON_CMP + PICOJSON_ASSERT(0); +#ifdef _MSC_VER + __assume(0); +#endif + return false; +} + +inline bool operator!=(const value& x, const value& y) { return !(x == y); } +} // namespace picojson + +#if !PICOJSON_USE_RVALUE_REFERENCE +namespace std { +template <> inline void swap(picojson::value& x, picojson::value& y) { x.swap(y); } +} // namespace std +#endif + +inline std::istream& operator>>(std::istream& is, picojson::value& x) { + picojson::set_last_error(std::string()); + const std::string err(picojson::parse(x, is)); + if (!err.empty()) { + picojson::set_last_error(err); + is.setstate(std::ios::failbit); + } + return is; +} + +inline std::ostream& operator<<(std::ostream& os, const picojson::value& x) { + x.serialize(std::ostream_iterator(os)); + return os; +} +#ifdef _MSC_VER +#pragma warning(pop) +#endif + +#endif diff --git a/External/HIP/catch/hipTestMain/hip_test_context.cc b/External/HIP/catch/hipTestMain/hip_test_context.cc new file mode 100644 index 0000000000..840a08e450 --- /dev/null +++ b/External/HIP/catch/hipTestMain/hip_test_context.cc @@ -0,0 +1,362 @@ +/* + * Copyright (C) Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a + * copy of this software and associated documentation files (the "Software"), + * to deal in the Software without restriction, including without limitation + * the rights to use, copy, modify, merge, publish, distribute, sublicense, + * and/or sell copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included + * in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS + * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE COPYRIGHT HOLDER(S) BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN + * AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN + * CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + */ + +#include +#include +#include +#include +#include +#include +#include "hip_test_context.hh" +#include "hip_test_filesystem.hh" +#include "hip_test_features.hh" + +void TestContext::detectOS() { +#if (HT_WIN == 1) + p_windows = true; +#elif (HT_LINUX == 1) + p_linux = true; +#endif +} + +void TestContext::detectPlatform() { +#if (HT_AMD == 1) + amd = true; +#elif (HT_NVIDIA == 1) + nvidia = true; +#endif +} + +std::string TestContext::substringFound(std::vector list, std::string filename) { + std::string match = ""; + for (unsigned int i = 0; i < list.size(); i++) { + if (filename.find(list.at(i)) != std::string::npos) { + match = list.at(i); + break; + } + } + return match; +} + +std::string TestContext::getCurrentArch() { +#if HT_LINUX + const char* cmd = + "/opt/rocm/bin/rocm_agent_enumerator | awk '$0 != \"gfx000\"' | xargs | sed -e 's/ /;/g' | " + "tr -d '\n'"; + std::array buffer; + std::string result; + std::unique_ptr pipe(popen(cmd, "r"), pclose); + if (!pipe) { + printf("popen() failed!"); + return ""; + } + while (fgets(buffer.data(), buffer.size(), pipe.get()) != nullptr) { + std::string res = buffer.data(); + result = res; + } + + std::string s_visible_devices = TestContext::getEnvVar("HIP_VISIBLE_DEVICES"); + + auto parser = [](std::string input, char c) -> std::vector { + std::vector ret; + auto loc = input.find(c); + while (loc != std::string::npos) { + auto t_str = input.substr(0, loc); + ret.push_back(t_str); + input.erase(0, loc + 1); + loc = input.find(c); + } + if (input.size() > 0) { + ret.push_back(input); + } + return ret; + }; + + std::vector archs = parser(result, ';'); + std::vector v_visible_devices = parser(s_visible_devices, ','); + std::vector visible_devices; + std::for_each(v_visible_devices.begin(), v_visible_devices.end(), + [&](const std::string& in) { visible_devices.push_back(std::stoi(in)); }); + + if (archs.size() == 0) { + return ""; // rocm_agent_enum gave us garbage + } + + auto first_arch = archs[0]; + if (!std::all_of(archs.begin(), archs.end(), + [&](const std::string& in) { return in == first_arch; })) { + // We have multiple archs in rocm_agent_enum + // Check if they are same or not by applying HIP_VISIBLE_DEVICES filter + std::vector filtered_archs; + if (visible_devices.size() > 0) { + for (size_t i = 0; i < visible_devices.size(); i++) { + filtered_archs.push_back(archs[visible_devices[i]]); + } + } else { + filtered_archs = archs; + } + auto first_filtered_arch = filtered_archs[0]; + if (!std::all_of(filtered_archs.begin(), filtered_archs.end(), + [&](const std::string& in) { return in == first_filtered_arch; })) { + LogPrintf("%s", + "[ERROR] Cannot run tests on Hetrogenous Architecture. Please set " + "HIP_VISIBLE_DEVICES with devices of same arch"); + std::abort(); + } + return first_filtered_arch; + } + return first_arch; +#else + return ""; +#endif +} + +std::string TestContext::getMatchingConfigFile(std::string config_dir) { + std::string configFileToUse = ""; + if (isLinux() && isAmd()) { + std::string cur_arch = getCurrentArch(); + LogPrintf("The arch present: %s", cur_arch.c_str()); + configFileToUse = config_dir + "/config_" + getConfig().platform + "_" + getConfig().os + "_" + + cur_arch + ".json"; + } else { + configFileToUse = + config_dir + "/config_" + getConfig().platform + "_" + getConfig().os + ".json"; + } + if (fs::exists(configFileToUse)) { + return configFileToUse; + } + return ""; +} + +std::string& TestContext::getCommonJsonFile() { + fs::path config_dir = exe_path; + config_dir = config_dir.parent_path(); + int levels = 0; + bool configFolderFound = false; + std::vector configList; + std::string configFile; + // check a max of 5 levels down the executable path + while (levels < 5) { + fs::path temp_path = config_dir; + temp_path /= "config"; + if (fs::exists(temp_path)) { + config_dir = fs::absolute(temp_path); + configFolderFound = true; + break; + } else { + config_dir = config_dir.parent_path(); + levels++; + } + } + + // get config.json files if config folder. + if (configFolderFound) { + json_file_ = getMatchingConfigFile(config_dir.string()); + } + return json_file_; +} + + +void TestContext::getConfigFiles() { + config_.platform = (amd ? "amd" : (nvidia ? "nvidia" : "unknown")); + config_.os = (p_windows ? "windows" : (p_linux ? "linux" : "unknown")); + + if (config_.os == "unknown" || config_.platform == "unknown") { + LogPrintf("%s", "Either Config or Os is unknown, this wont end well"); + abort(); + } + + std::string env_config = TestContext::getEnvVar("HIP_CATCH_EXCLUDE_FILE"); + LogPrintf("Env Config file: %s", (!env_config.empty()) ? env_config.c_str() : "Not found"); + // HIP_CATCH_EXCLUDE_FILE is set for custom file path + if (!env_config.empty()) { + if (fs::exists(env_config)) { + config_.json_files.push_back(env_config); + } + } else { + std::string jsonFile = getCommonJsonFile(); + // get common json file + if (jsonFile != "") { + config_.json_files.push_back(getCommonJsonFile()); + } + } + + for (const auto& fl : config_.json_files) { + LogPrintf("Config file path: %s", fl.c_str()); + } +} + +TestContext::TestContext(int argc, char** argv) { + detectOS(); + detectPlatform(); + setExePath(argc, argv); + getConfigFiles(); + parseJsonFiles(); + parseOptions(argc, argv); +} + +void TestContext::setExePath(int argc, char** argv) { + if (argc == 0) return; + fs::path p = std::string(argv[0]); + if (p.has_filename()) p.remove_filename(); + exe_path = p.string(); +} + +bool TestContext::isWindows() const { return p_windows; } +bool TestContext::isLinux() const { return p_linux; } + +bool TestContext::isNvidia() const { return nvidia; } +bool TestContext::isAmd() const { return amd; } + +void TestContext::parseOptions(int argc, char** argv) { + // Test name is at [1] position + if (argc != 2) return; + current_test = std::string(argv[1]); +} + +bool TestContext::skipTest() const { + // Direct Match + auto flags = std::regex::ECMAScript; + for (const auto& i : skip_test) { + auto regex = std::regex(i.c_str(), flags); + if (std::regex_match(current_test, regex)) { + return true; + } + } + // TODO add test case skip as well + return false; +} + +std::string TestContext::currentPath() const { return fs::current_path().string(); } + +bool TestContext::parseJsonFiles() { + // Check if file exists + for (const auto& fl : config_.json_files) { + if (!fs::exists(fl)) { + LogPrintf("Unable to find the file: %s", fl.c_str()); + return true; + } + // Open the file + std::ifstream js_file(fl); + std::string json_str((std::istreambuf_iterator(js_file)), + std::istreambuf_iterator()); + LogPrintf("Json contents:: %s", json_str.data()); + + picojson::value v; + std::string err = picojson::parse(v, json_str); + if (err.size() > 1) { + LogPrintf("Error from PicoJson: %s", err.data()); + return false; + } + + if (!v.is()) { + LogPrintf("%s", "Data in json is not in correct format, it should be an object"); + return false; + } + + const picojson::object& o = v.get(); + for (picojson::object::const_iterator i = o.begin(); i != o.end(); ++i) { + // Processing for DisabledTests + if (i->first == "DisabledTests") { + // Value should contain list of values + if (!i->second.is()) return false; + + auto& val = i->second.get(); + for (auto ai = val.begin(); ai != val.end(); ai++) { + std::string tmp = ai->get(); + std::string newRegexName; + for (const auto& c : tmp) { + if (c == '*') + newRegexName += ".*"; + else + newRegexName += c; + } + skip_test.insert(newRegexName); + } + } + } + } + return true; +} + +void TestContext::cleanContext() { + for (auto& pair : compiledKernels) { + hipError_t error = hipModuleUnload(pair.second.module); + if (error != hipSuccess) { + throw std::runtime_error("Unable to unload rtc module"); + } + } +} + +void TestContext::trackRtcState(std::string kernelNameExpression, hipModule_t loadedModule, + hipFunction_t kernelFunction) { + rtcState state{loadedModule, kernelFunction}; + compiledKernels[kernelNameExpression] = state; +} + +hipFunction_t TestContext::getFunction(const std::string kernelNameExpression) { + auto it{compiledKernels.find(kernelNameExpression)}; + + if (it != compiledKernels.end()) { + return it->second.kernelFunction; + } else { + return nullptr; + } +} + +void TestContext::addResults(HCResult r) { + std::unique_lock lock(resultMutex); + results.push_back(r); + if ((!r.conditionsResult) || + ((r.result != hipSuccess) && (r.result != hipErrorPeerAccessAlreadyEnabled))) { + hasErrorOccured_.store(true); + } +} + +void TestContext::finalizeResults() { + std::unique_lock lock(resultMutex); + // clear the results whatever happens + std::shared_ptr emptyVec(nullptr, [this](auto) { results.clear(); }); + + for (const auto& i : results) { + INFO("HIP API Result check\n File:: " + << i.file << "\n Line:: " << i.line << "\n API:: " << i.call + << "\n Result:: " << i.result << "\n Result Str:: " << hipGetErrorString(i.result)); + REQUIRE(((i.result == hipSuccess) || (i.result == hipErrorPeerAccessAlreadyEnabled) || + (i.result == hipErrorNotSupported))); + REQUIRE(i.conditionsResult); + } + hasErrorOccured_.store(false); // Clear the flag +} + +bool TestContext::hasErrorOccured() { return hasErrorOccured_.load(); } + +TestContext::~TestContext() { + // Show this message when there are unchecked results + if (results.size() != 0) { + std::cerr << "HIP_CHECK_THREAD_FINALIZE() has not been called after HIP_CHECK_THREAD\n" + << "Please call HIP_CHECK_THREAD_FINALIZE after joining threads\n" + << "There is/are " << results.size() << " unchecked results from threads." + << std::endl; + std::abort(); // Crash to bring users attention to this message and avoid accidental passing of + // tests without checking for errors + } +} diff --git a/External/HIP/catch/hipTestMain/hip_test_features.cc b/External/HIP/catch/hipTestMain/hip_test_features.cc new file mode 100644 index 0000000000..bb6d23cee6 --- /dev/null +++ b/External/HIP/catch/hipTestMain/hip_test_features.cc @@ -0,0 +1,159 @@ +/* + * Copyright (C) Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a + * copy of this software and associated documentation files (the "Software"), + * to deal in the Software without restriction, including without limitation + * the rights to use, copy, modify, merge, publish, distribute, sublicense, + * and/or sell copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included + * in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS + * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE COPYRIGHT HOLDER(S) BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN + * AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN + * CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + */ + +#include "hip_test_features.hh" + +#include +#include +#include +#include "hip_test_context.hh" + +std::vector> GCNArchFeatMap = { + {"gfx90a", "gfx942", "gfx950"}, // CT_FEATURE_FINEGRAIN_HWSUPPORT + {"gfx90a", "gfx942", "gfx950"}, // CT_FEATURE_HMM + {"gfx90a", "gfx942", "gfx950"}, // CT_FEATURE_TEXTURES_NOT_SUPPORTED +}; + +#if HT_AMD +std::string TrimAndGetGFXName(const std::string& full_gfx_name) { + std::string gfx_name(""); + + // Split the first part of the delimiter + std::string delimiter = ":"; + auto pos = full_gfx_name.find(delimiter); + if (pos == std::string::npos) { + gfx_name = full_gfx_name; + } else { + gfx_name = full_gfx_name.substr(0, pos); + } + + assert(gfx_name.substr(0, 3) == "gfx"); + return gfx_name; +} +#endif + +// Check if the GCN Maps +bool CheckIfFeatSupported(enum CTFeatures test_feat, std::string gcn_arch) { +#if HT_NVIDIA + return true; // returning true since feature check does not exist for NV. +#elif HT_AMD + assert(test_feat >= 0 && test_feat < CTFeatures::CT_FEATURE_LAST); + gcn_arch = TrimAndGetGFXName(gcn_arch); + assert(gcn_arch != ""); + return (GCNArchFeatMap[test_feat].find(gcn_arch) != GCNArchFeatMap[test_feat].cend()); +#else + std::cout << "Platform has to be either AMD or NVIDIA, asserting..." << std::endl; + assert(false); +#endif +} + +// Return true if agentTarget has corresponding generic target which will be returned in +// genericTarget; +// false, otherwise. +// Note: it will naturely return false on Nvidia device +bool getGenericTarget(const std::string& agentTarget, std::string& genericTarget) { + // The map is subject to change per removing policy + static std::map genericTargetMap{ + // "gfx9-generic" + {"gfx900", "gfx9-generic"}, + {"gfx902", "gfx9-generic"}, + {"gfx904", "gfx9-generic"}, + {"gfx906", "gfx9-generic"}, + {"gfx909", "gfx9-generic"}, + {"gfx90c", "gfx9-generic"}, + // "gfx9-4-generic + {"gfx940", "gfx9-4-generic"}, + {"gfx941", "gfx9-4-generic"}, + {"gfx942", "gfx9-4-generic"}, + {"gfx950", "gfx9-4-generic"}, + // "gfx10-1-generic" + {"gfx1010", "gfx10-1-generic"}, + {"gfx1011", "gfx10-1-generic"}, + {"gfx1012", "gfx10-1-generic"}, + {"gfx1013", "gfx10-1-generic"}, + // "gfx10-3-generic" + {"gfx1030", "gfx10-3-generic"}, + {"gfx1031", "gfx10-3-generic"}, + {"gfx1032", "gfx10-3-generic"}, + {"gfx1033", "gfx10-3-generic"}, + {"gfx1034", "gfx10-3-generic"}, + {"gfx1035", "gfx10-3-generic"}, + {"gfx1036", "gfx10-3-generic"}, + // "gfx11-generic" + {"gfx1100", "gfx11-generic"}, + {"gfx1101", "gfx11-generic"}, + {"gfx1102", "gfx11-generic"}, + {"gfx1103", "gfx11-generic"}, + {"gfx1150", "gfx11-generic"}, + {"gfx1151", "gfx11-generic"}, + // "gfx12-generic" + {"gfx1200", "gfx12-generic"}, + {"gfx1201", "gfx12-generic"}, + }; + auto search = genericTargetMap.find(agentTarget); + if (search == genericTargetMap.end()) return false; + genericTarget = search->second; + return true; +} + +/* +Return true, if gcnArchName has corresponding generic target; + false, otherwise. +If gcnArchName is nullptr, it will be queried from deviceId; + otherwise, deviceId will be ignored. + +The specific arches have the following mapping to generic targets, + +Generic GFX11 + +--offload-arch=gfx11-generic - includes [gfx1100-gfx1103], gfx1150, gfx1151 + +Generic GFX10.3 + +--offload-arch=gfx10.3-generic - includes [gfx1030-gfx1036] + +Generic GFX10.1 + +--offload-arch=gfx10.1-generic - includes [gfx1010-gfx1013] + +Generic GFX9 / Consumer + +--offload-arch=gfx9-generic - includes gfx900, gfx902, gfx904, gfx906, gfx909, gfx90c + +Generic GFX9.4 / Data center + +--offload-arch=gfx9-4-generic - includes gfx940, gfx941, gfx942, gfx950 +*/ +bool isGenericTargetSupported(char* gcnArchName, int deviceId) { + hipDeviceProp_t props{}; + if (gcnArchName == nullptr) { + if (hipGetDeviceProperties(&props, deviceId) != hipSuccess) return false; + gcnArchName = props.gcnArchName; + } + std::string target{gcnArchName}; + std::string genericTarget{}; + auto pos = target.find(':'); + if (pos != std::string::npos) { + target[pos] = 0; + target.resize(pos); + } + return getGenericTarget(target, genericTarget); +} diff --git a/External/HIP/catch/hipTestMain/main.cc b/External/HIP/catch/hipTestMain/main.cc new file mode 100644 index 0000000000..a142992cc9 --- /dev/null +++ b/External/HIP/catch/hipTestMain/main.cc @@ -0,0 +1,77 @@ +/* + * Copyright (C) Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a + * copy of this software and associated documentation files (the "Software"), + * to deal in the Software without restriction, including without limitation + * the rights to use, copy, modify, merge, publish, distribute, sublicense, + * and/or sell copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included + * in all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS + * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE COPYRIGHT HOLDER(S) BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN + * AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN + * CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + */ + +#define CATCH_CONFIG_RUNNER +#include +#include +#include + +CmdOptions cmd_options; + +int main(int argc, char** argv) { + auto& context = TestContext::get(argc, argv); + if (context.skipTest()) { + // CTest uses this regex to figure out if the test has been skipped + std::cout << "HIP_SKIP_THIS_TEST" << std::endl; + return 0; + } + + Catch::Session session; + + using namespace Catch::clara; + // clang-format off + auto cli = session.cli() + | Opt(cmd_options.iterations, "iterations") + ["-I"]["--iterations"] + ("Number of iterations used for performance tests (default: 1000)") + | Opt(cmd_options.warmups, "warmups") + ["-W"]["--warmups"] + ("Number of warmup iterations used for performance tests (default: 100)") + | Opt(cmd_options.no_display) + ["-S"]["--no-display"] + ("Do not display the output of performance tests") + | Opt(cmd_options.progress) + ["-P"]["--progress"] + ("Show progress bar when running performance tests") + | Opt(cmd_options.cg_iterations, "cg_iterations") + ["-C"]["--cg-iterations"] + ("Number of iterations used for cooperative groups sync tests (default: 5)") + | Opt(cmd_options.accuracy_iterations, "accuracy_iterations") + ["-A"]["--accuracy-iterations"] + ("Number of iterations used for math accuracy tests with randomly generated inputs (default: 2^32)") + | Opt(cmd_options.accuracy_max_memory, "accuracy_max_memory") + ["-M"]["--accuracy-max-memory"] + ("Percentage of global device memory allowed for math accuracy tests (default: 80%)") + | Opt(cmd_options.reduce_iterations, "reduce_iterations") + ["-R"]["--reduce-iterations"] + ("Number of iterations for fuzzing reduce operations (default: 1)") + | Opt(cmd_options.reduce_input_size, "reduce_input_size") + ["-Z"]["--reduce-input-size"] + ("Size of the input for the reduce sync operations performance test (megabytes) (default: 50)") + ; + // clang-format on + + session.cli(cli); + + int out = session.run(argc, argv); + TestContext::get().cleanContext(); + return out; +} diff --git a/External/HIP/catch/include/cmd_options.hh b/External/HIP/catch/include/cmd_options.hh new file mode 100644 index 0000000000..21b535dfa8 --- /dev/null +++ b/External/HIP/catch/include/cmd_options.hh @@ -0,0 +1,41 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include +#include + +struct CmdOptions { + int iterations = 5; + int warmups = 5; + int cg_extended_run = 5; + int cg_iterations = 2; + bool no_display = false; + bool progress = false; + uint64_t accuracy_iterations = std::numeric_limits::max() + 1ull; + uint64_t reduce_iterations = 1; + uint64_t reduce_input_size = 50; + int accuracy_max_memory = 80; +}; + +extern CmdOptions cmd_options; diff --git a/External/HIP/catch/include/hip_test_common.hh b/External/HIP/catch/include/hip_test_common.hh new file mode 100644 index 0000000000..989f09e425 --- /dev/null +++ b/External/HIP/catch/include/hip_test_common.hh @@ -0,0 +1,726 @@ +/* +Copyright (c) 2021 - 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once +#pragma clang diagnostic ignored "-Wsign-compare" +#include "hip_test_context.hh" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "hip_test_features.hh" + +#if HT_LINUX +#include +#endif + +#if !defined(__HIP_ATOMIC_BACKWARD_COMPAT) +#define __HIP_ATOMIC_BACKWARD_COMPAT 1 +#endif + +#if HT_AMD +#if defined(__has_extension) && __has_extension(clang_atomic_attributes) && __HIP_ATOMIC_BACKWARD_COMPAT +#define HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY [[clang::atomic(fine_grained_memory, remote_memory)]] +#else +#define HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY +#endif +#elif HT_NVIDIA +#define HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY +#endif + +#ifdef TEST_CLOCK_CYCLE +#define clock_function() clock64() +#else +#define clock_function() wall_clock64() +#endif + +#define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__); + +#define CHAR_BUF_SIZE 512 + +#define CONSOLE_PRINT(fmt, ...) \ + do { \ + std::printf(fmt "\n", ##__VA_ARGS__); \ + } while (0) + +// DEBUG_PRINT: If ENABLE_DEBUG is defined, prints immediately to console. +// Otherwise, uses Catch2 INFO() - debug messages will only appear if the test fails. +#if defined(ENABLE_DEBUG) +#define DEBUG_PRINT(fmt, ...) CONSOLE_PRINT("[DEBUG]: " fmt, ##__VA_ARGS__) +#else +#define DEBUG_PRINT(fmt, ...) \ + do { \ + char buf[CHAR_BUF_SIZE]; \ + std::snprintf(buf, CHAR_BUF_SIZE, "[INFO]: " fmt, ##__VA_ARGS__); \ + INFO(buf); \ + } while (0) +#endif + +// Not thread-safe +#define HIP_CHECK(error) \ + { \ + hipError_t localError = error; \ + if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ + INFO("Error: " << hipGetErrorString(localError) << "\n Code: " << localError \ + << "\n Str: " << #error << "\n In File: " << __FILE__ \ + << "\n At line: " << __LINE__); \ + REQUIRE(false); \ + } \ + } + +#define HIP_CHECK_IGNORED_RETURN(error, ignoredError) \ + { \ + hipError_t localError = error; \ + if ((localError == ignoredError)) { \ + INFO("Skipped: " << hipGetErrorString(localError) << "\n Code: " << localError \ + << "\n Str: " << #error << "\n In File: " << __FILE__ \ + << "\n At line: " << __LINE__); \ + return; \ + } \ + if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ + INFO("Error: " << hipGetErrorString(localError) << "\n Code: " << localError \ + << "\n Str: " << #error << "\n In File: " << __FILE__ \ + << "\n At line: " << __LINE__); \ + REQUIRE(false); \ + } \ + } + +// Threaded HIP_CHECKs +#define HIP_CHECK_THREAD(error) \ + { \ + /*To see if error has occured in previous threads, stop execution */ \ + if (TestContext::get().hasErrorOccured() == true) { \ + return; /*This will only work with std::thread and not with std::async*/ \ + } \ + auto localError = error; \ + HCResult result(__LINE__, __FILE__, localError, #error); \ + TestContext::get().addResults(result); \ + } + +#define REQUIRE_THREAD(condition) \ + { \ + /*To see if error has occured in previous threads, stop execution */ \ + if (TestContext::get().hasErrorOccured() == true) { \ + return; /*This will only work with std::thread and not with std::async*/ \ + } \ + auto localResult = (condition); \ + HCResult result(__LINE__, __FILE__, hipSuccess, #condition, localResult); \ + TestContext::get().addResults(result); \ + } + +// Do not call before all threads have joined +#define HIP_CHECK_THREAD_FINALIZE() \ + { TestContext::get().finalizeResults(); } + + +// Check that an expression, errorExpr, evaluates to the expected error_t, expectedError. +#define HIP_CHECK_ERROR(errorExpr, expectedError) \ + { \ + hipError_t localError = errorExpr; \ + INFO("Matching Errors: " \ + << "\n Expected Error: " << hipGetErrorString(expectedError) \ + << "\n Expected Code: " << expectedError << '\n' \ + << " Actual Error: " << hipGetErrorString(localError) \ + << "\n Actual Code: " << localError << "\nStr: " << #errorExpr \ + << "\n In File: " << __FILE__ << "\n At line: " << __LINE__); \ + REQUIRE(localError == expectedError); \ + } + +// Check that an expression, errorExpr, evaluates to the expected error_t, expectedError or +// expectedError1. +#define HIP_CHECK_ERRORS(errorExpr, expectedError, expectedError1) \ + { \ + hipError_t localError = errorExpr; \ + INFO("Matching Errors: " \ + << "\n Expected Error: " << hipGetErrorString(expectedError) \ + << "\n Expected Code: " << expectedError << " or " << expectedError << '\n' \ + << " Actual Error: " << hipGetErrorString(localError) \ + << "\n Actual Code: " << localError << "\nStr: " << #errorExpr \ + << "\n In File: " << __FILE__ << "\n At line: " << __LINE__); \ + REQUIRE((localError == expectedError || localError == expectedError1)); \ + } + +// Not thread-safe +#define HIPRTC_CHECK(error) \ + { \ + auto localError = error; \ + if (localError != HIPRTC_SUCCESS) { \ + INFO("Error: " << hiprtcGetErrorString(localError) << "\n Code: " << localError \ + << "\n Str: " << #error << "\n In File: " << __FILE__ \ + << "\n At line: " << __LINE__); \ + REQUIRE(false); \ + } \ + } + +// Check that an expression, errorExpr, evaluates to the expected error_t, expectedError. +#define HIPRTC_CHECK_ERROR(errorExpr, expectedError) \ + { \ + auto localError = errorExpr; \ + INFO("Matching Errors: " \ + << "\n Expected Error: " << hiprtcGetErrorString(expectedError) \ + << "\n Expected Code: " << expectedError << '\n' \ + << " Actual Error: " << hiprtcGetErrorString(localError) \ + << "\n Actual Code: " << localError << "\nStr: " << #errorExpr \ + << "\n In File: " << __FILE__ << "\n At line: " << __LINE__); \ + REQUIRE(localError == expectedError); \ + } + +// Although its assert, it will be evaluated at runtime +#define HIP_ASSERT(x) \ + { REQUIRE((x)); } + +#define HIPCHECK(error) \ + { \ + hipError_t localError = error; \ + if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ + printf("error: '%s'(%d) from %s at %s:%d\n", hipGetErrorString(localError), localError, \ + #error, __FILE__, __LINE__); \ + abort(); \ + } \ + } + +// Check that an expression, errorExpr, evaluates to the expected error_t, expectedError. +#define HIPRTC_CHECK_ERROR(errorExpr, expectedError) \ + { \ + auto localError = errorExpr; \ + INFO("Matching Errors: " \ + << "\n Expected Error: " << hiprtcGetErrorString(expectedError) \ + << "\n Expected Code: " << expectedError << '\n' \ + << " Actual Error: " << hiprtcGetErrorString(localError) \ + << "\n Actual Code: " << localError << "\nStr: " << #errorExpr \ + << "\n In File: " << __FILE__ << "\n At line: " << __LINE__); \ + REQUIRE(localError == expectedError); \ + } + +#define HIPASSERT(condition) \ + if (!(condition)) { \ + printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \ + abort(); \ + } + +#if HT_NVIDIA +#define CTX_CREATE() \ + hipCtx_t context; \ + initHipCtx(&context); +#define CTX_DESTROY() HIPCHECK(hipCtxDestroy(context)); +#define ARRAY_DESTROY(array) HIPCHECK(hipArrayDestroy(array)); +#define HIP_TEX_REFERENCE hipTexRef +#define HIP_ARRAY hipArray_t +static void initHipCtx(hipCtx_t* pcontext) { + HIPCHECK(hipInit(0)); + hipDevice_t device; + HIPCHECK(hipDeviceGet(&device, 0)); + HIPCHECK(hipCtxCreate(pcontext, 0, device)); +} +#else +#define CTX_CREATE() +#define CTX_DESTROY() +#define ARRAY_DESTROY(array) HIPCHECK(hipFreeArray(array)); +#define HIP_TEX_REFERENCE textureReference* +#define HIP_ARRAY hipArray_t +#endif + +static inline int getWarpSize() { +#if HT_NVIDIA + return 32; +#elif HT_AMD + int device = -1; + int warpSize = -1; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipDeviceGetAttribute(&warpSize, hipDeviceAttributeWarpSize, device)); + return warpSize; +#else + std::cout<<"Have to be either Nvidia or AMD platform, asserting"< N) { + blocks = (N + threadsPerBlock - 1) / threadsPerBlock; + } +} + +static inline int RAND_R(unsigned* rand_seed) { +#if defined(_WIN32) || defined(_WIN64) + srand(*rand_seed); + return rand(); +#else + return rand_r(rand_seed); +#endif +} + +inline bool isImageSupported() { + int imageSupport = 1; +#if HT_AMD + int device; + HIP_CHECK(hipGetDevice(&device)); + HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, device)); +#endif + return imageSupport != 0; +} + +inline bool isPcieAtomicSupported() { + int pcieAtomic = 1; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, device)); + return pcieAtomic; +} + +inline bool isP2PSupported(int& d1, int& d2) { + int num_devices = HipTest::getDeviceCount(); + int supported = 1; + for (auto i = 0u; i < num_devices; ++i) { + int canAccess = 0; + for (auto j = 0u; j < num_devices; ++j) { + if (i != j) { + HIP_CHECK(hipDeviceCanAccessPeer(&canAccess, i, j)); + if (!canAccess) { + supported = 0; + d1 = i; + d2 = j; + break; + } + } + } + } + return supported; +} + +inline bool checkConcurrentKernels(int num_devices) { + for (auto i = 0; i < num_devices; ++i) { + HIP_CHECK(hipSetDevice(i)); + int concurrent_kernels = 0; + HIP_CHECK(hipDeviceGetAttribute(&concurrent_kernels, hipDeviceAttributeConcurrentKernels, i)); + if (!concurrent_kernels) { + return false; + } + } + if (num_devices > 1) { + HIP_CHECK(hipSetDevice(0)); + } + return true; +} + +inline bool isXnackOn() { + hipDeviceProp_t prop; + int device = 0; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + return gfxName.find("xnack+") != std::string::npos; +} + +inline bool areWarpMatchFunctionsSupported() { + int matchFunctionsSupported = 1; +#if HT_NVIDIA + int device; + hipDeviceProp_t prop; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + if (prop.major < 7) { + matchFunctionsSupported = 0; + } +#endif + return matchFunctionsSupported != 0; +} + +/** + * Causes the test to stop and be skipped at runtime. + * reason: Message describing the reason the test has been skipped. + */ +static inline void HIP_SKIP_TEST(char const* const reason) noexcept { + // ctest is setup to parse for "HIP_SKIP_THIS_TEST", at which point it will skip the test. + std::cout << "Skipping test. Reason: " << reason << '\n' << "HIP_SKIP_THIS_TEST" << std::endl; +} + +/** + * @brief Helper template that returns the expected arguments of a kernel. + * + * @return constexpr std::tuple the expected arguments of the kernel. + */ +template std::tuple getExpectedArgs(void(FArgs...)){}; + +/** + * @brief Asserts that the types of the arguments of a function match exactly with the types in the + * function signature. + * This is necessary because HIP RTC does not do implicit casting of the kernel + * parameters. + * In order to get the kernel function signature, this function should only called when + * RTC is disabled. + * + * @tparam F the kernel function + * @tparam Args the parameters that will be passed to the kernel. + */ +template void validateArguments(F f, Args...) { + using expectedArgsTuple = decltype(getExpectedArgs(f)); + static_assert(std::is_same>::value, + "Kernel arguments types must match exactly!"); +} + +/** + * @brief Launch a kernel using either HIP or HIP RTC. + * + * @tparam Typenames A list of typenames used by the kernel (unused if the kernel is not a + * template). + * @tparam K The kernel type. Expects a function or template when RTC is disabled. Expects a + * function pointer instead when RTC is enabled. + * @tparam Dim Can be either dim3 or int. + * @tparam Args A list of kernel arguments to be forwarded. + * @param kernel The kernel to be launched (defined in kernels.hh) + * @param numBlocks + * @param numThreads + * @param memPerBlock + * @param stream + * @param packedArgs A list of kernel arguments to be forwarded. + */ +template +void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock, + hipStream_t stream, Args&&... packedArgs) { +#ifndef RTC_TESTING + validateArguments(kernel, packedArgs...); + kernel<<>>(std::forward(packedArgs)...); +#else + launchRTCKernel(kernel, numBlocks, numThreads, memPerBlock, stream, + std::forward(packedArgs)...); +#endif + HIP_CHECK(hipGetLastError()); +} + +//--- +struct Pinned { + static const bool isPinned = true; + static const char* str() { return "Pinned"; }; + + static void* Alloc(size_t sizeBytes) { + void* p; + HIPCHECK(hipHostMalloc((void**)&p, sizeBytes)); + return p; + }; +}; + + +//--- +struct Unpinned { + static const bool isPinned = false; + static const char* str() { return "Unpinned"; }; + + static void* Alloc(size_t sizeBytes) { + void* p = malloc(sizeBytes); + HIPASSERT(p); + return p; + }; +}; + + +struct Memcpy { + static const char* str() { return "Memcpy"; }; +}; + +struct MemcpyAsync { + static const char* str() { return "MemcpyAsync"; }; +}; + + +template struct MemTraits; + + +template <> struct MemTraits { + static void Copy(void* dest, const void* src, size_t sizeBytes, hipMemcpyKind kind, + hipStream_t stream) { + (void)stream; + HIPCHECK(hipMemcpy(dest, src, sizeBytes, kind)); + } +}; + + +template <> struct MemTraits { + static void Copy(void* dest, const void* src, size_t sizeBytes, hipMemcpyKind kind, + hipStream_t stream) { + HIPCHECK(hipMemcpyAsync(dest, src, sizeBytes, kind, stream)); + } +}; + +class BlockingContext { + std::atomic_bool blocked{true}; + hipStream_t stream; + + public: + BlockingContext(hipStream_t s) : blocked(true), stream(s) {} + + BlockingContext(const BlockingContext& in) { + blocked = in.blocked_val(); + stream = in.stream_val(); + } + + BlockingContext(const BlockingContext&& in) { + blocked = in.blocked_val(); + stream = in.stream_val(); + } + + void reset() { blocked = true; } + + BlockingContext& operator=(const BlockingContext& in) { + blocked = in.blocked_val(); + stream = in.stream_val(); + return *this; + } + + void block_stream() { + blocked = true; + auto blocking_callback = [](hipStream_t, hipError_t, void* data) { + auto blocked = reinterpret_cast(data); + while (blocked->load()) { + // Yield this thread till we are waiting + std::this_thread::yield(); + } + }; + HIP_CHECK(hipStreamAddCallback(stream, blocking_callback, (void*)&blocked, 0)); + } + + void unblock_stream() { + blocked = false; + } + + bool is_blocked() const { return hipStreamQuery(stream) == hipErrorNotReady; } + + bool blocked_val() const { return blocked.load(); } + hipStream_t stream_val() const { return stream; } +}; +} // namespace HipTest + +// This must be called in the beginning of image test app's main() to indicate whether image +// is supported. +#define CHECK_IMAGE_SUPPORT \ + if (!HipTest::isImageSupported()) { \ + INFO("Texture is not support on the device. Skipped."); \ + return; \ + } + +#define CHECK_PCIE_ATOMIC_SUPPORT \ + if (!HipTest::isPcieAtomicSupported()) { \ + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); \ + return; \ + } + +#define CHECK_P2P_SUPPORT \ + int d1, d2; \ + if (!HipTest::isP2PSupported(d1,d2)) { \ + std::string msg = "P2P access check failed between dev1:" + std::to_string(d1) + ",dev2:" + \ + std::to_string(d2); \ + HipTest::HIP_SKIP_TEST(msg.c_str()); \ + return; \ + } \ +// This must be called in the beginning of warp test app's main() to indicate warp match functions +// are supported. +#define CHECK_WARP_MATCH_FUNCTIONS_SUPPORT \ + if (!HipTest::areWarpMatchFunctionsSupported()) { \ + INFO("Warp Match Functions are not support on the device. Skipped."); \ + return; \ + } + +// Call GENERATE_CAPTURE macro at the start of the test, before using BEGIN/END_CAPTURE. +// Use BEGIN/END_CAPTURE macros to execute APIs in both stream capturing and non-capturing modes. +// Place BEGIN_CAPTURE before the API call and END_CAPTURE after the call. +#define GENERATE_CAPTURE() bool capture = GENERATE(true, false); + +#define BEGIN_CAPTURE(stream) \ + if (capture && stream != nullptr) { \ + hipStreamCaptureMode flags = GENERATE( \ + hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); \ + HIP_CHECK(hipStreamBeginCapture(stream, flags)); \ + } + +#define END_CAPTURE(stream) \ + if (capture && stream != nullptr) { \ + hipGraph_t graph = nullptr; \ + hipGraphExec_t graph_exec = nullptr; \ + HIP_CHECK(hipStreamEndCapture(stream, &graph)); \ + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); \ + HIP_CHECK(hipGraphLaunch(graph_exec, stream)); \ + HIP_CHECK(hipGraphExecDestroy(graph_exec)); \ + HIP_CHECK(hipGraphDestroy(graph)); \ + } + +// These macros are used for testing behaviour when sync APIs are being captured. Before +// calling BEGIN_CAPTURE_SYNC, hipError_t variable (capture_err) should be initialized to hipSuccess +// and passed to this macro. The scenario with using this macro should look like this: +// 1. BEGIN_CAPTURE_SYNC(capture_err) +// 2. HIP_CHECK_ERROR(SyncAPI, capture_err) +// 3. END_CAPTURE_SYNC(capture_err) +// Some sync APIs are allowed in relaxed capture mode which is indicated with +// rlx_mode_allowed variable. For other two modes, those APIs return +// hipErrorStreamCaptureUnsupported. These macros shouldn't be used with hipStreamSync and +// hipDeviceSync during capture. +#define BEGIN_CAPTURE_SYNC(capture_err, rlx_mode_allowed) \ + hipStream_t stream; \ + GENERATE_CAPTURE(); \ + if (capture) { \ + HIP_CHECK(hipStreamCreate(&stream)); \ + hipStreamCaptureMode mode = GENERATE( \ + hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); \ + HIP_CHECK(hipStreamBeginCapture(stream, mode)); \ + if (!rlx_mode_allowed) { \ + capture_err = hipErrorStreamCaptureImplicit; \ + } else if (mode != hipStreamCaptureModeRelaxed) { \ + capture_err = hipErrorStreamCaptureUnsupported; \ + } \ + } + +// If test has other HIP API calls that depend on sync call that is captured and fails, the rest of +// the test (except freeing the memory) should be skipped after calling END_CAPTURE_SYNC() by +// testing if previously created hipError_t variable (capture_err) doesn't equal hipSuccess. +#define END_CAPTURE_SYNC(capture_err) \ + if (capture) { \ + hipGraph_t graph; \ + hipError_t stream_err = hipSuccess; \ + if (capture_err != hipSuccess) { \ + stream_err = hipErrorStreamCaptureInvalidated; \ + } \ + HIP_CHECK_ERROR(hipStreamEndCapture(stream, &graph), stream_err); \ + if (graph != nullptr) { \ + HIP_CHECK(hipGraphDestroy(graph)); \ + } \ + HIP_CHECK(hipStreamDestroy(stream)); \ + } + +// Manage core dumps in specific tests which require it disabled (e.g., hipGetLastErrorOnAbort.cc) +#if HT_LINUX +#define DISABLE_CORE_DUMPS() \ + struct rlimit core_limit; \ + bool rlimit_saved = false; \ + if (getrlimit(RLIMIT_CORE, &core_limit) == 0) { \ + if (core_limit.rlim_cur != 0) { \ + struct rlimit new_limit; \ + new_limit.rlim_cur = 0; \ + new_limit.rlim_max = core_limit.rlim_max; \ + if (setrlimit(RLIMIT_CORE, &new_limit) == 0) { \ + rlimit_saved = true; \ + } \ + } \ + } + +#define RESTORE_CORE_DUMPS() \ + if (rlimit_saved) { \ + setrlimit(RLIMIT_CORE, &core_limit); \ + rlimit_saved = false; \ + } +#else +#define DISABLE_CORE_DUMPS() +#define RESTORE_CORE_DUMPS() +#endif diff --git a/External/HIP/catch/include/hip_test_context.hh b/External/HIP/catch/include/hip_test_context.hh new file mode 100644 index 0000000000..8e06c3fbb8 --- /dev/null +++ b/External/HIP/catch/include/hip_test_context.hh @@ -0,0 +1,197 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +// OS Check +#if defined(_WIN32) +#define HT_WIN 1 +#define HT_LINUX 0 +#elif defined(__linux__) +#define HT_WIN 0 +#define HT_LINUX 1 +#else +#error "OS not recognized" +#endif + +// Platform check +#if defined(__HIP_PLATFORM_AMD__) +#define HT_AMD 1 +#define HT_NVIDIA 0 +#elif defined(__HIP_PLATFORM_NVIDIA__) +#define HT_AMD 0 +#define HT_NVIDIA 1 +#else +#error "Platform not recognized" +#endif + +typedef struct Config_ { + std::vector json_files; // Json files + std::string platform; // amd/nvidia + std::string os; // windows/linux +} Config; + +// Store Multi threaded results +struct HCResult { + size_t line; // Line of check (HIP_CHECK_THREAD or REQUIRE_THREAD) + std::string file; // File name of the check + hipError_t result; // hipResult for HIP_CHECK_THREAD, for conditions its hipSuccess + std::string call; // Call of HIP API or a bool condition + bool conditionsResult; // If bool condition, result of call. For HIP Calls its true + HCResult(size_t l, std::string f, hipError_t r, std::string c, bool b = true) + : line(l), file(f), result(r), call(c), conditionsResult(b) {} +}; + + +class TestContext { + bool p_windows = false, p_linux = false; // OS + bool amd = false, nvidia = false; // HIP Platform + std::string exe_path; + std::string current_test; + std::set skip_test; + std::string json_file_; + std::vector platform_list_ = {"amd", "nvidia"}; + std::vector os_list_ = {"windows", "linux", "all"}; + std::vector amd_arch_list_ = {}; + + struct rtcState { + hipModule_t module; + hipFunction_t kernelFunction; + }; + + std::unordered_map compiledKernels{}; + + Config config_; + std::string& getCommonJsonFile(); + std::string substringFound(std::vector list, std::string filename); + void detectOS(); + void detectPlatform(); + void getConfigFiles(); + void setExePath(int, char**); + void parseOptions(int, char**); + bool parseJsonFiles(); + std::string getMatchingConfigFile(std::string config_dir); + std::string getCurrentArch(); + const Config& getConfig() const { return config_; } + + + TestContext(int argc, char** argv); + + // Multi threaded checks helpers + std::mutex resultMutex; + std::vector results; // Multi threaded test results buffer + std::atomic hasErrorOccured_{false}; + + public: + static TestContext& get(int argc = 0, char** argv = nullptr) { + static TestContext instance(argc, argv); + return instance; + } + + static std::string getEnvVar(std::string var) { + #if defined(_WIN32) + constexpr rsize_t MAX_LEN = 4096; + char dstBuf[MAX_LEN]; + size_t dstSize; + if (!::getenv_s(&dstSize, dstBuf, MAX_LEN, var.c_str())) { + return std::string(dstBuf); + } + #elif defined(__linux__) + char* val = std::getenv(var.c_str()); + if (val != NULL) { + return std::string(val); + } + #else + #error "OS not recognized" + #endif + return std::string(""); + } + + + bool isWindows() const; + bool isLinux() const; + bool isNvidia() const; + bool isAmd() const; + bool skipTest() const; + + const std::string& getCurrentTest() const { return current_test; } + std::string currentPath() const; + + // Multi threaded results helpers + void addResults(HCResult r); // Add multi threaded results + void finalizeResults(); // Validate on all results + bool hasErrorOccured(); // Query if error has occured + + /** + * @brief Unload all loaded modules. + * Note: This function needs to be called at the end of each test that uses RTC. + * It is not possible to unload the loaded modules without adding explicit code to the end + * of each test. This function exists only to provide a clean way to exit a test when using RTC. + * However, not unloading a module explicitly shouldn't have any effect on the outcome of + * the test. + */ + void cleanContext(); + + /** + * @brief Keeps track of all the already compiled rtc kernels. + * + * @param kernelNameExpression The name expression (e.g. hipTest::vectorADD). + * @param loadedModule The loaded module. + * @param kernelFunction The hipFunction that will be used to run the kernel in the future. + */ + void trackRtcState(std::string kernelNameExpression, hipModule_t loadedModule, + hipFunction_t kernelFunction); + + /** + * @brief Get the already compiled hip rtc kernel function if it exists. + * + * @param kernelNameExpression The name expression (e.g. hipTest::vectorADD). + * @return the hipFunction if it exists. nullptr otherwise + */ + hipFunction_t getFunction(const std::string kernelNameExpression); + + TestContext(const TestContext&) = delete; + void operator=(const TestContext&) = delete; + + ~TestContext(); +}; + +static bool _log_enable = (!TestContext::getEnvVar("HT_LOG_ENABLE").empty() ? true : false); + +// printing logs +#define LogPrintf(format, ...) \ +{ \ + if(_log_enable) { \ + printf(format, __VA_ARGS__); \ + printf("%c", '\n'); \ + } \ +} diff --git a/External/HIP/catch/include/hip_test_features.hh b/External/HIP/catch/include/hip_test_features.hh new file mode 100644 index 0000000000..0534e9b954 --- /dev/null +++ b/External/HIP/catch/include/hip_test_features.hh @@ -0,0 +1,40 @@ +/* +Copyright (c) 2021 - 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include +#include +#include +#include + +// Catch Test Features +typedef enum CTFeatures { + CT_FEATURE_FINEGRAIN_HWSUPPORT = 0x0, // FINEGRAIN Supported Hardware. + CT_FEATURE_HMM = 0x1, // HMM Enabled + CT_FEATURE_TEXTURES_NOT_SUPPORTED = 0x2, // Textures not supported + CT_FEATURE_LAST = 0x3 +} CTFeatures; + +bool CheckIfFeatSupported(enum CTFeatures test_feat, std::string gcn_arch); +bool getGenericTarget(const std::string& agentTarget, std::string& genericTarget); +bool isGenericTargetSupported(char* gcnArchName = nullptr, int deviceId = 0); diff --git a/External/HIP/catch/include/hip_test_filesystem.hh b/External/HIP/catch/include/hip_test_filesystem.hh new file mode 100644 index 0000000000..fffda4146b --- /dev/null +++ b/External/HIP/catch/include/hip_test_filesystem.hh @@ -0,0 +1,89 @@ + +/* +Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once +// We haven't checked which filesystem to include yet +#ifndef INCLUDE_STD_FILESYSTEM_EXPERIMENTAL +// Check for feature test macro for +#if defined(__cpp_lib_filesystem) +#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 0 +// Check for feature test macro for +#elif defined(__cpp_lib_experimental_filesystem) +#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 1 +// We can't check if headers exist... +// Let's assume experimental to be safe +#elif !defined(__has_include) +#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 1 +// Check if the header "" exists +#elif __has_include() +// If we're compiling on Visual Studio and are not compiling with C++17, +// we need to use experimental +#ifdef _MSC_VER +// Check and include header that defines "_HAS_CXX17" +#if __has_include() +#include + +// Check for enabled C++17 support +#if defined(_HAS_CXX17) && _HAS_CXX17 +// We're using C++17, so let's use the normal version +#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 0 +#endif + +#endif + +// If the marco isn't defined yet, that means any of the other +// VS specific checks failed, so we need to use experimental +#ifndef INCLUDE_STD_FILESYSTEM_EXPERIMENTAL +#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 1 +#endif + +// Not on Visual Studio. Let's use the normal version +#else // #ifdef _MSC_VER +#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 0 +#endif + +// Check if the header "" exists +#elif __has_include() +#define INCLUDE_STD_FILESYSTEM_EXPERIMENTAL 1 + +// Fail if neither header is available with a nice error message +#else +#error Could not find system header "" || + "" +#endif + +// We priously determined that we need the exprimental version +#if INCLUDE_STD_FILESYSTEM_EXPERIMENTAL +// Include it +#define _SILENCE_EXPERIMENTAL_FILESYSTEM_DEPRECATION_WARNING 1; +#include +// We need the alias from std::experimental::filesystem to std::filesystem +namespace fs = std::experimental::filesystem; +// We have a decent compiler and can use the normal version +#else +// Include it +#include +namespace fs = std::filesystem; +#endif + +#endif // #ifndef INCLUDE_STD_FILESYSTEM_EXPERIMENTAL \ No newline at end of file diff --git a/External/HIP/catch/unit/compiler/hipClassKernel.cc b/External/HIP/catch/unit/compiler/hipClassKernel.cc new file mode 100644 index 0000000000..39cee07da7 --- /dev/null +++ b/External/HIP/catch/unit/compiler/hipClassKernel.cc @@ -0,0 +1,172 @@ +/* +Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hipClassKernel.h" + +__global__ void ovrdClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testOvrD tobj1; + result_ecd[tid] = (tobj1.ovrdFunc1() == 30); +} + +__global__ void ovldClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testFuncOvld tfo1; + result_ecd[tid] = (tfo1.func1(10) == 20) && (tfo1.func1(10, 10) == 30); +} + +TEST_CASE("Unit_hipClassKernel_Overload_Override") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(ovrdClassKernel, dim3(BLOCKS), dim3(THREADS_PER_BLOCK), 0, 0, result_ecd); + + VerifyResult(result_ech, result_ecd); + FreeMem(result_ech, result_ecd); + + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(ovldClassKernel, dim3(BLOCKS), dim3(THREADS_PER_BLOCK), 0, 0, result_ecd); + + VerifyResult(result_ech, result_ecd); + FreeMem(result_ech, result_ecd); +} + +// check for friend +__global__ void friendClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testFrndB tfb1; + result_ecd[tid] = (tfb1.showA() == 10); +} + +TEST_CASE("Unit_hipClassKernel_Friend") { + bool* result_ecd; + result_ecd = AllocateDeviceMemory(); + hipLaunchKernelGGL(friendClassKernel, dim3(BLOCKS), dim3(THREADS_PER_BLOCK), 0, 0, result_ecd); + HIP_CHECK(hipStreamSynchronize(nullptr)); + HIP_CHECK(hipFree(result_ecd)); +} + +// check sizeof empty class is 1 +__global__ void emptyClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testClassEmpty ob1, ob2; + result_ecd[tid] = (sizeof(testClassEmpty) == 1) && (&ob1 != &ob2); +} + +TEST_CASE("Unit_hipClassKernel_Empty") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(emptyClassKernel, dim3(BLOCKS), dim3(THREADS_PER_BLOCK), 0, 0, result_ecd); + + VerifyResult(result_ech, result_ecd); + FreeMem(result_ech, result_ecd); +} + +// tests for classes >8 bytes +__global__ void sizeClassBKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeB) == 12) && (sizeof(testSizeC) == 16) && + (sizeof(testSizeP1) == 6) && (sizeof(testSizeP2) == 13) && + (sizeof(testSizeP3) == 8); +} + +TEST_CASE("Unit_hipClassKernel_BSize") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeClassBKernel, dim3(BLOCKS), dim3(THREADS_PER_BLOCK), 0, 0, result_ecd); + + VerifyResult(result_ech, result_ecd); + FreeMem(result_ech, result_ecd); +} + +__global__ void sizeClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeA) == 16) && (sizeof(testSizeDerived) == 24) && + (sizeof(testSizeDerived2) == 20); +} + +TEST_CASE("Unit_hipClassKernel_Size") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeClassKernel, dim3(BLOCKS), dim3(THREADS_PER_BLOCK), 0, 0, result_ecd); + + VerifyResult(result_ech, result_ecd); + FreeMem(result_ech, result_ecd); +} + +__global__ void sizeVirtualClassKernel(bool* result_ecd, refStructSizes structSizes) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (structSizes.sizeOftestSizeDV == sizeof(testSizeDV)) && + (structSizes.sizeOftestSizeDerivedDV == sizeof(testSizeDerivedDV)) && + (structSizes.sizeOftestSizeVirtDer = sizeof(testSizeVirtDer)) && + (structSizes.sizeOftestSizeVirtDerPack = sizeof(testSizeVirtDerPack)) && + (structSizes.sizeOftestSizeDerMulti = sizeof(testSizeDerMulti)); +} + +TEST_CASE("Unit_hipClassKernel_Virtual") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + struct refStructSizes structSizes; + structSizes.sizeOftestSizeDV = sizeof(testSizeDV); + structSizes.sizeOftestSizeDerivedDV = sizeof(testSizeDerivedDV); + structSizes.sizeOftestSizeVirtDer = sizeof(testSizeVirtDer); + structSizes.sizeOftestSizeVirtDerPack = sizeof(testSizeVirtDerPack); + structSizes.sizeOftestSizeDerMulti = sizeof(testSizeDerMulti); + + hipLaunchKernelGGL(sizeVirtualClassKernel, dim3(BLOCKS), dim3(THREADS_PER_BLOCK), 0, 0, + result_ecd, structSizes); + + VerifyResult(result_ech, result_ecd); + FreeMem(result_ech, result_ecd); +} + +// check pass by value +__global__ void passByValueKernel(testPassByValue obj, bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (obj.exI == 10) && (obj.exC == 'C'); +} + +TEST_CASE("Unit_hipClassKernel_Value") { + bool *result_ecd, *result_ech; + result_ech = AllocateHostMemory(); + result_ecd = AllocateDeviceMemory(); + + testPassByValue exObj; + exObj.exI = 10; + exObj.exC = 'C'; + hipLaunchKernelGGL(passByValueKernel, dim3(BLOCKS), dim3(THREADS_PER_BLOCK), 0, 0, exObj, + result_ecd); + + VerifyResult(result_ech, result_ecd); + FreeMem(result_ech, result_ecd); +} diff --git a/External/HIP/catch/unit/compiler/hipClassKernel.h b/External/HIP/catch/unit/compiler/hipClassKernel.h new file mode 100644 index 0000000000..674f0c381c --- /dev/null +++ b/External/HIP/catch/unit/compiler/hipClassKernel.h @@ -0,0 +1,220 @@ +/* +Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#ifndef _COMPILER_HIPCLASSKERNEL_H_ +#define _COMPILER_HIPCLASSKERNEL_H_ + +#include + +static const int BLOCKS = 512; +static const int THREADS_PER_BLOCK = 1; +size_t NBOOL = BLOCKS * sizeof(bool); + +class testFuncOvld { + public: + int __host__ __device__ func1(int a) { return a + 10; } + + int __host__ __device__ func1(int a, int b) { return a + b + 10; } +}; + +class testOvrB { + public: + int __host__ __device__ ovrdFunc1() { return 10; } +}; + +class testOvrD : public testOvrB { + public: + int __host__ __device__ ovrdFunc1() { + int x = testOvrB::ovrdFunc1(); + return x + 20; + } +}; + +class testFrndA { + private: + int fa = 10; + + public: + friend class testFrndB; +}; + +class testFrndB { + public: + __host__ __device__ int showA() { + testFrndA x; + return x.fa; + } +}; + +class testClassEmpty {}; + +class testPassByValue { + public: + int exI; + char exC; +}; + +class testSizeA { + public: + float xa; + int ia; + double da; + static char ca; +}; + +class testSizeDerived : testSizeA { + public: + float fd; +}; + +#pragma pack(push, 4) +class testSizeDerived2 : testSizeA { + public: + float fd; +}; +#pragma pack(pop) + +class testSizeB { + public: + char ab; + int ib; + char cb; +}; + +class testSizeVirtDer : public virtual testSizeB { + public: + int ivd; +}; + +class testSizeVirtDer1 : public virtual testSizeB { + public: + int ivd1; +}; + +class testSizeDerMulti : public testSizeVirtDer, public testSizeVirtDer1 { + public: + int ivd2; +}; + +#pragma pack(push, 4) +class testSizeVirtDerPack : public virtual testSizeB { + public: + int ivd; +}; +#pragma pack(pop) + +class testSizeC { + public: + char ac; + int ic; + int bc[2]; +}; + +class testSizeDV { + public: + virtual void __host__ __device__ func1(); + + private: + int iDV; +}; + +class testSizeDerivedDV : testSizeDV { + public: + virtual void __host__ __device__ funcD1(); + + private: + int iDDV; +}; + +struct refStructSizes { + size_t sizeOftestSizeDV; + size_t sizeOftestSizeDerivedDV; + size_t sizeOftestSizeVirtDer; + size_t sizeOftestSizeVirtDerPack; + size_t sizeOftestSizeDerMulti; +}; + +#pragma pack(push, 1) +class testSizeP1 { + public: + char ap; + int ip; + char cp; +}; +#pragma pack(pop) + +#pragma pack(push, 1) +class testSizeP2 { + public: + char ap1; + int ip1; + int bp1[2]; +}; +#pragma pack(pop) + +#pragma pack(push, 2) +class testSizeP3 { + public: + char ap2; + int ip2; + char cp2; +}; +#pragma pack(pop) + +class testDeviceClass { + public: + int iVar; + __host__ __device__ testDeviceClass(); + __host__ __device__ testDeviceClass(int a); + __host__ __device__ ~testDeviceClass(); +}; + +__host__ __device__ testDeviceClass::testDeviceClass() { iVar = 5; } + +__host__ __device__ testDeviceClass::testDeviceClass(int a) { iVar = a; } + +bool* AllocateHostMemory(void) { + bool* result_ech; + HIPCHECK(hipHostMalloc(&result_ech, NBOOL, hipHostMallocDefault)); + return result_ech; +} + +bool* AllocateDeviceMemory(void) { + bool* result_ecd; + HIPCHECK(hipMalloc(&result_ecd, NBOOL)); + HIPCHECK(hipMemset(result_ecd, false, NBOOL)); + return result_ecd; +} + +void VerifyResult(bool* result_ech, bool* result_ecd) { + HIPCHECK(hipMemcpy(result_ech, result_ecd, BLOCKS * sizeof(bool), hipMemcpyDeviceToHost)); + // validation on host side + for (int i = 0; i < BLOCKS; i++) { + HIPASSERT(result_ech[i] == true); + } +} + +void FreeMem(bool* result_ech, bool* result_ecd) { + HIPCHECK(hipHostFree(result_ech)); + HIPCHECK(hipFree(result_ecd)); +} + +#endif // _HIPCLASSKERNEL_H_ diff --git a/External/HIP/catch/unit/compiler/hipSquare.cc b/External/HIP/catch/unit/compiler/hipSquare.cc new file mode 100644 index 0000000000..dd7dc3dfe5 --- /dev/null +++ b/External/HIP/catch/unit/compiler/hipSquare.cc @@ -0,0 +1,86 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include + +/* + * Square each element in the array A and write to array C. + */ +template __global__ void vector_square(T* C_d, const T* A_d, size_t N) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < N; i += stride) { + C_d[i] = A_d[i] * A_d[i]; + } +} + +TEST_CASE("Unit_test_compressed_codeobject") { + float *A_d, *C_d; + float *A_h, *C_h; + size_t N = 1000000; + size_t Nbytes = N * sizeof(float); + static int device = 0; + HIP_CHECK(hipSetDevice(device)); + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, device /*deviceID*/)); + printf("info: running on device %s\n", props.name); +#ifdef __HIP_PLATFORM_AMD__ + printf("info: architecture on AMD GPU device is: %s\n", props.gcnArchName); +#endif + printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); + A_h = (float*)malloc(Nbytes); + HIP_CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess); + C_h = (float*)malloc(Nbytes); + HIP_CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess); + // Fill with Phi + i + for (size_t i = 0; i < N; i++) { + A_h[i] = 1.618f + i; + } + printf("info: allocate device mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + + printf("info: copy Host2Device\n"); + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + + const unsigned blocks = 512; + const unsigned threadsPerBlock = 256; + + printf("info: launch 'vector_square' kernel\n"); + hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); + + printf("info: copy Device2Host\n"); + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + printf("info: check result\n"); + for (size_t i = 0; i < N; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + HIP_CHECK(hipErrorUnknown); + } + } + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); + free(A_h); + free(C_h); + printf("PASSED!\n"); + REQUIRE(true); +} diff --git a/External/HIP/catch/unit/compiler/hipSquareGenericTarget.cc b/External/HIP/catch/unit/compiler/hipSquareGenericTarget.cc new file mode 100644 index 0000000000..5902cbd759 --- /dev/null +++ b/External/HIP/catch/unit/compiler/hipSquareGenericTarget.cc @@ -0,0 +1,120 @@ +/* +Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include + +/* + * Square each element in the array A and write to array C. + */ +template static __global__ void vector_square_generic(T* C_d, const T* A_d, size_t N) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < N; i += stride) { + C_d[i] = A_d[i] * A_d[i]; + } +} + +#ifdef GENERIC_COMPRESSED +TEST_CASE("Unit_test_generic_target_in_compressed_fatbin") { +#else +TEST_CASE("Unit_test_generic_target_in_regular_fatbin") { +#endif + if (!isGenericTargetSupported()) { + fprintf(stderr, "Generic target test is skipped\n"); + return; + } + float *A_d, *C_d; + float *A_h, *C_h; + size_t N = 1000000; + size_t Nbytes = N * sizeof(float); + static int device = 0; + HIP_CHECK(hipSetDevice(device)); + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, device /*deviceID*/)); + printf("info: running on device %s\n", props.name); +#ifdef __HIP_PLATFORM_AMD__ + printf("info: architecture on AMD GPU device is: %s\n", props.gcnArchName); +// check the scope of supportted types +#endif + printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); + A_h = (float*)malloc(Nbytes); + HIP_CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess); + C_h = (float*)malloc(Nbytes); + HIP_CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess); + // Fill with Phi + i + for (size_t i = 0; i < N; i++) { + A_h[i] = 1.618f + i; + } + printf("info: allocate device mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + + printf("info: copy Host2Device\n"); + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + + const unsigned blocks = 512; + const unsigned threadsPerBlock = 256; + + printf("info: launch 'vector_square' kernel\n"); + hipLaunchKernelGGL(vector_square_generic, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); + + printf("info: copy Device2Host\n"); + HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + + printf("info: check result\n"); + for (size_t i = 0; i < N; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + HIP_CHECK(hipErrorUnknown); + } + } + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); + free(A_h); + free(C_h); + printf("PASSED!\n"); + REQUIRE(true); +} + +#ifndef NO_GENERIC_TARGET_ONLY_TEST +#ifdef GENERIC_COMPRESSED +TEST_CASE("Unit_test_generic_target_only_in_compressed_fatbin") { +#ifdef __linux__ + char* cmd = + "chmod u+x ./hipSquareGenericTargetOnlyCompressed && ./hipSquareGenericTargetOnlyCompressed"; +#else + char* cmd = "hipSquareGenericTargetOnlyCompressed.exe"; +#endif +#else // else GENERIC_COMPRESSED +TEST_CASE("Unit_test_generic_target_only_in_regular_fatbin ") { +#ifdef __linux__ + char* cmd = "chmod u+x ./hipSquareGenericTargetOnly && ./hipSquareGenericTargetOnly"; +#else + char* cmd = "hipSquareGenericTargetOnly.exe"; +#endif +#endif // GENERIC_COMPRESSED + + printf("Run %s\n", cmd); + REQUIRE(std::system(cmd) == 0); + printf("PASSED!\n"); +} +#endif // NO_GENERIC_TARGET_ONLY_TEST