From 5589832d8ef3a993974754be0ea89184d2f458c7 Mon Sep 17 00:00:00 2001 From: Tyler Zhao Date: Tue, 6 Jun 2023 11:23:54 -0700 Subject: [PATCH] Fuzz Testing Implementation --- .github/workflows/fuzz-test.yml | 71 + CMakeLists.txt | 7 + fuzz_testing/CMakeLists.txt | 6 + fuzz_testing/README.md | 42 + fuzz_testing/fuzz_src/fuzz_testing.h | 195 +++ .../acl_auto_configure_fuzz_test.yml | 33 + fuzz_testing/script/fuzz_test.py | 269 +++ fuzz_testing/script/mutator.py | 59 + fuzz_testing/test/CMakeLists.txt | 41 + .../test/acl_auto_configure_fuzz_test.cpp | 1487 +++++++++++++++++ fuzz_testing/test/acl_fuzz_test.cpp | 745 +++++++++ fuzz_testing/test/acl_fuzz_test.h | 114 ++ fuzz_testing/test/acl_globals_fuzz_test.cpp | 740 ++++++++ fuzz_testing/test/acl_globals_fuzz_test.h | 12 + fuzz_testing/test/acl_hal_fuzz_test.cpp | 694 ++++++++ fuzz_testing/test/acl_hal_fuzz_test.h | 27 + 16 files changed, 4542 insertions(+) create mode 100644 .github/workflows/fuzz-test.yml create mode 100755 fuzz_testing/CMakeLists.txt create mode 100755 fuzz_testing/README.md create mode 100644 fuzz_testing/fuzz_src/fuzz_testing.h create mode 100644 fuzz_testing/original_inputs/acl_auto_configure_fuzz_test.yml create mode 100644 fuzz_testing/script/fuzz_test.py create mode 100644 fuzz_testing/script/mutator.py create mode 100755 fuzz_testing/test/CMakeLists.txt create mode 100644 fuzz_testing/test/acl_auto_configure_fuzz_test.cpp create mode 100644 fuzz_testing/test/acl_fuzz_test.cpp create mode 100755 fuzz_testing/test/acl_fuzz_test.h create mode 100644 fuzz_testing/test/acl_globals_fuzz_test.cpp create mode 100755 fuzz_testing/test/acl_globals_fuzz_test.h create mode 100644 fuzz_testing/test/acl_hal_fuzz_test.cpp create mode 100755 fuzz_testing/test/acl_hal_fuzz_test.h diff --git a/.github/workflows/fuzz-test.yml b/.github/workflows/fuzz-test.yml new file mode 100644 index 00000000..4c4bc555 --- /dev/null +++ b/.github/workflows/fuzz-test.yml @@ -0,0 +1,71 @@ +name: Runtime Fuzz Testing + +on: + workflow_dispatch: + inputs: + num_of_iterations: + description: Number of iterations per fuzzable variable + required: False + default: 5 + +jobs: + build: + runs-on: + - self-hosted + - linux + - x64 + - container + + container: + image: ghcr.io/intel/fpga-runtime-for-opencl/ubuntu-22.04-dev:main + + name: Fuzz Testing + steps: + - name: Set up Python + uses: actions/setup-python@v4 + with: + python-version: '3.10' + - name: Clone Radamsa + run: | + mkdir radamsa_repo + cd radamsa_repo + git clone https://gitlab.com/akihe/radamsa.git . + - name: Install Radamsa + run: | + cd radamsa_repo + make + sudo make install + cd .. + - name: Install PyYAML + run: pip install pyyaml + - name: Checkout runtime + uses: actions/checkout@v3 + - name: Build + run: | + mkdir -p build/fuzz_testing + cd build/fuzz_testing + CC="${CC:-gcc}" CXX="${CXX:-g++}" cmake -G Ninja ../.. -DCMAKE_BUILD_TYPE=Debug -DACL_CODE_COVERAGE=ON -DACL_TSAN=OFF -DACL_WITH_ASAN=ON -DFUZZ_TESTING=ON "$@" + ninja -v + - name: Fuzz testing + run: | + cd build/fuzz_testing + ls + cd fuzz_testing/script + export AOCL_BOARD_PACKAGE_ROOT="$(git rev-parse --show-toplevel)/test/board/a10_ref" + NUM_OF_ITERATIONS=${{ github.event.inputs.num_of_iterations }} + # This if block is only used during testing, because if this workflow is triggered via pull_request, ${{ github.event.inputs.num_of_iterations }} would be empty + if [ -z "${NUM_OF_ITERATIONS}" ]; then + NUM_OF_ITERATIONS=1 + fi + python3 fuzz_test.py --all -n $NUM_OF_ITERATIONS + - name: Peek results + run: | + cat build/fuzz_testing/fuzz_testing/results/results.yml + - name: Upload results + uses: actions/upload-artifact@v3 + with: + name: fpga-runtime-for-opencl-${{ github.sha }}-fuzz-test-results-${{ github.run_id }} + path: | + /__w/fpga-runtime-for-opencl/fpga-runtime-for-opencl/build/fuzz_testing/fuzz_testing/results/results.yml + /__w/fpga-runtime-for-opencl/fpga-runtime-for-opencl/build/fuzz_testing/fuzz_testing/test_outputs + if-no-files-found: error diff --git a/CMakeLists.txt b/CMakeLists.txt index 28967256..29f10434 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -404,3 +404,10 @@ install(FILES add_subdirectory(lib) add_subdirectory(test) + +# Flag for building fuzz tests +option(FUZZ_TESTING "Build fuzz tests" OFF) +message(STATUS "Build fuzz tests: ${FUZZ_TESTING}") +if(FUZZ_TESTING) + add_subdirectory(fuzz_testing) +endif() diff --git a/fuzz_testing/CMakeLists.txt b/fuzz_testing/CMakeLists.txt new file mode 100755 index 00000000..b536972f --- /dev/null +++ b/fuzz_testing/CMakeLists.txt @@ -0,0 +1,6 @@ +# Copyright (C) 2021 Intel Corporation +# SPDX-License-Identifier: BSD-3-Clause + +add_subdirectory(test) +file(COPY original_inputs DESTINATION ${CMAKE_CURRENT_BINARY_DIR}) +file(COPY script DESTINATION ${CMAKE_CURRENT_BINARY_DIR}) diff --git a/fuzz_testing/README.md b/fuzz_testing/README.md new file mode 100755 index 00000000..3930c19d --- /dev/null +++ b/fuzz_testing/README.md @@ -0,0 +1,42 @@ +# Fuzz Testing + +## Context +"In programming and software development, fuzzing or fuzz testing is an automated software testing technique that involves providing invalid, unexpected, or random data as inputs to a computer program. The program is then monitored for exceptions such as crashes, failing built-in code assertions, or potential memory leaks." + +## How to do fuzz testing on Github +1. Click on the `Actions` tab +2. Click on `Runtime Fuzz Testing` workflow +3. Click on `Run workflow` (Dropdown button) +4. Enter the number of iterations you want a single variable to be mutated +5. Click on `Run workflow` (Green button) +6. Wait until the test finishes +7. Inside the workflow run, click `Peek results` to look at the results +8. Download the artifact from the `Job Summary` to view the full output & result + +## How to do fuzz testing locally +1. Follow [Prerequisites](https://github.com/intel/fpga-runtime-for-opencl#prerequisites) & [Build](https://github.com/intel/fpga-runtime-for-opencl#building-the-runtime) & [Test](https://github.com/intel/fpga-runtime-for-opencl#building-the-runtime) instructions to get a working build. Make sure that the unit tests finishes successfully. +2. Install Radamsa 0.6+, Python 3.10.7+ and Pyyaml/6.0+ +3. Run the following at the top level in the runtime repo +``` +mkdir -p build/fuzz_testing +cd build/fuzz_testing +CC="${CC:-gcc}" CXX="${CXX:-g++}" cmake -G Ninja ../.. -DCMAKE_BUILD_TYPE=Debug -DACL_CODE_COVERAGE=ON -DACL_TSAN=OFF -DACL_WITH_ASAN=ON -DFUZZ_TESTING=ON "$@" +ninja -v +cd fuzz_testing/script +export AOCL_BOARD_PACKAGE_ROOT="$(git rev-parse --show-toplevel)/test/board/a10_ref" +python3 fuzz_test.py --all +``` +4. A results directory and a test_output directory will be created after the fuzz test finishes + +## Classifying Failures +- Address Sanitizer Error (ASAN Errors): Any errors caught by address sanitizer such as memory leaks (Not acceptable) +- Aborted Runs: The test aborted/crashed during execution (Not acceptable) +- Assertion Failures: The test failed due to an assertion (Acceptable) +- Failed Runs: The test failed because the CHECK statements failed (Acceptable) +- Hangs: The program did not terminate (Not acceptable) +- Successful runs: The test succeeded (Acceptable) +- Test errors: The test failed due to bugs in the fuzz testing infrastructure (Acceptable) + +## Additional Notes +- Currently we are doing most of the fuzz testing on auto discovery strings because these strings are indirectly given by the user. +- The fuzz tests are initialized based on the current unit tests, however they should be regarded as separate tests. (i.e. You can make fuzz_test/unit_test only changes) diff --git a/fuzz_testing/fuzz_src/fuzz_testing.h b/fuzz_testing/fuzz_src/fuzz_testing.h new file mode 100644 index 00000000..011ab2e1 --- /dev/null +++ b/fuzz_testing/fuzz_src/fuzz_testing.h @@ -0,0 +1,195 @@ +/* + This header file contains all the functions that are needed for fuzz testing +*/ + +#ifndef FUZZ_TESTING_H +#define FUZZ_TESTING_H +#include +#include +#include +#include +#include +#include + +using namespace std; + +#define YML ".yml" +#define TAB " " +#define TAB_LENGTH_FOR_VAR 4 + +#define FUZZ_TEST_ERROR "Fuzz test error: " + +// groupName-testname-varName => data +inline unordered_map data_map; + +// Basically find terminating '"' that concludes the variable value +inline bool parseVariableValue(ifstream &inputFile, string &returnVal) { + string line; + while (getline(inputFile, line)) { + returnVal = returnVal + "\n" + line; + // Terminating '"' found + if (line[line.length() - 1] == '"') { + // Remove double quotes + returnVal = returnVal.substr(1, returnVal.length() - 2); + return true; + } + } + // EOF + return false; +} + +// Generate a string that combines the groupName testName and varName (i.e. +// groupName-testname-varName) +inline string generate_key(const vector &names) { + if (names.size() != 3) { + cout << FUZZ_TEST_ERROR << "Incorrect key size" << endl; + exit(1); + } + return names[0] + "--" + names[1] + "--" + names[2]; +} + +// A utility function that counts how many "TAB"s are in the current line +// A line should only have at most 2 tabs in the current schema +inline int countTabs(const string &line) { + int count = 0; + for (unsigned int i = 0; i < line.size(); i++) { + if (line[i] == ' ') { + count++; + } else { + if (count % 2 != 0) { + cout << FUZZ_TEST_ERROR << "Yaml does not have correct indentation " + << endl; + exit(1); + } + int indent = count / 2; + if (indent > 2) { + cout << FUZZ_TEST_ERROR << "Yaml has too much indentation " << endl; + exit(1); + } + return indent; + } + } + // A line of spaces, yaml error + cout << FUZZ_TEST_ERROR << "Yaml should not have an empty line " << endl; + exit(1); +} + +// A function that loads fuzzed data from a input yaml file to data_map +inline void preload(ifstream &inputFile) { + string line; + vector names; + // 0 => group, 1 => test, 2 => variable + int lastIndent = 0; + bool first = true; + + while (getline(inputFile, line)) { + // Parsing yaml file manually + int currentIndent = countTabs(line); + if (currentIndent - 1 > lastIndent) { + cout << FUZZ_TEST_ERROR << "Yaml does not have correct indentation " + << endl; + exit(1); + } + int diff = lastIndent - currentIndent; + // Delete leaf if exiting + if (diff >= 0 && !first) { + names.resize(names.size() - diff - 1); + } + stringstream ss_line(line); + string curr_name; + ss_line >> curr_name; + // If current indent level is at group/test + if (currentIndent <= 1) { + // Remove : at the end and add them to names + names.push_back(curr_name.substr(0, curr_name.length() - 1)); + } else { + // Else the current indent level must be at variable + stringstream ss(line); + string varName; + string varValue; + // Note: varName contains ':' + ss >> varName; + names.push_back(varName.substr(0, varName.length() - 1)); + // Note that varValue contains "" + getline(ss, varValue); + // Bypass leading space + varValue = varValue.substr(1, varValue.size() - 1); + string key = generate_key(names); + if (varValue[varValue.length() - 1] == '"') { + // Filter out double quote characters + data_map.insert(pair( + key, varValue.substr(1, varValue.length() - 2))); + } else { + // The variable value contains newline character + // Needs to parse multiple lines until finding the terminating '"' + if (parseVariableValue(inputFile, varValue)) { + data_map.insert(pair(key, varValue)); + } + // EOF reached but still couldn't find terminating " + else { + cout << FUZZ_TEST_ERROR + << "EOF reached but still couldn't find terminating \"" << endl; + exit(1); + } + } + } + lastIndent = currentIndent; + first = false; + } +} + +// A top level function that preloads all fuzzed data into the data_map +inline void preload_data(string fileName) { + cout << "\nPreloading fuzzed data starts: " << endl; + string pathPrefix = "../mutated_inputs/"; + string path = pathPrefix + fileName + YML; + cout << "Opening input file: " << path << endl; + ifstream inputFile(path.c_str()); + string value; + if (inputFile.is_open()) { + preload(inputFile); + cout << "Preloading fuzzed data ends " << endl; + } else { + cout << FUZZ_TEST_ERROR << "Unable to open fuzz test file! " << endl; + cout << "Make sure you run the fuzz test in the test directory" << endl; + exit(1); + } +} + +// Parameters: group name, test name, variable name +inline string load_fuzzed_value(string groupName, string testName, + string varName) { + vector names = {groupName, testName, varName}; + string key = generate_key(names); + if (data_map.find(key) == data_map.end()) { + cout << FUZZ_TEST_ERROR << "Unable to find " << key << endl; + exit(1); + } + return data_map[key]; +} + +// Cast data to corresponding data types +template +inline castType load_fuzzed_value_cast(string groupName, string testName, + string varName) { + try { + // Convert to unsigned long long then convert to castType + string value = load_fuzzed_value(groupName, testName, varName); + int base = 10; + if (value.size() > 2 && value.substr(0, 2) == "0x") { + base = 16; + } + return (castType)stoull(value, 0, base); + } catch (string e) { + cout << FUZZ_TEST_ERROR << "Data type mismatch " << endl; + cout << e << endl; + exit(1); + } +} + +// Utility functions +inline bool check_condition(bool condition, bool &check) { + check = condition; + return check; +} +#endif diff --git a/fuzz_testing/original_inputs/acl_auto_configure_fuzz_test.yml b/fuzz_testing/original_inputs/acl_auto_configure_fuzz_test.yml new file mode 100644 index 00000000..24d04c7c --- /dev/null +++ b/fuzz_testing/original_inputs/acl_auto_configure_fuzz_test.yml @@ -0,0 +1,33 @@ +auto_configure: + simple: + autodiscovery: "23 38 sample40byterandomhash000000000000000000 de4_gen2x4_swdimm 0 1 10 DDR 2 1 2 0 2048 0 - 0 0 1 5 pipe_name 1 0 32 32768 0 2 7 kernel15_dev_global 0x1000 2048 3 0 0 0 kernel15_dev_global2 0x800 1024 1 0 1 0 1 82 foo 64 128 0 192 256 0 0 1 5 7 8 1 1 4 1024 0 5 16768 0 6 2 1 4 1024 0 0 6 0 0 8 1 0 0 6 2 1 4 1024 0 0 6 0 0 4 1 0 0 6 2 1 4 1024 0 0 6 2 1 4 1024 0 0 2 2 0 d 1 d 2 2 5 1024 6 2048 0 0 0 1 0 3 1 1" + autodiscovery2: "23 23 sample40byterandomhash000000000000000000 pcie385_a7 1 1 10 DDR 2 1 2 0 2048 0 - 0 0 1 5 pipe_name 1 0 32 32768 0 1 52 bar 64 128 0 192 256 1 1 2 5 3 8 1 1 4 1024 0 5 16768 0 6 2 1 4 1024 0 0 6 0 0 4 1 0 0 2 2 0 d 1 d 0 0 2 3 5 3 32 8 4 0 0 0" + autodiscovery3: "23 18 sample40byterandomhash000000000000000000 de4_gen2x4_swdimm 0 1 6 DDR 2 1 2 0 2048 1 5 pipe_name 1 0 32 32768 1 81 foo 64 128 0 192 256 0 0 1 5 7 8 1 1 4 1024 0 5 16768 0 6 2 1 4 1024 0 0 6 0 0 8 1 0 0 6 2 1 4 1024 0 0 6 0 0 4 1 0 0 6 2 1 4 1024 0 0 6 2 1 4 1024 0 0 2 2 0 d 1 d 2 2 5 1024 6 2048 0 0 0 1 0 3 1" + many_ok_forward_compatibility: + str: "23 51 sample40byterandomhash000000000000000000 a10gx 0 1 17 DDR 2 1 6 0 2147483648 100 100 100 100 0 - 0 200 200 200 200 0 0 0 2 10 ms_dev_global1 0x800 1024 3 0 0 0 300 300 300 ms_dev_global2 0x1000 1024 1 1 1 0 300 300 300 0 0 400 400 47 40 external_sort_stage_0 0 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 external_sort_stage_1 256 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 external_sort_stage_2 512 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 external_sort_stage_3 768 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 external_sort_stage_4 1024 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 external_sort_stage_5 1280 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 external_sort_stage_6 1536 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 38 external_stream_writer0 1792 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 external_stream_writer1 2048 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 external_stream_writer2 2304 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 external_stream_writer3 2560 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 external_stream_writer4 2816 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 external_stream_writer5 3072 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 external_stream_writer6 3328 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 input_reader 3584 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 output_writer 3840 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 40 sort_stage_1 4096 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_10 4352 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_11 4608 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_12 4864 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_13 5120 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_14 5376 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_15 5632 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_16 5888 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_17 6144 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_2 6400 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_3 6656 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_4 6912 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_5 7168 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_6 7424 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_7 7680 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_8 7936 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 40 sort_stage_9 8192 128 1 0 0 1 0 1 0 1 10 0 0 4 1 0 0 0 500 500 500 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 38 stream_reader_A0 8448 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_A1 8704 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_A2 8960 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_A3 9216 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_A4 9472 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_A5 9728 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_A6 9984 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_B0 10240 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_B1 10496 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_B2 10752 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_B3 11008 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_B4 11264 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_B5 11520 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 38 stream_reader_B6 11776 256 1 0 0 0 0 1 0 1 10 2 1 8 1024 0 0 0 500 500 500 0 0 0 0 0 0 0 1 2147483647 3 1 0 0 800 800 800 900 900 900 900 900" + many_limit_check: + str: "23 19 sample40byterandomhash000000000000000000 a10gx 0 1 9 DDR 2 1 2 0 2147483648 0 - 0 0 0 0 0 0 75 31 external_sort_stage_0 0 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_1 256 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_2 512 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_3 768 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_4 1024 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_5 1280 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_6 1536 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 29 external_stream_writer0 1792 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer1 2048 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer2 2304 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer3 2560 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer4 2816 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer5 3072 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer6 3328 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 input_reader 3584 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 output_writer 2931 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 31 sort_stage_1 3196 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_10 4352 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_11 4608 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_12 4864 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_13 5120 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_14 5376 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_15 5632 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_16 5888 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_17 6144 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_2 6310 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_3 6656 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_4 6912 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_5 7168 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_6 7424 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_7 7680 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_8 7936 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_9 8192 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 29 stream_reader_A0 8448 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_A1 8704 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_A2 8960 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_A3 9216 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_A4 9472 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_A5 9728 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_A6 9984 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_B0 10231 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_B1 10496 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_B2 10752 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_B3 11008 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_B4 11264 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_B5 11520 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 stream_reader_B6 11776 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 31 external_sort_stage_0 0 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_1 256 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_2 512 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_3 768 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_4 1024 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_5 1280 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 external_sort_stage_6 1536 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 29 external_stream_writer0 1792 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer1 2048 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer2 2304 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer3 2560 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer4 2816 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer5 3072 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 external_stream_writer6 3328 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 input_reader 3584 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 29 output_writer 2931 256 1 0 0 0 0 1 0 1 6 2 1 8 1024 0 0 0 0 0 0 0 0 0 1 2147483647 3 1 31 sort_stage_1 3196 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_10 4352 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_11 4608 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_12 4864 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_13 5120 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_14 5376 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_15 5632 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_16 5888 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_17 6144 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_2 6310 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_3 6656 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_4 6912 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1" + bad_config: + str1: "0 pcie385n_a7 0 0 2 1024 0 2147483648 2147483648 4294967296 0 0 0 2 sort_stage_1 0 128 0 0 0 1 0 1 0 1 0 0 4 1 0 0 1 1 1 1 sort_stage_2 128 128 0 0 0 1 0 1 0 1 0 0 4 1 0 0 1 1 1 1" + str3: "23 15 sample40byterandomhash000000000000000000 a10gx 0 1 7 DDR 2 1 2 0 2147483648 0 0 0 0 1 31 external_sort_stage_0 0 128 1 0 0 0 1 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1" + multi_mem_config: + str: "23 56 sample40byterandomhash000000000000000000 pcie385n_a7 0 4 10 SVM 0 1 2 0 1073741824 0 - 1 SVM2 11 DDR 2 2 24 1 2 1073741824 3221225472 3221225472 5368709120 0 15 QDR 2 4 48 0 2 5368709120 5369757696 5369757696 5370806272 5370806272 5371854848 5371854848 5372903424 4 9 SVM2 0 1 2 0 1073741824 0 SVM 0 0 0 0 2 31 sort_stage_1 0 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 31 sort_stage_2 128 128 1 0 0 1 0 1 0 1 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1" + kernel_arg_info: + str1: "23 15 sample40byterandomhash000000000000000000 a10gx 0 1 7 DDR 2 1 2 0 2147483648 0 0 0 1 2 34 external_sort_stage_0 0 128 1 0 0 1 0 1 0 2 9 0 0 4 1 0 0 arg_one type_one 1 9 0 0 4 1 0 0 arg_two type_two 2 0 0 0 0 1 1 1 3 1 1 1 3 1 34 external_sort_stage_1 256 128 1 0 0 1 0 1 0 2 9 0 0 4 1 0 0 arg_three type_three 1 9 0 0 4 1 0 0 arg_four type_four 2 0 0 0 0 1 1 1 3 1 1 1 3 1" + str2: "23 15 sample40byterandomhash000000000000000000 a10gx 0 1 7 DDR 2 1 2 0 2147483648 0 0 0 0 2 34 external_sort_stage_0 0 128 1 0 0 1 0 1 0 2 6 0 0 4 1 0 0 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 34 external_sort_stage_1 256 128 1 0 0 1 0 1 0 2 6 0 0 4 1 0 0 6 0 0 4 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1" + hostpipe_basic: + str: "23 50 sample40byterandomhash000000000000000000 a10gx_hostpipe 0 1 15 DDR 2 1 6 0 2147483648 0 100 100 100 100 200 200 200 200 2 9 host_to_dev 1 0 32 32768 300 300 300 300 dev_to_host 0 1 32 32768 300 300 300 300 400 1 7 dev_global_3 0x400 2048 0 0 0 0 1 29 foo 0 128 1 0 0 1 0 1 0 0 0 0 0 0 1 1 1 3 1 1 1 3 1 0 0 800 800 800 900 900" + streaming_basic: + config_str: "23 30 sample40byterandomhash000000000000000000 pac_a10 0 1 13 DDR 2 2 24 1 2 0 4294967296 4294967296 8589934592 0 - 0 0 0 0 1 7 device_global_name 0x100 128 0 0 0 0 1 105 _ZTS3CRCILi0EE 0 256 1 0 0 1 0 1 0 9 8 0 0 8 1 0 0 1 k0_ZTS3CRCILi0EE_arg0 8 2 1 8 1024 0 3 1 k0_ZTS3CRCILi0EE_arg1 8 0 0 8 1 0 0 1 k0_ZTS3CRCILi0EE_arg2 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 7 2 1 8 1024 0 2 0 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 0 0 1 2 64 4096 1 1 1 3 1 1 1 3 1 0 1 k0_ZTS3CRCILi0EE_streaming_start k0_ZTS3CRCILi0EE_streaming_done" + one_streaming_arg_and_streaming_kernel: + config_str: "23 27 531091a097f0d7096b21f349b4b283f9e206ebc0 pac_s10 0 1 17 DDR 2 4 24 1 2 0 8589934592 8589934592 17179869184 17179869184 25769803776 25769803776 34359738368 0 - 0 0 0 0 0 0 1 125 _ZTS15binomial_kernel 0 256 0 0 0 0 0 1 0 8 7 2 1 8 1024 0 2 0 8 0 0 8 1 0 0 1 k0_ZTS15binomial_kernel_arg1 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 7 2 1 8 1024 0 2 0 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 0 0 16 2 64 8196 65 8196 66 8196 67 8196 68 8196 69 8196 70 8196 71 8196 72 8196 73 8196 74 8196 75 8196 76 8196 77 8196 78 8196 79 8196 1 1 1 3 1 1 1 3 1 1 1 k0_ZTS15binomial_kernel_streaming_start k0_ZTS15binomial_kernel_streaming_done" + two_streaming_args_and_streaming_kernel: + config_str: "23 27 531091a097f0d7096b21f349b4b283f9e206ebc0 pac_s10 0 1 17 DDR 2 4 24 1 2 0 8589934592 8589934592 17179869184 17179869184 25769803776 25769803776 34359738368 0 - 0 0 0 0 0 0 1 126 _ZTS15binomial_kernel 0 256 0 0 0 0 0 1 0 8 8 2 1 8 1024 0 2 1 k0_ZTS15binomial_kernel_arg0 8 0 0 8 1 0 0 1 k0_ZTS15binomial_kernel_arg1 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 7 2 1 8 1024 0 2 0 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 0 0 16 2 64 8196 65 8196 66 8196 67 8196 68 8196 69 8196 70 8196 71 8196 72 8196 73 8196 74 8196 75 8196 76 8196 77 8196 78 8196 79 8196 1 1 1 3 1 1 1 3 1 1 1 k0_ZTS15binomial_kernel_streaming_start k0_ZTS15binomial_kernel_streaming_done" + two_streaming_args_and_non_streaming_kernel: + config_str: "23 27 531091a097f0d7096b21f349b4b283f9e206ebc0 pac_s10 0 1 17 DDR 2 4 24 1 2 0 8589934592 8589934592 17179869184 17179869184 25769803776 25769803776 34359738368 0 - 0 0 0 0 0 0 1 124 _ZTS15binomial_kernel 0 256 0 0 0 0 0 1 0 8 8 2 1 8 1024 0 2 1 k0_ZTS15binomial_kernel_arg0 8 0 0 8 1 0 0 1 k0_ZTS15binomial_kernel_arg1 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 7 2 1 8 1024 0 2 0 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 7 0 0 8 1 0 0 0 0 0 16 2 64 8196 65 8196 66 8196 67 8196 68 8196 69 8196 70 8196 71 8196 72 8196 73 8196 74 8196 75 8196 76 8196 77 8196 78 8196 79 8196 1 1 1 3 1 1 1 3 1 1 0" + cra_ring_root_not_exist: + config_str: "23 50 2ccb683dee8e34a004c1a27e6d722090a8cc684d custom_ipa 0 1 9 0 2 1 2 0 2199023255552 3 - 0 6 5 ZTSZ4mainE4MyIP_arg_input_a 1 0 8 32768 ZTSZ4mainE4MyIP_arg_input_b 1 0 8 32768 ZTSZ4mainE4MyIP_arg_input_c 1 0 8 32768 ZTSZ4mainE4MyIP_arg_n 1 0 4 32768 ZTSZ4mainE4MyIP_streaming_start 1 0 0 32768 ZTSZ4mainE4MyIP_streaming_done 0 1 0 32768 0 0 0 0 1 64 _ZTSZ4mainE4MyIP 0 128 1 0 0 1 0 1 0 4 8 2 1 8 4 0 0 1 ZTSZ4mainE4MyIP_arg_input_a 8 2 1 8 4 0 0 1 ZTSZ4mainE4MyIP_arg_input_b 8 2 1 8 4 0 0 1 ZTSZ4mainE4MyIP_arg_input_c 8 0 0 4 1 0 0 1 ZTSZ4mainE4MyIP_arg_n 0 0 0 0 1 1 1 3 1 1 1 3 1 1 1 ZTSZ4mainE4MyIP_streaming_start ZTSZ4mainE4MyIP_streaming_done" + cra_ring_root_exist: + config_str: "23 50 2ccb683dee8e34a004c1a27e6d722090a8cc684d custom_ipa 0 1 9 0 2 1 2 0 2199023255552 3 - 0 6 5 ZTSZ4mainE4MyIP_arg_input_a 1 0 8 32768 ZTSZ4mainE4MyIP_arg_input_b 1 0 8 32768 ZTSZ4mainE4MyIP_arg_input_c 1 0 8 32768 ZTSZ4mainE4MyIP_arg_n 1 0 4 32768 ZTSZ4mainE4MyIP_streaming_start 1 0 0 32768 ZTSZ4mainE4MyIP_streaming_done 0 1 0 32768 0 0 0 1 1 64 _ZTSZ4mainE4MyIP 0 128 1 0 0 1 0 1 0 4 8 2 1 8 4 0 0 1 ZTSZ4mainE4MyIP_arg_input_a 8 2 1 8 4 0 0 1 ZTSZ4mainE4MyIP_arg_input_b 8 2 1 8 4 0 0 1 ZTSZ4mainE4MyIP_arg_input_c 8 0 0 4 1 0 0 1 ZTSZ4mainE4MyIP_arg_n 0 0 0 0 1 1 1 3 1 1 1 3 1 1 1 ZTSZ4mainE4MyIP_streaming_start ZTSZ4mainE4MyIP_streaming_done" + hostpipe_mappings: + config_str: "23 66 sample40byterandomhash000000000000000000 pac_a10 0 1 13 DDR 2 2 24 1 2 0 4294967296 4294967296 8589934592 0 - 0 0 0 0 0 0 1 5 8 pipe_logical_name1 pipe_physical_name1 1 12345 0 1 4 10 pipe_logical_name2 pipe_physical_name2 0 12323 1 0 8 20 pipe_logical_name3 pipe_physical_name1 1 12313 0 1 4 10 pipe_logical_name5 pipe_physical_name1 0 12316 1 0 8 20 pipe_logical_name4 pipe_physical_name3 0 12342 0 1 4 10 3 90 _ZTS3CRCILi0EE 512 256 1 0 0 1 0 1 0 9 6 0 0 8 1 0 0 6 2 1 8 1024 0 3 6 0 0 8 1 0 0 6 0 0 8 1 0 0 6 0 0 8 1 0 0 6 2 1 8 1024 0 2 6 0 0 8 1 0 0 6 0 0 8 1 0 0 6 0 0 8 1 0 0 0 0 1 2 64 4096 1 1 1 3 1 1 1 3 1 0 64 _ZTS11LZReductionILi0EE 0 256 1 0 0 0 0 1 0 5 6 0 0 8 1 0 0 6 2 1 8 1024 0 3 6 0 0 8 1 0 0 6 0 0 8 1 0 0 6 0 0 8 1 0 0 0 0 2 2 64 131072 65 32768 1 1 1 3 1 1 1 3 1 0 125 _ZTS13StaticHuffmanILi0EE 256 256 1 0 0 1 0 1 0 10 6 0 0 8 1 0 0 6 0 0 4 1 0 0 6 2 1 8 1024 0 2 6 0 0 8 1 0 0 6 0 0 8 1 0 0 6 0 0 8 1 0 0 6 2 1 8 1024 0 2 6 0 0 8 1 0 0 6 0 0 8 1 0 0 6 0 0 8 1 0 0 0 0 15 2 64 116 65 116 66 1152 67 512 68 256 69 120 70 120 71 1152 72 116 73 1152 74 512 75 256 76 120 77 120 78 1152 1 1 1 3 1 1 1 3 1 0" diff --git a/fuzz_testing/script/fuzz_test.py b/fuzz_testing/script/fuzz_test.py new file mode 100644 index 00000000..7428c09f --- /dev/null +++ b/fuzz_testing/script/fuzz_test.py @@ -0,0 +1,269 @@ +# Please run this script at the shell directory in build + +import yaml +import argparse +import subprocess +import os +import re +import time + +parser = argparse.ArgumentParser(description='Fuzz testing') +parser.add_argument("file_name", type=str, nargs='*', help="Unit test file name") +parser.add_argument("--group", type=str, nargs='?', help="Group name") +parser.add_argument("--test", type=str, nargs='?', help="Test name") +parser.add_argument("-n", type=int, nargs='?', help="Number of times a single variable is mutated and tested") +parser.add_argument('--all', action='store_true') +parser.add_argument("--coverage", action='store_true', help="Generate Coverage report using lcov") +args = parser.parse_args() + +all_tests = [ + 'acl_auto_configure_fuzz_test', + ] + +# First level: unit test file name => table of attributes +# Second level: attribute name => value +# Attributes: +# - Total runs +# - Successful runs +# - Failed runs +# - ASAN errors +# - Hangs +# - Test errors +results_dictionary = dict() + +TOTAL_RUNS = "Total runs" +SUCCESSFUL_RUNS = "Successful runs" +FAILED_RUNS = "Failed runs" +ABORTED_RUNS = "Aborted runs" +ASSERTION_FAILURES = "Assertion failures" +ASAN_ERRORS = "ASAN errors" +HANGS = "Hangs" +TEST_ERRORS = "Test errors" + +SUCCESS_TEST_PATTERN = "OK .* tests, .* ran, .* checks, .* ignored, .* filtered out, .* ms" +FAILED_TEST_PATTERN = "Errors .* tests, .* ran, .* checks, .* ignored, .* filtered out, .* ms" +ASAN_ERROR_PATTERN = "SUMMARY: AddressSanitizer" +FUZZ_TEST_ERROR_PATTERN = "Fuzz test error" +ASSERTION_FAILURE_PATTERN = "acl_fuzz_test:.*Assertion" + +# Give it 60 seconds to finish a single unit test +TIMEOUT = 60 +TIMEOUT_COMMAND = "timeout " + str(TIMEOUT) +TIMEOUT_MESSAGE = "TIMEOUT! " + str(TIMEOUT) + " seconds have passed! " + +if args.all and args.file_name != []: + print("You should not have file_name arguments when using --all flag") + exit(1) + +def initialize_attribute_table(): + return { + TOTAL_RUNS: 0, + SUCCESSFUL_RUNS: 0, + FAILED_RUNS: 0, + ABORTED_RUNS : 0, + ASSERTION_FAILURES: 0, + ASAN_ERRORS: 0, + HANGS: 0, + TEST_ERRORS: 0, + } + +def parse_output(output): + # Errors (1 failures, 230 tests, 1 ran, 5 checks, 0 ignored, 229 filtered out, 3 ms) + return + +def encode_if_condition(pattern, condition): + if condition: + return pattern.encode() + return pattern + +def fuzz_test_main(test_file_name, group, test, var, original_value, all_outputs): + # Initialize results + if test_file_name not in results_dictionary: + results_dictionary[test_file_name] = initialize_attribute_table() + + # Mutate variable + fuzz_var_str = "Fuzzing variable: " + group + "--" + test + "--" + var + print(" " + fuzz_var_str) + mutation_command = ["python", "mutator.py", test_file_name, group, test, var] + # mutated_value = subprocess.check_output(mutation_command) + original_value_message = "Original value: \n" + original_value + subprocess.check_output(mutation_command) + print(" Mutation finished! ") + with open("temp.txt", "rb") as file: + mutated_value = file.read() + mutated_value_message = "Mutated value: " + os.chdir("../test") + # Run test with timeout + start = time.time() + end = start + time_taken = end - start + run_test_command = TIMEOUT_COMMAND + " ./acl_fuzz_test -v -g " + group + " -n " + test + " " + print(" Test command: " + run_test_command) + # Save output + try: + output = subprocess.check_output(run_test_command.split(), stderr=subprocess.STDOUT) + except Exception as e: + output = str(e.output) + end = time.time() + time_taken = end - start + os.chdir("../script") + print(" Test finished! ") + # print(output) + all_outputs.append(fuzz_var_str) + all_outputs.append(original_value_message) + all_outputs.append(mutated_value_message) + all_outputs.append(mutated_value) + all_outputs.append("\n") + encoded = type(output) != str + all_outputs.append(output) + all_outputs.append("\n") + timeout = False + if time_taken >= TIMEOUT: + timeout = True + all_outputs.append(TIMEOUT_MESSAGE) + + # Update results + # Total run + results_dictionary[test_file_name][TOTAL_RUNS] = results_dictionary[test_file_name][TOTAL_RUNS] + 1 + # If successful test message found + if re.search(encode_if_condition(SUCCESS_TEST_PATTERN, encoded), output): + # If failed test message found + if re.search(encode_if_condition(FAILED_TEST_PATTERN, encoded), output): + # Test error + 1 + results_dictionary[test_file_name][TEST_ERRORS] = results_dictionary[test_file_name][TEST_ERRORS] + 1 + test_error = True + # If failed test message not found + else: + # Successful test + 1 + results_dictionary[test_file_name][SUCCESSFUL_RUNS] = results_dictionary[test_file_name][SUCCESSFUL_RUNS] + 1 + # If successful test message not found + else: + # If failed test message found + if re.search(encode_if_condition(FAILED_TEST_PATTERN, encoded), output): + # Failed test + 1 + results_dictionary[test_file_name][FAILED_RUNS] = results_dictionary[test_file_name][FAILED_RUNS] + 1 + # If failed test message not found + else: + # If timeout + if timeout: + # Hang + 1 + results_dictionary[test_file_name][HANGS] = results_dictionary[test_file_name][HANGS] + 1 + # If not timeout + else: + # If assertion message found + if re.search(encode_if_condition(ASSERTION_FAILURE_PATTERN, encoded), output): + # Assertion failures + 1 + results_dictionary[test_file_name][ASSERTION_FAILURES] = results_dictionary[test_file_name][ASSERTION_FAILURES] + 1 + # If assertion message not found + else: + # Aborted run + 1 + results_dictionary[test_file_name][ABORTED_RUNS] = results_dictionary[test_file_name][ABORTED_RUNS] + 1 + # ASAN errors + if re.search(encode_if_condition(ASAN_ERROR_PATTERN, encoded), output): + results_dictionary[test_file_name][ASAN_ERRORS] = results_dictionary[test_file_name][ASAN_ERRORS] + 1 + # Test errors + if re.search(encode_if_condition(FUZZ_TEST_ERROR_PATTERN, encoded), output) and not test_error: + results_dictionary[test_file_name][TEST_ERRORS] = results_dictionary[test_file_name][TEST_ERRORS] + 1 + +def load_yaml(test_file_name): + # Fetch data from original_inputs + original_file_path = "../original_inputs/" + test_file_name + ".yml" + with open(original_file_path, 'r') as file: + inputs = yaml.safe_load(file) + return inputs + +def store_outputs(test_file_name, outputs): + # Print outputs to test_outputs + test_outputs_path = "../test_outputs" + test_outputs_file_path = test_outputs_path + "/" + test_file_name + ".txt" + + # Output to test_outputs directory + if not os.path.exists(test_outputs_path): + os.makedirs(test_outputs_path) + with open(test_outputs_file_path, 'w') as file: + for out in outputs: + if type(out) != str: + file.write(out.decode("utf-8", errors='replace')) + else: + file.write(out) + file.write("\n") + +def fuzz_test_iterations(test_file_name, group, test, var, original_value, all_outputs, iterations=1, indents=""): + for i in range(iterations): + print(indents + "Iteration: " + str(i+1) + " / " + str(iterations)) + fuzz_test_main(test_file_name, group, test, var, original_value, all_outputs) + +def fuzz_test(test_file_name, iterations=1): + inputs = load_yaml(test_file_name) + + # Save all output strings + all_outputs = [] + + print("Running Fuzz tests for " + test_file_name) + + group_total = len(inputs) + group_count = 1 + + for group in inputs: + print("Group: {} ({} / {})".format(group, group_count, group_total)) + test_total = len(inputs[group]) + test_count = 1 + for test in inputs[group]: + print(" Test: {} ({} / {})".format(test, test_count, test_total)) + for var in inputs[group][test]: + fuzz_test_iterations(test_file_name, group, test, var, inputs[group][test][var], all_outputs, iterations, " ") + test_count += 1 + group_count += 1 + + store_outputs(test_file_name, all_outputs) + +def fuzz_test_single(test_file_name, group, test, iterations=1): + all_outputs = [] + inputs = load_yaml(test_file_name) + for var in inputs[group][test]: + fuzz_test_iterations(test_file_name, group, test, var, inputs[group][test][var], all_outputs, iterations, "") + store_outputs(test_file_name, all_outputs) + +def generate_coverage_report(): + print("Generating Coverage Report...") + os.chdir("../test") + clean_command = "rm -f coverage.info && rm -rf coverage_report" + subprocess.check_output(clean_command.split()) + lcov_command = "lcov --capture --directory CMakeFiles/acl_fuzz_test.dir --output-file=coverage.info" + subprocess.check_output(lcov_command.split()) + generate_html_command = "genhtml coverage.info --output-directory=coverage_report" + subprocess.check_output(generate_html_command.split()) + print("Coverage Report generated at " + os.getcwd() + "/coverage_report/index.html ") + os.chdir("../script") + +def generate_results_yml(): + print("Generating results yml...") + results_directory_path = "../results" + if not os.path.exists(results_directory_path): + os.makedirs(results_directory_path) + os.chdir(results_directory_path) + results_yml_path = "results.yml" + with open(results_yml_path, 'w') as outfile: + yaml.safe_dump(results_dictionary, outfile, default_flow_style=False) + print("Results yaml file generated at " + os.getcwd() + "/" + results_yml_path) + os.chdir("../script") + +def main(): + tests = args.file_name + iterations = 1 + if args.n: + iterations = args.n + if args.all: + tests = all_tests + + # TODO: Error checking + if args.group and args.test and len(tests) == 1: + fuzz_test_single(tests[0], args.group, args.test, iterations) + else: + for test in tests: + fuzz_test(test, iterations) + if args.coverage: + generate_coverage_report() + generate_results_yml() + +main() diff --git a/fuzz_testing/script/mutator.py b/fuzz_testing/script/mutator.py new file mode 100644 index 00000000..3bf1a633 --- /dev/null +++ b/fuzz_testing/script/mutator.py @@ -0,0 +1,59 @@ +import yaml +import argparse +import subprocess +import os + +parser = argparse.ArgumentParser(description='Mutate single variable') +parser.add_argument("file_name", type=str, help="Unit test file name") +parser.add_argument("group_name", type=str, help="Unit test group name") +parser.add_argument("test_name", type=str, help="Specific unit test name") +parser.add_argument("variable_name", type=str, help="Variable name") +args = parser.parse_args() + +# Fetch data from original_inputs +original_file_path = "../original_inputs/" + args.file_name + ".yml" + +with open(original_file_path, 'r') as file: + inputs = yaml.safe_load(file) + +value = inputs[args.group_name][args.test_name][args.variable_name] + +# Mutation +# Pipe the output to a txt file, we can not save this output as a variable in python3 because +# as it might be converted to a binary +if os.path.exists("temp.txt"): + os.remove("temp.txt") +subprocess.check_output("echo {0} | radamsa > temp.txt".format(value), shell=True) +# Remove last newline character +subprocess.check_output("perl -pi -e 'chomp if eof' temp.txt", shell=True) +mutated_key = args.group_name + "--" + args.test_name + "--" + args.variable_name + +# Output to mutated_inputs +mutated_inputs_path = "../mutated_inputs" +mutated_inputs_file_path = mutated_inputs_path + "/" + args.file_name + ".yml" +if not os.path.exists(mutated_inputs_path): + os.makedirs(mutated_inputs_path) +tab = " " +col = ":" +nl = "\n" +if os.path.exists(mutated_inputs_file_path): + os.remove(mutated_inputs_file_path) +for group in inputs: + with open(mutated_inputs_file_path, 'a+') as file: + file.write(group + col + nl) + for test in inputs[group]: + with open(mutated_inputs_file_path, 'a+') as file: + file.write(tab + test + col + nl) + for var in inputs[group][test]: + with open(mutated_inputs_file_path, 'a+') as file: + file.write(tab + tab + var + col + " \"") + # file.write(inputs[group][test][var]) + current_key = group + "--" + test + "--" + var + if current_key == mutated_key: + command = "cat temp.txt >> " + mutated_inputs_file_path + subprocess.check_output(command, shell=True) + else: + with open(mutated_inputs_file_path, 'a+') as file: + file.write(inputs[group][test][var]) + with open(mutated_inputs_file_path, 'a+') as file: + file.write("\"" + nl) diff --git a/fuzz_testing/test/CMakeLists.txt b/fuzz_testing/test/CMakeLists.txt new file mode 100755 index 00000000..f7e3ab82 --- /dev/null +++ b/fuzz_testing/test/CMakeLists.txt @@ -0,0 +1,41 @@ +# Copyright (C) 2021 Intel Corporation +# SPDX-License-Identifier: BSD-3-Clause + +add_library(test_acl_fuzz_test_header INTERFACE) +target_include_directories(test_acl_fuzz_test_header INTERFACE .) + +add_executable(acl_fuzz_test + $ + acl_auto_configure_fuzz_test.cpp + acl_globals_fuzz_test.cpp + acl_hal_fuzz_test.cpp + acl_fuzz_test.cpp + ) +set_target_properties(acl_fuzz_test PROPERTIES CXX_EXTENSIONS OFF) +target_compile_features(acl_fuzz_test PRIVATE cxx_std_11) +target_compile_definitions(acl_fuzz_test PRIVATE + _GLIBCXX_USE_CXX11_ABI=0 + "ACL_TARGET_BIT=${ACL_TARGET_BIT}" + CL_USE_DEPRECATED_OPENCL_1_0_APIS=1 + CL_USE_DEPRECATED_OPENCL_1_1_APIS=1 + CL_USE_DEPRECATED_OPENCL_1_2_APIS=1 + CL_TARGET_OPENCL_VERSION=300 + ) +target_include_directories(acl_fuzz_test PRIVATE + "${CMAKE_BINARY_DIR}/include" + "${CMAKE_SOURCE_DIR}/src" + ) +target_link_libraries(acl_fuzz_test PRIVATE + acl_headers + acl_check_sys_cmd + acl_hash + acl_threadsupport + CppUTest + pkg_editor + ) + +# Fuzz tests should not be added to regular unit testing, it should be explicitly called +# add_test(NAME acl_fuzz_test COMMAND acl_fuzz_test -v) +# set_property(TEST acl_fuzz_test PROPERTY ENVIRONMENT +# "AOCL_BOARD_PACKAGE_ROOT=${CMAKE_CURRENT_SOURCE_DIR}/board/a10_ref" +# ) diff --git a/fuzz_testing/test/acl_auto_configure_fuzz_test.cpp b/fuzz_testing/test/acl_auto_configure_fuzz_test.cpp new file mode 100644 index 00000000..fe20a95f --- /dev/null +++ b/fuzz_testing/test/acl_auto_configure_fuzz_test.cpp @@ -0,0 +1,1487 @@ +// Copyright (C) 2011-2021 Intel Corporation +// SPDX-License-Identifier: BSD-3-Clause + +/* + This is a fuzz test that is built based off the original unit test version. + Many CHECK() are commented out because if these failed, then the program would + exit immediately without correctly destructing non-primitive type variables + (i.e. strings), which causes Address Sanitizer (ASAN) errors. However, in fuzz + testing, it is very common for these CHECKs to fail as the input data is + mutated. Therefore, CHECK is replaced with check_condition() from + fuzz_testing.h. For most tests, we capture the content of the test inside a + scope (i.e. {}), then use "break" to break out of scope when check_condition + fails. Non-primitive variables defined inside the scope will be destructed at + that time. Finally, we use a final CHECK to see if the test passes, as it is + required by the unit test framework (CppUTest). This workaround prevents the + fuzz testing infrastructure to catch ASAN errors that are introduced by the + unit test framework instead of the source code. +*/ + +#ifdef _MSC_VER +#pragma warning(push) +#pragma warning(disable : 4100) // unreferenced formal parameter +#endif +#include +#ifdef _MSC_VER +#pragma warning(pop) +#endif + +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#include "acl_fuzz_test.h" + +#include "../fuzz_src/fuzz_testing.h" +#include + +TEST_GROUP(auto_configure) { +public: + // preload_data is a function that is specific to fuzz test only (See + // fuzz_testing.h) + void setup() { preload_data("acl_auto_configure_fuzz_test"); } + void teardown() { acl_test_run_standard_teardown_checks(); } + +protected: + acl_device_def_t m_device_def; +}; + +TEST(auto_configure, simple) { + +#define VERSIONIDSTRINGIFY(x) #x +#define VERSIONIDTOSTR(x) VERSIONIDSTRINGIFY(x) +#define DEVICE_FIELDS " 23" +#define DEVICE_FIELDS_DEV_GLOBAL " 38" +#define DEVICE_FIELDS_OLD " 18" +#define BOARDNAME "de4_gen2x4_swdimm" +#define BOARDNAME2 "pcie385_a7" +#define RANDOM_HASH " sample40byterandomhash000000000000000000" +#define IS_BIG_ENDIAN " 1" +#define IS_NOT_BIG_ENDIAN " 0" +#define MEM_BACKWARDS_COMP " 1 6 DDR 2 1 2 0 2048" +#define MEM " 1 10 DDR 2 1 2 0 2048 0 - 0 0" +#define HOSTPIPE " 1 5 pipe_name 1 0 32 32768" +#define KERNEL_ARG_INFO_NONE " 0" +#define ARG_INT " 6 0 0 4 1 0 0" +#define ARG_LONG " 6 0 0 8 1 0 0" +#define ARG_LOCAL " 8 1 1 4 1024 0 5 16768 0" +#define ARG_GLOBAL " 6 2 1 4 1024 0 0" +#define ARG_CONST " 6 3 1 4 1024 0 0" +#define ARG_PROF ARG_GLOBAL +#define ARGS_LOCAL_GLOBAL_INT " 3" ARG_LOCAL ARG_GLOBAL ARG_INT + // the last four args are for printf(2) and profiling(2) +#define ARGS_LOCAL_GLOBAL_LONG_PROF \ + " 7" ARG_LOCAL ARG_GLOBAL ARG_LONG ARG_GLOBAL ARG_INT ARG_PROF ARG_PROF + +#define KERNEL_FIELDS " 72" +#define KERNEL_CRA " 64 128" // Address(offset) and number of bytes +#define KERNEL_FAST_LAUNCH_DEPTH " 0" +#define KERNEL_PERF_MON " 192 256" // Address(offset) and number of bytes +#define KERNEL_WORKGROUP_INVARIANT " 1" +#define KERNEL_WORKITEM_INVARIANT " 1" +#define KERNEL_WORKGROUP_VARIANT " 0" +#define KERNEL_WORKITEM_VARIANT " 0" +#define KERNEL_NUM_VECTOR_LANES1 " 1" +#define KERNEL_NUM_VECTOR_LANES2 " 2" + +#define KERNEL_PRINTF_FORMAT1 " 0 d" +#define KERNEL_PRINTF_FORMAT2 " 1 d" +#define KERNEL_PRINTF_FORMATSTRINGS \ + " 2 2" KERNEL_PRINTF_FORMAT1 KERNEL_PRINTF_FORMAT2 +#define KERNEL_PROFILE_SCANCHAIN_LENGTH " 5" + +// statically determined demand for local memory +#define LD_1024 " 2 2 5 1024 6 2048" +#define LD_0 " 0 0" + +// kernel attribute reqd_work_group_size(x,y,z). 0,0,0 means not specified. +#define KERNEL_REQD_WORK_GROUP_SIZE_NONE " 0 0 0" +#define KERNEL_REQD_WORK_GROUP_SIZE_235 " 2 3 5" + +// kernel attribute max_work_group_size(x). 0 means not specified. +#define KERNEL_MAX_WORK_GROUP_SIZE_NONE " 1 0" +#define KERNEL_MAX_WORK_GROUP_SIZE_1024 " 3 32 8 4" + +// kernel attribute max_global_work_dim(n). 3 means not specified (or 3). +#define KERNEL_MAX_GLOBAL_WORK_DIM_NONE " 3" +#define KERNEL_MAX_GLOBAL_WORK_DIM_ZERO " 0" +#define KERNEL_USES_GLOBAL_WORK_OFFSET_ENABLED " 1" +#define KERNEL_USES_GLOBAL_WORK_OFFSET_DISABLED " 0" + +// sycl compile +#define IS_SYCL_COMPILE " 1" +#define IS_NOT_SYCL_COMPILE " 0" + +// Device global autodiscovery entries +#define NUM_DEV_GLOBAL " 2" +#define NUM_DEV_GLOBAL_FIELD " 7" +// The 7 fields are dev_globa_name, address, size, host_access, +// can_skip_programming, implement_in_csr, reset_on_reuse +#define DEV_GLOBAL_1 " kernel15_dev_global 0x1000 2048 3 0 0 0" +#define DEV_GLOBAL_2 " kernel15_dev_global2 0x800 1024 1 0 1 0" + int parsed; + char *err_str_c; + { + std::string autodiscovery = + load_fuzzed_value("auto_configure", "simple", "autodiscovery"); + std::string err_str; + ACL_LOCKED(parsed = acl_load_device_def_from_str( + autodiscovery, m_device_def.autodiscovery_def, err_str)); + // Using char* to save err_str because char* is primitive type which + // will be destructed properly when CHECK fails + err_str_c = const_cast(err_str.c_str()); + } + + CHECK_EQUAL(1, parsed); + + CHECK_EQUAL(1, m_device_def.autodiscovery_def.num_global_mem_systems); + CHECK_EQUAL(0, m_device_def.autodiscovery_def.global_mem_defs[0].range.begin); + CHECK_EQUAL((void *)2048, + m_device_def.autodiscovery_def.global_mem_defs[0].range.next); + CHECK_EQUAL( + (acl_system_global_mem_allocation_type_t)0, + m_device_def.autodiscovery_def.global_mem_defs[0].allocation_type); + CHECK("" == + m_device_def.autodiscovery_def.global_mem_defs[0].primary_interface); + CHECK_EQUAL( + 0, + m_device_def.autodiscovery_def.global_mem_defs[0].can_access_list.size()); + + CHECK(BOARDNAME == m_device_def.autodiscovery_def.name); + + CHECK_EQUAL(0, (int)m_device_def.autodiscovery_def.is_big_endian); + + CHECK_EQUAL(1, (int)m_device_def.autodiscovery_def.accel.size()); + CHECK_EQUAL(1, (int)m_device_def.autodiscovery_def.hal_info.size()); + + // Check HAL's view + CHECK("foo" == m_device_def.autodiscovery_def.hal_info[0].name); + CHECK_EQUAL(64, (int)m_device_def.autodiscovery_def.hal_info[0].csr.address); + CHECK_EQUAL(128, + (int)m_device_def.autodiscovery_def.hal_info[0].csr.num_bytes); + CHECK_EQUAL(192, + (int)m_device_def.autodiscovery_def.hal_info[0].perf_mon.address); + CHECK_EQUAL( + 256, (int)m_device_def.autodiscovery_def.hal_info[0].perf_mon.num_bytes); + + // Check hostpipe info + CHECK_EQUAL(1, m_device_def.autodiscovery_def.acl_hostpipe_info.size()); + CHECK("pipe_name" == + m_device_def.autodiscovery_def.acl_hostpipe_info[0].name); + CHECK_EQUAL( + true, m_device_def.autodiscovery_def.acl_hostpipe_info[0].is_host_to_dev); + CHECK_EQUAL( + false, + m_device_def.autodiscovery_def.acl_hostpipe_info[0].is_dev_to_host); + CHECK_EQUAL(32, + m_device_def.autodiscovery_def.acl_hostpipe_info[0].data_width); + CHECK_EQUAL( + 32768, + m_device_def.autodiscovery_def.acl_hostpipe_info[0].max_buffer_depth); + + // Check ACL's view + CHECK_EQUAL(0, m_device_def.autodiscovery_def.accel[0].id); + CHECK_EQUAL(0, m_device_def.autodiscovery_def.accel[0].mem.begin); + CHECK_EQUAL( + (void *)0x020000, + m_device_def.autodiscovery_def.accel[0] + .mem.next); // Not sure why this isn't 16KB like OpenCL spec minimum + + CHECK_EQUAL( + 2, (int)m_device_def.autodiscovery_def.accel[0].local_aspaces.size()); + CHECK_EQUAL( + 5, + (int)m_device_def.autodiscovery_def.accel[0].local_aspaces[0].aspace_id); + CHECK_EQUAL(1024, (int)m_device_def.autodiscovery_def.accel[0] + .local_aspaces[0] + .static_demand); + CHECK_EQUAL( + 6, + (int)m_device_def.autodiscovery_def.accel[0].local_aspaces[1].aspace_id); + CHECK_EQUAL(2048, (int)m_device_def.autodiscovery_def.accel[0] + .local_aspaces[1] + .static_demand); + + CHECK("foo" == m_device_def.autodiscovery_def.accel[0].iface.name); + CHECK_EQUAL( + 0, (int)m_device_def.autodiscovery_def.accel[0].is_workgroup_invariant); + CHECK_EQUAL( + 0, (int)m_device_def.autodiscovery_def.accel[0].is_workitem_invariant); + CHECK_EQUAL(3, + (int)m_device_def.autodiscovery_def.accel[0].max_global_work_dim); + CHECK_EQUAL( + 5, + (int)m_device_def.autodiscovery_def.accel[0].profiling_words_to_readback); + CHECK_EQUAL(7, + (int)m_device_def.autodiscovery_def.accel[0].iface.args.size()); + + CHECK_EQUAL(1, + m_device_def.autodiscovery_def.accel[0].iface.args[0].addr_space); + CHECK_EQUAL(1, + m_device_def.autodiscovery_def.accel[0].iface.args[0].category); + CHECK_EQUAL(4, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[0].size); + CHECK_EQUAL( + 5, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[0].aspace_number); + CHECK_EQUAL(16768, (int)m_device_def.autodiscovery_def.accel[0] + .iface.args[0] + .lmem_size_bytes); + + CHECK_EQUAL(2, + m_device_def.autodiscovery_def.accel[0].iface.args[1].addr_space); + CHECK_EQUAL(1, + m_device_def.autodiscovery_def.accel[0].iface.args[1].category); + CHECK_EQUAL(4, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[1].size); + + CHECK_EQUAL(0, + m_device_def.autodiscovery_def.accel[0].iface.args[2].addr_space); + CHECK_EQUAL(0, + m_device_def.autodiscovery_def.accel[0].iface.args[2].category); + CHECK_EQUAL(8, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[2].size); + + // printf buffer start address + CHECK_EQUAL(2, + m_device_def.autodiscovery_def.accel[0].iface.args[3].addr_space); + CHECK_EQUAL(1, + m_device_def.autodiscovery_def.accel[0].iface.args[3].category); + CHECK_EQUAL(4, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[3].size); + + // printf buffer size + CHECK_EQUAL(0, + m_device_def.autodiscovery_def.accel[0].iface.args[4].addr_space); + CHECK_EQUAL(0, + m_device_def.autodiscovery_def.accel[0].iface.args[4].category); + CHECK_EQUAL(4, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[4].size); + + CHECK_EQUAL(2, + m_device_def.autodiscovery_def.accel[0].iface.args[5].addr_space); + CHECK_EQUAL(1, + m_device_def.autodiscovery_def.accel[0].iface.args[5].category); + CHECK_EQUAL(4, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[5].size); + + CHECK_EQUAL( + 0, + (int)m_device_def.autodiscovery_def.accel[0].compile_work_group_size[0]); + CHECK_EQUAL( + 0, + (int)m_device_def.autodiscovery_def.accel[0].compile_work_group_size[1]); + CHECK_EQUAL( + 0, + (int)m_device_def.autodiscovery_def.accel[0].compile_work_group_size[2]); + + CHECK_EQUAL(0, + (int)m_device_def.autodiscovery_def.accel[0].max_work_group_size); + CHECK_EQUAL(1, (int)m_device_def.autodiscovery_def.accel[0].is_sycl_compile); + + // Checks for device global entry. + CHECK_EQUAL(2, m_device_def.autodiscovery_def.device_global_mem_defs.size()); + const auto kernel15_dev_global = + m_device_def.autodiscovery_def.device_global_mem_defs.find( + "kernel15_dev_global"); + const auto kernel15_dev_global2 = + m_device_def.autodiscovery_def.device_global_mem_defs.find( + "kernel15_dev_global2"); + CHECK(kernel15_dev_global != + m_device_def.autodiscovery_def.device_global_mem_defs.end()); + CHECK(kernel15_dev_global2 != + m_device_def.autodiscovery_def.device_global_mem_defs.end()); + CHECK_EQUAL(4096, kernel15_dev_global->second.address); + CHECK_EQUAL(2048, kernel15_dev_global->second.size); + CHECK_EQUAL(ACL_DEVICE_GLOBAL_HOST_ACCESS_NONE, + kernel15_dev_global->second.host_access); + CHECK_EQUAL(false, kernel15_dev_global->second.can_skip_programming); + CHECK_EQUAL(false, kernel15_dev_global->second.implement_in_csr); + CHECK_EQUAL(false, kernel15_dev_global->second.reset_on_reuse); + CHECK_EQUAL(2048, kernel15_dev_global2->second.address); + CHECK_EQUAL(1024, kernel15_dev_global2->second.size); + CHECK_EQUAL(ACL_DEVICE_GLOBAL_HOST_ACCESS_WRITE_ONLY, + kernel15_dev_global2->second.host_access); + CHECK_EQUAL(false, kernel15_dev_global2->second.can_skip_programming); + CHECK_EQUAL(true, kernel15_dev_global2->second.implement_in_csr); + CHECK_EQUAL(false, kernel15_dev_global2->second.reset_on_reuse); + + // Check a second parsing. + // It should allocate a new string for the name. + { + std::string autodiscovery2 = + load_fuzzed_value("auto_configure", "simple", "autodiscovery2"); + std::string err_str = std::string(err_str_c); + ACL_LOCKED(parsed = acl_load_device_def_from_str( + autodiscovery2, m_device_def.autodiscovery_def, err_str)); + } + + CHECK_EQUAL(1, parsed); + + CHECK(BOARDNAME2 == m_device_def.autodiscovery_def.name); + + CHECK_EQUAL( + 1, (int)m_device_def.autodiscovery_def.accel[0].is_workgroup_invariant); + CHECK_EQUAL( + 1, (int)m_device_def.autodiscovery_def.accel[0].is_workitem_invariant); + CHECK("bar" == m_device_def.autodiscovery_def.hal_info[0].name); + + CHECK_EQUAL(1, (int)m_device_def.autodiscovery_def.is_big_endian); + + CHECK_EQUAL( + 0, (int)m_device_def.autodiscovery_def.accel[0].local_aspaces.size()); + + CHECK_EQUAL( + 2, + (int)m_device_def.autodiscovery_def.accel[0].compile_work_group_size[0]); + CHECK_EQUAL( + 3, + (int)m_device_def.autodiscovery_def.accel[0].compile_work_group_size[1]); + CHECK_EQUAL( + 5, + (int)m_device_def.autodiscovery_def.accel[0].compile_work_group_size[2]); + + CHECK_EQUAL(1024, + (int)m_device_def.autodiscovery_def.accel[0].max_work_group_size); + CHECK_EQUAL( + 32, + (int)m_device_def.autodiscovery_def.accel[0].max_work_group_size_arr[0]); + CHECK_EQUAL( + 8, + (int)m_device_def.autodiscovery_def.accel[0].max_work_group_size_arr[1]); + CHECK_EQUAL( + 4, + (int)m_device_def.autodiscovery_def.accel[0].max_work_group_size_arr[2]); + + CHECK_EQUAL(0, + (int)m_device_def.autodiscovery_def.accel[0].max_global_work_dim); + CHECK_EQUAL(0, (int)m_device_def.autodiscovery_def.accel[0].is_sycl_compile); + + CHECK_EQUAL( + 5, + (int)m_device_def.autodiscovery_def.accel[0].profiling_words_to_readback); + + // Backwards-compatibility test (last backward compatible aoc version: 20.1, + // version id: 23) + { + std::string autodiscovery3 = + load_fuzzed_value("auto_configure", "simple", "autodiscovery3"); + std::string err_str = std::string(err_str_c); + ACL_LOCKED(parsed = acl_load_device_def_from_str( + autodiscovery3, m_device_def.autodiscovery_def, err_str)); + } + + CHECK_EQUAL(1, parsed); + + CHECK_EQUAL(1, m_device_def.autodiscovery_def.num_global_mem_systems); + CHECK_EQUAL(0, m_device_def.autodiscovery_def.global_mem_defs[0].range.begin); + CHECK_EQUAL((void *)2048, + m_device_def.autodiscovery_def.global_mem_defs[0].range.next); + + CHECK(BOARDNAME == m_device_def.autodiscovery_def.name); + + CHECK_EQUAL(0, (int)m_device_def.autodiscovery_def.is_big_endian); + + CHECK_EQUAL(1, (int)m_device_def.autodiscovery_def.accel.size()); + CHECK_EQUAL(1, (int)m_device_def.autodiscovery_def.hal_info.size()); + + CHECK_EQUAL(0, m_device_def.autodiscovery_def.accel[0].id); + CHECK_EQUAL(0, m_device_def.autodiscovery_def.accel[0].mem.begin); + CHECK_EQUAL( + (void *)0x020000, + m_device_def.autodiscovery_def.accel[0] + .mem.next); // Not sure why this isn't 16KB like OpenCL spec minimum + + CHECK_EQUAL( + 2, (int)m_device_def.autodiscovery_def.accel[0].local_aspaces.size()); + CHECK_EQUAL( + 5, + (int)m_device_def.autodiscovery_def.accel[0].local_aspaces[0].aspace_id); + CHECK_EQUAL(1024, (int)m_device_def.autodiscovery_def.accel[0] + .local_aspaces[0] + .static_demand); + CHECK_EQUAL( + 6, + (int)m_device_def.autodiscovery_def.accel[0].local_aspaces[1].aspace_id); + CHECK_EQUAL(2048, (int)m_device_def.autodiscovery_def.accel[0] + .local_aspaces[1] + .static_demand); + + CHECK("foo" == m_device_def.autodiscovery_def.accel[0].iface.name); + CHECK_EQUAL( + 0, (int)m_device_def.autodiscovery_def.accel[0].is_workgroup_invariant); + CHECK_EQUAL( + 0, (int)m_device_def.autodiscovery_def.accel[0].is_workitem_invariant); + CHECK_EQUAL( + 5, + (int)m_device_def.autodiscovery_def.accel[0].profiling_words_to_readback); + CHECK_EQUAL(7, + (int)m_device_def.autodiscovery_def.accel[0].iface.args.size()); + + CHECK_EQUAL(1, + m_device_def.autodiscovery_def.accel[0].iface.args[0].addr_space); + CHECK_EQUAL(1, + m_device_def.autodiscovery_def.accel[0].iface.args[0].category); + CHECK_EQUAL(4, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[0].size); + CHECK_EQUAL( + 5, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[0].aspace_number); + CHECK_EQUAL(16768, (int)m_device_def.autodiscovery_def.accel[0] + .iface.args[0] + .lmem_size_bytes); + + CHECK_EQUAL(2, + m_device_def.autodiscovery_def.accel[0].iface.args[1].addr_space); + CHECK_EQUAL(1, + m_device_def.autodiscovery_def.accel[0].iface.args[1].category); + CHECK_EQUAL(4, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[1].size); + + CHECK_EQUAL(0, + m_device_def.autodiscovery_def.accel[0].iface.args[2].addr_space); + CHECK_EQUAL(0, + m_device_def.autodiscovery_def.accel[0].iface.args[2].category); + CHECK_EQUAL(8, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[2].size); + + // printf buffer start address + CHECK_EQUAL(2, + m_device_def.autodiscovery_def.accel[0].iface.args[3].addr_space); + CHECK_EQUAL(1, + m_device_def.autodiscovery_def.accel[0].iface.args[3].category); + CHECK_EQUAL(4, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[3].size); + + // printf buffer size + CHECK_EQUAL(0, + m_device_def.autodiscovery_def.accel[0].iface.args[4].addr_space); + CHECK_EQUAL(0, + m_device_def.autodiscovery_def.accel[0].iface.args[4].category); + CHECK_EQUAL(4, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[4].size); + + CHECK_EQUAL(2, + m_device_def.autodiscovery_def.accel[0].iface.args[5].addr_space); + CHECK_EQUAL(1, + m_device_def.autodiscovery_def.accel[0].iface.args[5].category); + CHECK_EQUAL(4, + (int)m_device_def.autodiscovery_def.accel[0].iface.args[5].size); + + CHECK_EQUAL( + 0, + (int)m_device_def.autodiscovery_def.accel[0].compile_work_group_size[0]); + CHECK_EQUAL( + 0, + (int)m_device_def.autodiscovery_def.accel[0].compile_work_group_size[1]); + CHECK_EQUAL( + 0, + (int)m_device_def.autodiscovery_def.accel[0].compile_work_group_size[2]); + + CHECK_EQUAL(0, + (int)m_device_def.autodiscovery_def.accel[0].max_work_group_size); +} + +TEST(auto_configure, many_ok_forward_compatibility) { + bool check = true; + { + // From example_designs/merge_sort with extra fields at the end of each + // sections and subsections to check forward compatibility + + std::string str = load_fuzzed_value("auto_configure", + "many_ok_forward_compatibility", "str"); + std::vector device_defs(ACL_MAX_DEVICE); + for (auto &device_def : device_defs) { + int parsed; + std::string err_str; + ACL_LOCKED(parsed = acl_load_device_def_from_str( + str, device_def.autodiscovery_def, err_str)); + if (!check_condition(1 == parsed, check)) + break; + + if (!check_condition("a10gx" == device_def.autodiscovery_def.name, check)) + break; + if (!check_condition("sample40byterandomhash000000000000000000" == + device_def.autodiscovery_def.binary_rand_hash, + check)) + break; + + if (!check_condition(47 == (int)device_def.autodiscovery_def.accel.size(), + check)) + break; + if (!check_condition( + 47 == (int)device_def.autodiscovery_def.hal_info.size(), check)) + break; + } + } + CHECK(check); +} + +TEST(auto_configure, many_limit_check) { + bool check = true; + { + std::string str = + load_fuzzed_value("auto_configure", "many_limit_check", "str"); + // Verify that we can handle 75 kernels + std::vector device_defs(ACL_MAX_DEVICE); + for (auto &device_def : device_defs) { + int parsed; + std::string err; + ACL_LOCKED(parsed = acl_load_device_def_from_str( + str, device_def.autodiscovery_def, err)); + if (!check_condition(1 == parsed, check)) + break; + + // But we should write 75 kernels to the requested physical_device_id + // entry. + if (!check_condition(75 == (int)device_def.autodiscovery_def.accel.size(), + check)) + break; + if (!check_condition( + 75 == (int)device_def.autodiscovery_def.hal_info.size(), check)) + break; + } + } + CHECK(check); +} + +TEST(auto_configure, bad_config) { + bool check = true; + { + std::string str1 = + load_fuzzed_value("auto_configure", "bad_config", "str1"); + std::string str2 = + "Error: The accelerator hardware currently programmed is " + "incompatible with this\nversion of the runtime (" ACL_VERSION + " Commit " ACL_GIT_COMMIT + "). Recompile the hardware with\nthe same version of " + "the compiler and program that onto the board.\n"; + std::string str3 = + load_fuzzed_value("auto_configure", "bad_config", "str3"); + std::string str4 = + "FAILED to read auto-discovery string at byte 132: kernel cannot be " + "workitem-invariant while it is workgroup-variant. Full " + "auto-discovery string value is " VERSIONIDTOSTR( + ACL_AUTO_CONFIGURE_VERSIONID) " 15 " + "sample40byterandomhash000000000000" + "000000 a10gx 0 1 7 DDR 2 1 2 0 " + "2147483648 0 0 0 0 1 " + "31 external_sort_stage_0 0 128 1 " + "0 0 0 1 1 0 1 6 0 0 4 1 0 0 0 0 0 " + "0 1 1 1 3 1 1 1 3 1\n"; + std::vector strs = {str1, str2, str3, str4}; + + for (unsigned istr = 0; istr < strs.size(); istr += 2) { + int parsed; + std::string err; + ACL_LOCKED(parsed = acl_load_device_def_from_str( + strs[istr], m_device_def.autodiscovery_def, err)); + if (!check_condition(0 == parsed, check)) + break; + + const auto &expect = strs[istr + 1]; + for (unsigned ichar = 0; err[ichar] && expect[ichar] && + ichar < err.length() && ichar < expect.length(); + ichar++) { + if (err[ichar] != expect[ichar]) { + std::cout << "Failed at char " << ichar << " '" << err[ichar] + << "' vs '" << expect[ichar] << "'\n"; + } + } + + if (expect != err) { + std::cout << istr / 2 << ": err string is " << err << "\n"; + std::cout << istr / 2 << ": exp string is " << expect << "\n"; + } + + if (!check_condition(expect == err, check)) + break; + + if (!check_condition( + 0 == (int)m_device_def.autodiscovery_def.accel.size(), check)) + break; + if (!check_condition( + 0 == (int)m_device_def.autodiscovery_def.hal_info.size(), check)) + break; + } + } + CHECK(check); +} + +TEST(auto_configure, multi_mem_config) { + bool check = true; + { + std::string str = + load_fuzzed_value("auto_configure", "multi_mem_config", "str"); + + std::vector device_defs(ACL_MAX_DEVICE); + for (auto &device_def : device_defs) { + int parsed; + std::string err_str; + ACL_LOCKED(parsed = acl_load_device_def_from_str( + str, device_def.autodiscovery_def, err_str)); + + if (!check_condition(1 == parsed, check)) + break; + + if (!check_condition("pcie385n_a7" == device_def.autodiscovery_def.name, + check)) + break; + if (!check_condition("sample40byterandomhash000000000000000000" == + device_def.autodiscovery_def.binary_rand_hash, + check)) + break; + + if (!check_condition( + 4 == device_def.autodiscovery_def.num_global_mem_systems, check)) + break; + + if (!check_condition( + "SVM" == device_def.autodiscovery_def.global_mem_defs[0].name, + check)) + break; + if (!check_condition( + ACL_GLOBAL_MEM_SHARED_VIRTUAL == + device_def.autodiscovery_def.global_mem_defs[0].type, + check)) + break; + + if (!check_condition(1 == device_def.autodiscovery_def.global_mem_defs[0] + .burst_interleaved, + check)) + break; + if (!check_condition( + 0 == device_def.autodiscovery_def.global_mem_defs[0].config_addr, + check)) + break; + if (!check_condition(1 == device_def.autodiscovery_def.global_mem_defs[0] + .num_global_banks, + check)) + break; + if (!check_condition("" == device_def.autodiscovery_def.global_mem_defs[0] + .primary_interface, + check)) + break; + if (!check_condition(1 == device_def.autodiscovery_def.global_mem_defs[0] + .can_access_list.size(), + check)) + break; + if (!check_condition("SVM2" == + device_def.autodiscovery_def.global_mem_defs[0] + .can_access_list.at(0), + check)) + break; + + if (!check_condition( + "DDR" == device_def.autodiscovery_def.global_mem_defs[1].name, + check)) + break; + if (!check_condition( + ACL_GLOBAL_MEM_DEVICE_PRIVATE == + device_def.autodiscovery_def.global_mem_defs[1].type, + check)) + break; + if (!check_condition(1 == device_def.autodiscovery_def.global_mem_defs[1] + .burst_interleaved, + check)) + break; + if (!check_condition( + 0x18 == + device_def.autodiscovery_def.global_mem_defs[1].config_addr, + check)) + break; + if (!check_condition(2 == device_def.autodiscovery_def.global_mem_defs[1] + .num_global_banks, + check)) + break; + if (!check_condition(ACL_GLOBAL_MEM_UNDEFINED_ALLOCATION == + device_def.autodiscovery_def.global_mem_defs[1] + .allocation_type, + check)) + break; + + if (!check_condition( + "QDR" == device_def.autodiscovery_def.global_mem_defs[2].name, + check)) + break; + if (!check_condition( + ACL_GLOBAL_MEM_DEVICE_PRIVATE == + device_def.autodiscovery_def.global_mem_defs[2].type, + check)) + break; + if (!check_condition(0 == device_def.autodiscovery_def.global_mem_defs[2] + .burst_interleaved, + check)) + break; + if (!check_condition( + 0x30 == + device_def.autodiscovery_def.global_mem_defs[2].config_addr, + check)) + break; + if (!check_condition(4 == device_def.autodiscovery_def.global_mem_defs[2] + .num_global_banks, + check)) + break; + if (!check_condition(ACL_GLOBAL_MEM_DEVICE_ALLOCATION == + device_def.autodiscovery_def.global_mem_defs[2] + .allocation_type, + check)) + break; + + if (!check_condition( + "SVM2" == device_def.autodiscovery_def.global_mem_defs[3].name, + check)) + break; + if (!check_condition( + ACL_GLOBAL_MEM_SHARED_VIRTUAL == + device_def.autodiscovery_def.global_mem_defs[3].type, + check)) + break; + if (!check_condition(1 == device_def.autodiscovery_def.global_mem_defs[3] + .burst_interleaved, + check)) + break; + if (!check_condition( + 0 == device_def.autodiscovery_def.global_mem_defs[3].config_addr, + check)) + break; + if (!check_condition(1 == device_def.autodiscovery_def.global_mem_defs[3] + .num_global_banks, + check)) + break; + if (!check_condition("SVM" == + device_def.autodiscovery_def.global_mem_defs[3] + .primary_interface, + check)) + break; + if (!check_condition(0 == device_def.autodiscovery_def.global_mem_defs[3] + .can_access_list.size(), + check)) + break; + if (!check_condition(2 == (int)device_def.autodiscovery_def.accel.size(), + check)) + break; + if (!check_condition( + 2 == (int)device_def.autodiscovery_def.hal_info.size(), check)) + break; + } + } + CHECK(check); +} + +TEST(auto_configure, kernel_arg_info) { + bool check = true; + { + std::string str1 = + load_fuzzed_value("auto_configure", "kernel_arg_info", "str1"); + std::string str2 = + load_fuzzed_value("auto_configure", "kernel_arg_info", "str2"); + std::vector strs = {str1, str2}; + + // kernel arg info available + { + std::vector device_defs(ACL_MAX_DEVICE); + for (auto &device_def : device_defs) { + int parsed; + std::string err_str; + ACL_LOCKED(parsed = acl_load_device_def_from_str( + strs[0], device_def.autodiscovery_def, err_str)); + if (!check_condition(1 == parsed, check)) + break; + + if (!check_condition("a10gx" == device_def.autodiscovery_def.name, + check)) + break; + if (!check_condition("sample40byterandomhash000000000000000000" == + device_def.autodiscovery_def.binary_rand_hash, + check)) + break; + + if (!check_condition( + 1 == device_def.autodiscovery_def.num_global_mem_systems, + check)) + break; + if (!check_condition( + "DDR" == device_def.autodiscovery_def.global_mem_defs[0].name, + check)) + break; + + if (!check_condition( + 2 == (int)device_def.autodiscovery_def.accel.size(), check)) + break; + if (!check_condition( + 2 == (int)device_def.autodiscovery_def.hal_info.size(), check)) + break; + if (!check_condition(2 == (int)device_def.autodiscovery_def.accel[0] + .iface.args.size(), + check)) + break; + if (!check_condition(2 == (int)device_def.autodiscovery_def.accel[1] + .iface.args.size(), + check)) + break; + + if (!check_condition( + "arg_one" == + device_def.autodiscovery_def.accel[0].iface.args[0].name, + check)) + if (!check_condition( + check &= + "type_one" == + device_def.autodiscovery_def.accel[0].iface.args[0].type_name, + check)) + break; + if (!check_condition(1 == device_def.autodiscovery_def.accel[0] + .iface.args[0] + .access_qualifier, + check)) + break; + + if (!check_condition( + "arg_two" == + device_def.autodiscovery_def.accel[0].iface.args[1].name, + check)) + break; + if (!check_condition("type_two" == device_def.autodiscovery_def.accel[0] + .iface.args[1] + .type_name, + check)) + break; + if (!check_condition(2 == device_def.autodiscovery_def.accel[0] + .iface.args[1] + .access_qualifier, + check)) + break; + + if (!check_condition( + "arg_three" == + device_def.autodiscovery_def.accel[1].iface.args[0].name, + check)) + break; + if (!check_condition( + "arg_three" == + device_def.autodiscovery_def.accel[1].iface.args[0].name, + check)) + break; + if (!check_condition("type_three" == + device_def.autodiscovery_def.accel[1] + .iface.args[0] + .type_name, + check)) + break; + if (!check_condition(1 == device_def.autodiscovery_def.accel[1] + .iface.args[0] + .access_qualifier, + check)) + break; + + if (!check_condition( + "arg_four" == + device_def.autodiscovery_def.accel[1].iface.args[1].name, + check)) + break; + if (!check_condition("type_four" == + device_def.autodiscovery_def.accel[1] + .iface.args[1] + .type_name, + check)) + break; + if (!check_condition(2 == device_def.autodiscovery_def.accel[1] + .iface.args[1] + .access_qualifier, + check)) + break; + } + } + + // kernel arg info not available + { + std::vector device_defs(ACL_MAX_DEVICE); + for (auto &device_def : device_defs) { + int parsed; + std::string err_str; + ACL_LOCKED(parsed = acl_load_device_def_from_str( + strs[1], device_def.autodiscovery_def, err_str)); + if (!check_condition(1 == parsed, check)) + break; + + if (!check_condition("a10gx" == device_def.autodiscovery_def.name, + check)) + break; + if (!check_condition("sample40byterandomhash000000000000000000" == + device_def.autodiscovery_def.binary_rand_hash, + check)) + break; + + if (!check_condition( + 1 == device_def.autodiscovery_def.num_global_mem_systems, + check)) + break; + if (!check_condition( + "DDR" == device_def.autodiscovery_def.global_mem_defs[0].name, + check)) + break; + + if (!check_condition( + 2 == (int)device_def.autodiscovery_def.accel.size(), check)) + break; + if (!check_condition( + 2 == (int)device_def.autodiscovery_def.hal_info.size(), check)) + break; + if (!check_condition(2 == (int)device_def.autodiscovery_def.accel[0] + .iface.args.size(), + check)) + break; + if (!check_condition(2 == (int)device_def.autodiscovery_def.accel[1] + .iface.args.size(), + check)) + break; + + if (!check_condition( + "" == device_def.autodiscovery_def.accel[0].iface.args[0].name, + check)) + break; + if (!check_condition("" == device_def.autodiscovery_def.accel[0] + .iface.args[0] + .type_name, + check)) + break; + if (!check_condition(0 == device_def.autodiscovery_def.accel[0] + .iface.args[0] + .access_qualifier, + check)) + break; + + if (!check_condition( + "" == device_def.autodiscovery_def.accel[0].iface.args[1].name, + check)) + break; + if (!check_condition("" == device_def.autodiscovery_def.accel[0] + .iface.args[1] + .type_name, + check)) + break; + if (!check_condition(0 == device_def.autodiscovery_def.accel[0] + .iface.args[1] + .access_qualifier, + check)) + break; + + if (!check_condition( + "" == device_def.autodiscovery_def.accel[1].iface.args[0].name, + check)) + break; + if (!check_condition("" == device_def.autodiscovery_def.accel[1] + .iface.args[0] + .type_name, + check)) + break; + if (!check_condition(0 == device_def.autodiscovery_def.accel[1] + .iface.args[0] + .access_qualifier, + check)) + break; + + if (!check_condition( + "" == device_def.autodiscovery_def.accel[1].iface.args[1].name, + check)) + break; + if (!check_condition("" == device_def.autodiscovery_def.accel[1] + .iface.args[1] + .type_name, + check)) + break; + if (!check_condition(0 == device_def.autodiscovery_def.accel[1] + .iface.args[1] + .access_qualifier, + check)) + break; + } + } + } + CHECK(check); +} + +TEST(auto_configure, hostpipe_basic) { + bool check = true; + { + std::string str( + load_fuzzed_value("auto_configure", "hostpipe_basic", "str")); + + std::vector device_defs(ACL_MAX_DEVICE); + for (auto &device_def : device_defs) { + int parsed; + std::string err_str; + ACL_LOCKED(parsed = acl_load_device_def_from_str( + str, device_def.autodiscovery_def, err_str)); + + if (!check_condition(check &= 1 == parsed, check)) + break; + if (!check_condition(check &= "a10gx_hostpipe" == + device_def.autodiscovery_def.name, + check)) + break; + if (!check_condition(check &= + "sample40byterandomhash000000000000000000" == + device_def.autodiscovery_def.binary_rand_hash, + check)) + break; + + if (!check_condition( + check &= 1 == device_def.autodiscovery_def.num_global_mem_systems, + check)) + break; + if (!check_condition( + check &= "foo" == device_def.autodiscovery_def.hal_info[0].name, + check)) + break; + if (!check_condition(check &= + "DDR" == + device_def.autodiscovery_def.global_mem_defs[0].name, + check)) + break; + + if (!check_condition(check &= + 1 == (int)device_def.autodiscovery_def.accel.size(), + check)) + break; + if (!check_condition( + check &= 1 == (int)device_def.autodiscovery_def.hal_info.size(), + check)) + break; + } + } + CHECK(check); +} + +TEST(auto_configure, streaming_basic) { + bool check = 1; + // Dummy for loop + for (int i = 0; i < 1; i++) { + std::string config_str = + load_fuzzed_value("auto_configure", "streaming_basic", "config_str"); + acl_device_def_autodiscovery_t devdef; + + bool result; + std::string err_str; + ACL_LOCKED(result = + acl_load_device_def_from_str(config_str, devdef, err_str)); + std::cerr << err_str; + if (!check_condition(check &= result, check)) + break; + + if (!check_condition(check &= 1 == devdef.accel.size(), check)) + break; + + if (!check_condition(check &= !devdef.accel[0].is_sycl_compile, check)) + break; + if (!check_condition( + check &= devdef.accel[0].streaming_control_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS3CRCILi0EE_streaming_start" == + devdef.accel[0].streaming_control_info.start, + check)) + break; + if (!check_condition(check &= "k0_ZTS3CRCILi0EE_streaming_done" == + devdef.accel[0].streaming_control_info.done, + check)) + break; + + const auto &args = devdef.accel[0].iface.args; + if (!check_condition(check &= 9 == args.size(), check)) + break; + + if (!check_condition(check &= args[0].streaming_arg_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS3CRCILi0EE_arg0" == + args[0].streaming_arg_info.interface_name, + check)) + break; + + if (!check_condition(check &= args[1].streaming_arg_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS3CRCILi0EE_arg1" == + args[1].streaming_arg_info.interface_name, + check)) + break; + + if (!check_condition(check &= args[2].streaming_arg_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS3CRCILi0EE_arg2" == + args[2].streaming_arg_info.interface_name, + check)) + break; + + for (size_t i = 3; i < args.size(); ++i) { + if (!check_condition(check &= !args[i].streaming_arg_info_available, + check)) + break; + } + } + CHECK(check); +} + +TEST(auto_configure, one_streaming_arg_and_streaming_kernel) { + bool check = 1; + // Dummy for loop + for (int i = 0; i < 1; i++) { + std::string config_str = load_fuzzed_value( + "auto_configure", "one_streaming_arg_and_streaming_kernel", + "config_str"); + acl_device_def_autodiscovery_t devdef; + + bool result; + std::string err_str; + ACL_LOCKED(result = + acl_load_device_def_from_str(config_str, devdef, err_str)); + std::cerr << err_str; + if (!check_condition(check &= result, check)) + break; + + if (!check_condition(check &= 1 == devdef.accel.size(), check)) + break; + + if (!check_condition( + check &= devdef.accel[0].streaming_control_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS15binomial_kernel_streaming_start" == + devdef.accel[0].streaming_control_info.start, + check)) + break; + if (!check_condition(check &= "k0_ZTS15binomial_kernel_streaming_done" == + devdef.accel[0].streaming_control_info.done, + check)) + break; + + const auto &args = devdef.accel[0].iface.args; + if (!check_condition(check &= 8 == args.size(), check)) + break; + + if (!check_condition(check &= !args[0].streaming_arg_info_available, check)) + break; + + if (!check_condition(check &= args[1].streaming_arg_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS15binomial_kernel_arg1" == + args[1].streaming_arg_info.interface_name, + check)) + break; + + for (size_t i = 2; i < args.size(); ++i) { + if (!check_condition(check &= !args[i].streaming_arg_info_available, + check)) + break; + } + } + + CHECK(check); +} + +TEST(auto_configure, two_streaming_args_and_streaming_kernel) { + bool check = true; + // Dummy for loop + for (int i = 0; i < 1; i++) { + std::string config_str = load_fuzzed_value( + "auto_configure", "two_streaming_args_and_streaming_kernel", + "config_str"); + acl_device_def_autodiscovery_t devdef; + + bool result; + std::string err_str; + ACL_LOCKED(result = + acl_load_device_def_from_str(config_str, devdef, err_str)); + std::cerr << err_str; + if (!check_condition(check &= result, check)) + break; + + if (!check_condition(check &= 1 == devdef.accel.size(), check)) + break; + + if (!check_condition(check &= devdef.accel[0].is_sycl_compile, check)) + break; + if (!check_condition( + check &= devdef.accel[0].streaming_control_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS15binomial_kernel_streaming_start" == + devdef.accel[0].streaming_control_info.start, + check)) + break; + if (!check_condition(check &= "k0_ZTS15binomial_kernel_streaming_done" == + devdef.accel[0].streaming_control_info.done, + check)) + break; + + const auto &args = devdef.accel[0].iface.args; + if (!check_condition(check &= 8 == args.size(), check)) + break; + + if (!check_condition(check &= args[0].streaming_arg_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS15binomial_kernel_arg0" == + args[0].streaming_arg_info.interface_name, + check)) + break; + + if (!check_condition(check &= args[1].streaming_arg_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS15binomial_kernel_arg1" == + args[1].streaming_arg_info.interface_name, + check)) + break; + + for (size_t i = 2; i < args.size(); ++i) { + if (!check_condition(check &= !args[i].streaming_arg_info_available, + check)) + break; + } + } + CHECK(check); +} + +TEST(auto_configure, two_streaming_args_and_non_streaming_kernel) { + bool check = 1; + // Dummy for loop + for (int i = 0; i < 1; i++) { + std::string config_str = load_fuzzed_value( + "auto_configure", "two_streaming_args_and_non_streaming_kernel", + "config_str"); + acl_device_def_autodiscovery_t devdef; + + bool result; + std::string err_str; + ACL_LOCKED(result = + acl_load_device_def_from_str(config_str, devdef, err_str)); + std::cerr << err_str; + if (!check_condition(check &= result, check)) + break; + + if (!check_condition(check &= 1 == devdef.accel.size(), check)) + break; + + if (!check_condition(check &= devdef.accel[0].is_sycl_compile, check)) + break; + if (!check_condition( + check &= !devdef.accel[0].streaming_control_info_available, check)) + break; + + const auto &args = devdef.accel[0].iface.args; + if (!check_condition(check &= 8 == args.size(), check)) + break; + + if (!check_condition(check &= args[0].streaming_arg_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS15binomial_kernel_arg0" == + args[0].streaming_arg_info.interface_name, + check)) + break; + + if (!check_condition(check &= args[1].streaming_arg_info_available, check)) + break; + if (!check_condition(check &= "k0_ZTS15binomial_kernel_arg1" == + args[1].streaming_arg_info.interface_name, + check)) + break; + + for (size_t i = 2; i < args.size(); ++i) { + if (!check_condition(check &= !args[i].streaming_arg_info_available, + check)) + break; + } + } + CHECK(check); +} + +TEST(auto_configure, cra_ring_root_not_exist) { + bool check = true; + // Dummy for loop + for (int i = 0; i < 1; i++) { + std::string config_str = load_fuzzed_value( + "auto_configure", "cra_ring_root_not_exist", "config_str"); + acl_device_def_autodiscovery_t devdef; + + bool result; + std::string err_str; + ACL_LOCKED(result = acl_load_device_def_from_str(std::string(config_str), + devdef, err_str)); + std::cerr << err_str; + if (!check_condition(check &= result, check)) + break; + if (!check_condition(check &= 0 == devdef.cra_ring_root_exist, check)) + break; + } + CHECK(check); +} + +TEST(auto_configure, cra_ring_root_exist) { + bool check = true; + // Dummy for loop + for (int i = 0; i < 1; i++) { + std::string config_str = load_fuzzed_value( + "auto_configure", "cra_ring_root_exist", "config_str"); + acl_device_def_autodiscovery_t devdef; + + bool result; + std::string err_str; + ACL_LOCKED(result = + acl_load_device_def_from_str(config_str, devdef, err_str)); + std::cerr << err_str; + if (!check_condition(check &= result, check)) + break; + if (!check_condition(check &= 1 == devdef.cra_ring_root_exist, check)) + break; + } + CHECK(check); +} + +TEST(auto_configure, hostpipe_mappings) { + bool check = true; + // Dummy for loop + for (int i = 0; i < 1; i++) { + std::string config_str = + load_fuzzed_value("auto_configure", "hostpipe_mappings", "config_str"); + acl_device_def_autodiscovery_t devdef; + bool result; + std::string err_str; + ACL_LOCKED(result = + acl_load_device_def_from_str(config_str, devdef, err_str)); + std::cerr << err_str; + if (!check_condition(check &= result, check)) + break; + + if (!check_condition(check &= 5 == devdef.hostpipe_mappings.size(), check)) + break; + + if (!check_condition(check &= devdef.hostpipe_mappings[0].logical_name == + "pipe_logical_name1", + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[0].physical_name == + "pipe_physical_name1", + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[0].implement_in_csr, + check)) + break; + if (!check_condition( + check &= devdef.hostpipe_mappings[0].csr_address == "12345", check)) + break; + if (!check_condition(check &= !devdef.hostpipe_mappings[0].is_read, check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[0].is_write, check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[0].pipe_width == 4, + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[0].pipe_depth == 10, + check)) + break; + + if (!check_condition(check &= devdef.hostpipe_mappings[1].logical_name == + "pipe_logical_name2", + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[1].physical_name == + "pipe_physical_name2", + check)) + break; + if (!check_condition(check &= !devdef.hostpipe_mappings[1].implement_in_csr, + check)) + break; + if (!check_condition( + check &= devdef.hostpipe_mappings[1].csr_address == "12323", check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[1].is_read, check)) + break; + if (!check_condition(check &= !devdef.hostpipe_mappings[1].is_write, check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[1].pipe_width == 8, + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[1].pipe_depth == 20, + check)) + break; + + if (!check_condition(check &= devdef.hostpipe_mappings[2].logical_name == + "pipe_logical_name3", + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[2].physical_name == + "pipe_physical_name1", + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[2].implement_in_csr, + check)) + break; + if (!check_condition( + check &= devdef.hostpipe_mappings[2].csr_address == "12313", check)) + break; + if (!check_condition(check &= !devdef.hostpipe_mappings[2].is_read, check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[2].is_write, check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[2].pipe_width == 4, + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[2].pipe_depth == 10, + check)) + break; + + if (!check_condition(check &= devdef.hostpipe_mappings[3].logical_name == + "pipe_logical_name5", + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[3].physical_name == + "pipe_physical_name1", + check)) + break; + if (!check_condition(check &= !devdef.hostpipe_mappings[3].implement_in_csr, + check)) + break; + if (!check_condition( + check &= devdef.hostpipe_mappings[3].csr_address == "12316", check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[3].is_read, check)) + break; + if (!check_condition(check &= !devdef.hostpipe_mappings[3].is_write, check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[3].pipe_width == 8, + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[3].pipe_depth == 20, + check)) + break; + + if (!check_condition(check &= devdef.hostpipe_mappings[4].logical_name == + "pipe_logical_name4", + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[4].physical_name == + "pipe_physical_name3", + check)) + break; + if (!check_condition(check &= !devdef.hostpipe_mappings[4].implement_in_csr, + check)) + break; + if (!check_condition( + check &= devdef.hostpipe_mappings[4].csr_address == "12342", check)) + break; + if (!check_condition(check &= !devdef.hostpipe_mappings[4].is_read, check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[4].is_write, check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[4].pipe_width == 4, + check)) + break; + if (!check_condition(check &= devdef.hostpipe_mappings[4].pipe_depth == 10, + check)) + break; + } + CHECK(check); +} diff --git a/fuzz_testing/test/acl_fuzz_test.cpp b/fuzz_testing/test/acl_fuzz_test.cpp new file mode 100644 index 00000000..6652eec8 --- /dev/null +++ b/fuzz_testing/test/acl_fuzz_test.cpp @@ -0,0 +1,745 @@ +// Copyright (C) 2010-2021 Intel Corporation +// SPDX-License-Identifier: BSD-3-Clause + +#ifdef _MSC_VER +#pragma warning(push) +#pragma warning(disable : 4100) // unreferenced formal parameter +#pragma warning(disable : 4266) // no override available for virtual member + // function from base +#endif +#include +#include +#include +#ifdef _MSC_VER +#pragma warning(pop) +#endif + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "acl_fuzz_test.h" +#include "acl_globals_fuzz_test.h" +#include "acl_hal_fuzz_test.h" + +#ifdef _WIN32 +#include +#include +#else +#include +#endif + +// An example binary we cache for reuse many times, with offline-capture mode. +// It's built for the default board naned by ACLTEST_DEFAULT_BOARD +static unsigned char *acl_test_example_binary = 0; +static size_t acl_test_example_binary_len = 0; +// The corresponding sysdef. Used by some tests. +static acl_system_def_t acl_test_example_binary_sysdef{}; + +static void l_load_example_binary(); +static void l_run_benchmark(); + +int main(int argc, const char **argv) { + for (int i = 0; i < argc; ++i) { + if (std::string(argv[i]) == std::string("--benchmark")) { + l_run_benchmark(); + return 0; + } + } + + acl_test_unsetenv("CL_CONTEXT_OFFLINE_DEVICE_INTELFPGA"); + acl_test_unsetenv("CL_CONTEXT_COMPILER_MODE_INTELFPGA"); + printf("sizeof platform %llu\n", (unsigned long long)sizeof(acl_platform)); + printf("sizeof context %llu\n", + (unsigned long long)sizeof(struct _cl_context)); + printf("sizeof command queue %llu. Initially allocated %d queues\n", + (unsigned long long)sizeof(struct _cl_command_queue), + ACL_INIT_COMMAND_QUEUE_ALLOC); + printf("sizeof event %llu\n", (unsigned long long)sizeof(struct _cl_event)); + printf("sizeof cl_mem %llu\n", (unsigned long long)sizeof(struct _cl_mem)); + printf("sizeof kernel %llu\n", (unsigned long long)sizeof(struct _cl_kernel)); + printf("sizeof device_op_queue %llu\n", + (unsigned long long)sizeof(struct acl_device_op_queue_t)); + + if (getenv("ACL_SKIP_BB")) { // for faster test turnaround in some cases + printf("Skipping building binary\n"); + } else { + l_load_example_binary(); + } + + return CommandLineTestRunner::RunAllTests(argc, argv); +} + +void acl_test_setup_generic_system() { + acl_mutex_wrapper.lock(); + assert(1 == acl_set_hal(acl_test_get_simple_hal())); + assert(1 == acl_init(acl_test_get_complex_system_def())); + acl_mutex_wrapper.unlock(); +} + +void acl_test_setup_empty_system() { + acl_mutex_wrapper.lock(); + assert(1 == acl_set_hal(acl_test_get_simple_hal())); + assert(1 == acl_init(acl_test_get_empty_system_def())); + acl_mutex_wrapper.unlock(); +} + +void acl_test_setup_sample_default_board_system(void) { + acl_mutex_wrapper.lock(); + assert(1 == acl_set_hal(acl_test_get_simple_hal())); + assert(1 == acl_init(&acl_test_example_binary_sysdef)); + acl_mutex_wrapper.unlock(); +} + +void acl_test_teardown_sample_default_board_system(void) { + acl_test_teardown_system(); +} + +void acl_test_teardown_generic_system(void) { acl_test_teardown_system(); } +void acl_test_teardown_system(void) { + acl_mutex_wrapper.lock(); + acl_reset(); + acl_reset_hal(); + acltest_hal_teardown(); + acl_mutex_wrapper.unlock(); +} + +void acl_hal_test_setup_generic_system(void) { return; }; + +void acl_hal_test_teardown_generic_system(void) { return; }; + +void acl_test_run_standard_teardown_checks() { + CHECK(!acl_is_locked()); + if (acl_get_num_alloc_cl_mem() != 0) + printf("num aclloc cl_mem= %d\n", acl_get_num_alloc_cl_mem()); + CHECK_EQUAL(0, acl_get_num_alloc_cl_mem()); + CHECK_EQUAL(0, acl_get_num_alloc_cl_program()); + CHECK_EQUAL(0, acl_get_num_alloc_cl_context()); + acl_set_allow_invalid_type(0); + acl_set_allow_invalid_type(0); + acl_set_allow_invalid_type(0); + acl_set_allow_invalid_type(0); +} + +TEST_GROUP(Min){}; + +TEST(Min, basic) { + CHECK_EQUAL(-50, MIN(-50, 10)); + CHECK_EQUAL(10, MIN(10, 50)); + CHECK_EQUAL(19, MIN(19, 19)); +} + +#ifdef _WIN32 +#define snprintf sprintf_s +#endif + +SimpleString StringFrom(cl_uint x) { + char buf[30]; // probably 12 will do + snprintf(&buf[0], sizeof(buf) / sizeof(buf[0]), "%u", x); + const char *start_of_buf = &buf[0]; + return StringFrom(start_of_buf); +} +#if ACL_TARGET_BIT == 32 +SimpleString StringFrom(cl_ulong x) { + char buf[30]; // probably 12 will do + snprintf(&buf[0], sizeof(buf) / sizeof(buf[0]), "%lu", x); + const char *start_of_buf = &buf[0]; + return StringFrom(start_of_buf); +} +#endif +#ifdef _WIN64 +SimpleString StringFrom(intptr_t x) { + char buf[30]; // probably 12 will do + snprintf(&buf[0], sizeof(buf) / sizeof(buf[0]), "%Iu", x); + const char *start_of_buf = &buf[0]; + return StringFrom(start_of_buf); +} +#endif +// If ACL_TARGET_BIT is 32, then size_t == cl_ulong == cl_uint, and we've +// already got a body for that. +#if ACL_TARGET_BIT > 32 +SimpleString StringFrom(size_t x) { + char buf[30]; // probably 12 will do + snprintf(&buf[0], sizeof(buf) / sizeof(buf[0]), "%zd", + x); // format string might be platform dependent..? + const char *start_of_buf = &buf[0]; + return StringFrom(start_of_buf); +} +#endif + +void acl_test_unsetenv(const char *var) { +#ifdef _WIN32 + _putenv_s(var, ""); +#else + unsetenv(var); +#endif +} + +void acl_test_setenv(const char *var, const char *value) { +#ifdef _WIN32 + _putenv_s(var, value); +#else + setenv(var, value, 1); +#endif +} + +void CL_CALLBACK acl_test_notify_print(const char *errinfo, + const void *private_info, size_t cb, + void *user_data) { + printf("Context error: %s\n", errinfo); + cb = cb; // avoid warning on windows + private_info = private_info; // avoid warning on windows + user_data = user_data; // avoid warning on windows +} + +const unsigned char *acl_test_get_example_binary(size_t *binary_len) { + *binary_len = acl_test_example_binary_len; + return acl_test_example_binary; +} + +static void l_load_example_binary(void) { + const char *envvar_offline_device = "CL_CONTEXT_OFFLINE_DEVICE_INTELFPGA"; + const char *envvar_program_lib = + "CL_CONTEXT_PROGRAM_EXE_LIBRARY_ROOT_INTELFPGA"; + const char *offline_old_value = acl_getenv(envvar_offline_device); + const char *program_lib_old_value = acl_getenv(envvar_program_lib); + int system_ret = -1; + enum { MAX_DEVICES = 100 }; + cl_platform_id platform; + cl_device_id device[MAX_DEVICES]; + cl_context context; + cl_program program; + cl_int status; + + acl_test_setenv(envvar_offline_device, ACLTEST_DEFAULT_BOARD); + acl_test_setenv(envvar_program_lib, ".acltest_builtin_prog"); + system_ret = system("rm -rf .acltest_builtin_prog"); + assert(system_ret != -1); + + ACL_LOCKED(acl_test_setup_generic_system()); + + // Since this runs before the CppUTest runner is set up, we can't use + // the CHECK* macros. + // Just use asserts. + + assert(CL_SUCCESS == clGetPlatformIDs(1, &platform, 0)); + assert(CL_SUCCESS == clGetDeviceIDs(platform, CL_DEVICE_TYPE_ACCELERATOR, + MAX_DEVICES, device, 0)); + + cl_context_properties props[] = { + CL_CONTEXT_COMPILER_MODE_INTELFPGA, + CL_CONTEXT_COMPILER_MODE_OFFLINE_CREATE_EXE_LIBRARY_INTELFPGA, 0}; + context = clCreateContext(props, 1, device, acl_test_notify_print, 0, 0); + assert(context); + + const char *src = + "kernel void vecaccum(global int*A, global int*B) {\n" + " size_t gid = get_global_id(0);\n" + " A[gid] += B[gid];\n" + "};\n" + // This one has two constant arguments. + "kernel void vecsum(global int*A, constant int*B, constant int*C) {\n" + " size_t gid = get_global_id(0);\n" + " A[gid] = B[gid] + C[gid];\n" + "};\n" + + // This has a printf. + "kernel void printit(global int*A) {\n" + " printf(\"Hello world! %d\\n\", A[0]);\n" + "};\n"; + + program = clCreateProgramWithSource(context, 1, &src, 0, 0); + assert(program); + + status = clBuildProgram(program, 1, device, "-cl-kernel-arg-info", 0, 0); + if (status != CL_SUCCESS) { + printf("Compilation failed. Kernel source is:\n-----\n%s\n----\n", src); + size_t log_size = 0; + clGetProgramBuildInfo(program, device[0], CL_PROGRAM_BUILD_LOG, 0, 0, + &log_size); + char *log = (char *)acl_malloc(log_size); + clGetProgramBuildInfo(program, device[0], CL_PROGRAM_BUILD_LOG, log_size, + log, 0); + if (log) + printf("Build log is:\n-----\n%s\n----\n", log); + exit(1); + } + + // The build log should not be empty + size_t log_size = 0; + size_t empty_log_size = 1; + clGetProgramBuildInfo(program, device[0], CL_PROGRAM_BUILD_LOG, 0, 0, + &log_size); + assert(log_size > empty_log_size); + + acl_test_example_binary_len = 0; + assert(CL_SUCCESS == clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, + sizeof(size_t), + &acl_test_example_binary_len, 0)); + acl_test_example_binary = + (unsigned char *)acl_malloc(acl_test_example_binary_len); + assert(acl_test_example_binary); + assert(CL_SUCCESS == clGetProgramInfo(program, CL_PROGRAM_BINARIES, + sizeof(acl_test_example_binary), + &acl_test_example_binary, 0)); + + // Save the derived sysdef for later tests. + { + acl_pkg_file_t pkg; + size_t autodiscovery_len = 0; + static char autodiscovery[4097]; + pkg = acl_pkg_open_file_from_memory((char *)acl_test_example_binary, + acl_test_example_binary_len, 0); + assert(pkg); + + assert(acl_pkg_section_exists(pkg, ACL_PKG_SECTION_AUTODISCOVERY, + &autodiscovery_len)); + assert(autodiscovery_len < sizeof(autodiscovery)); // won't overflow. + assert(acl_pkg_read_section(pkg, ACL_PKG_SECTION_AUTODISCOVERY, + autodiscovery, autodiscovery_len + 1)); + + // Now parse it. + static std::string errstr; + auto result = false; + ACL_LOCKED(result = acl_load_device_def_from_str( + std::string(autodiscovery), + acl_test_example_binary_sysdef.device[0] + .autodiscovery_def, // populating this + errstr)); + assert(result); + ACL_LOCKED(acl_test_example_binary_sysdef.num_devices = 1); + + assert(ACLTEST_DEFAULT_BOARD == + acl_test_example_binary_sysdef.device[0].autodiscovery_def.name); + acl_pkg_close_file(pkg); + } + + // Don't leak + clReleaseProgram(program); + clReleaseContext(context); + + acl_test_unsetenv(envvar_offline_device); + if (offline_old_value) { + acl_test_setenv(envvar_offline_device, offline_old_value); + } + acl_test_unsetenv(envvar_program_lib); + if (program_lib_old_value) { + acl_test_setenv(envvar_program_lib, program_lib_old_value); + } + + ACL_LOCKED(acl_test_teardown_generic_system()); +} + +// Return a context properties array that specifies preloaded binary only. +// This was the default for all releases prior to 13.0, but is not +// conformant. +cl_context_properties *acl_test_context_prop_preloaded_binary_only(void) { + static cl_context_properties props[] = { + CL_CONTEXT_COMPILER_MODE_INTELFPGA, + (cl_context_properties) + CL_CONTEXT_COMPILER_MODE_PRELOADED_BINARY_ONLY_INTELFPGA, + 0}; + return &(props[0]); +} + +TEST_GROUP(envsets){void setup(){} void teardown(){}}; + +TEST(envsets, test) { + const char *foo = "ACLFOO"; + acl_test_setenv("ACLFOO", "abc"); + const char *fooenv = acl_getenv(foo); + CHECK(fooenv); + CHECK_EQUAL(0, strncmp("abc", fooenv, MAX_NAME_SIZE)); + acl_test_unsetenv("ACLFOO"); + CHECK_EQUAL(0, acl_getenv(foo)); +} + +// --- benchmark --------------------------------------------------------------- + +#ifdef _WIN32 +static LONGLONG l_ticks_per_second = 0; +#endif + +static inline cl_ulong l_get_timestamp() { +#ifdef __linux__ + struct timespec time; +#ifdef CLOCK_MONOTONIC_RAW + int ret = clock_gettime(CLOCK_MONOTONIC_RAW, &time); +#else + int ret = clock_gettime(CLOCK_MONOTONIC, &time); +#endif + + assert(ret == 0); + return time.tv_sec * 1000 * 1000 * 1000 + time.tv_nsec; +#else + LARGE_INTEGER li; + double seconds; + INT64 ticks; + + const INT64 NS_PER_S = 1000000000; + + QueryPerformanceCounter(&li); + ticks = li.QuadPart; + seconds = ticks / (double)l_ticks_per_second; + return (cl_ulong)((double)seconds * (double)NS_PER_S + 0.5); +#endif +} + +static inline void l_sleep(int milliseconds) { +#ifdef _WIN32 + Sleep((DWORD)milliseconds); +#else + struct timespec delay; + delay.tv_sec = milliseconds / 1000; + delay.tv_nsec = (milliseconds % 1000) * 1000 * 1000; + nanosleep(&delay, NULL); +#endif +} + +static void l_generic_context_callback(const char *errinfo, + const void *private_info, size_t cb, + void *user_data) { + UNREFERENCED_PARAMETER(private_info); + UNREFERENCED_PARAMETER(cb); + UNREFERENCED_PARAMETER(user_data); + std::cout << "Error from context callback: " << errinfo << std::endl; +} + +static void l_run_benchmark() { + std::cout << "Starting benchmark..." << std::endl; + + if (debug_mode > 0) { + std::cout << "WARNING! You are running this benchmark in debug mode!" + << std::endl; + } + +#ifdef _WIN32 + LARGE_INTEGER li; + QueryPerformanceFrequency(&li); + l_ticks_per_second = li.QuadPart; + assert(l_ticks_per_second != 0); +#endif + + acl_test_setup_generic_system(); + + const int INNER_REPS = 100000; + const int OUTER_REPS = 5; + + typedef std::deque times_t; + typedef std::map results_t; + + results_t results; + + for (int outer_rep = 0; outer_rep < OUTER_REPS; ++outer_rep) { + std::cout << "Iteration " << (outer_rep + 1) << std::endl; + + cl_int status; + cl_platform_id platform; + cl_device_id device; + cl_context context; + cl_program program; + cl_command_queue cq; + cl_mem mem; + cl_kernel kernel; + + cl_ulong start_time, end_time; + times_t *times; + times_t *create_times; + times_t *release_times; + + std::cout << "Measuring acl_lock/acl_unlock..." << std::endl; + times = &results["acl_lock/acl_unlock"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + start_time = l_get_timestamp(); + acl_mutex_wrapper.lock(); + acl_mutex_wrapper.unlock(); + end_time = l_get_timestamp(); + times->push_back(end_time - start_time); + } + + std::cout << "Measuring acl_assert_locked..." << std::endl; + times = &results["acl_assert_locked"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + acl_mutex_wrapper.lock(); + start_time = l_get_timestamp(); + acl_assert_locked(); + end_time = l_get_timestamp(); + acl_mutex_wrapper.unlock(); + times->push_back(end_time - start_time); + } + + std::cout << "Measuring clGetPlatformIDs..." << std::endl; + times = &results["clGetPlatformIDs"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + start_time = l_get_timestamp(); + status = clGetPlatformIDs(1, &platform, 0); + end_time = l_get_timestamp(); + times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + } + + std::cout << "Measuring clGetDeviceIDs..." << std::endl; + times = &results["clGetDeviceIDs"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + start_time = l_get_timestamp(); + status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, NULL); + end_time = l_get_timestamp(); + times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + } + + std::cout << "Measuring clCreateContext/clReleaseContext..." << std::endl; + create_times = &results["clCreateContext"]; + release_times = &results["clReleaseContext"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + start_time = l_get_timestamp(); + context = + clCreateContext(acl_test_context_prop_preloaded_binary_only(), 1, + &device, l_generic_context_callback, 0, &status); + end_time = l_get_timestamp(); + create_times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + assert(context); + + start_time = l_get_timestamp(); + status = clReleaseContext(context); + end_time = l_get_timestamp(); + release_times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + } + + context = clCreateContext(acl_test_context_prop_preloaded_binary_only(), 1, + &device, 0, 0, &status); + assert(status == CL_SUCCESS); + assert(context); + + const unsigned char *binary = (const unsigned char *)"0"; + size_t binary_length = 1; + program = clCreateProgramWithBinary(context, 1, &device, &binary_length, + &binary, NULL, &status); + assert(status == CL_SUCCESS); + + status = clBuildProgram(program, 1, &device, "", NULL, NULL); + assert(status == CL_SUCCESS); + + std::cout << "Measuring clCreateCommandQueue/clReleaseCommandQueue..." + << std::endl; + create_times = &results["clCreateCommandQueue"]; + release_times = &results["clReleaseCommandQueue"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + start_time = l_get_timestamp(); + cq = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, + &status); + end_time = l_get_timestamp(); + create_times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + assert(cq); + + start_time = l_get_timestamp(); + status = clReleaseCommandQueue(cq); + end_time = l_get_timestamp(); + release_times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + } + + cq = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, + &status); + assert(status == CL_SUCCESS); + assert(cq); + + std::cout << "Measuring clCreateBuffer/clReleaseMemObject..." << std::endl; + create_times = &results["clCreateBuffer"]; + release_times = &results["clReleaseMemObject"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + start_time = l_get_timestamp(); + mem = clCreateBuffer(context, CL_MEM_READ_ONLY, 64, 0, &status); + end_time = l_get_timestamp(); + create_times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + assert(mem); + + start_time = l_get_timestamp(); + status = clReleaseMemObject(mem); + end_time = l_get_timestamp(); + release_times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + } + + mem = clCreateBuffer(context, CL_MEM_READ_ONLY, 64, 0, &status); + assert(status == CL_SUCCESS); + assert(mem); + + std::cout << "Measuring clEnqueueWriteBuffer..." << std::endl; + times = &results["clEnqueueWriteBuffer"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + char host_buf[64] = {0}; + start_time = l_get_timestamp(); + status = clEnqueueWriteBuffer(cq, mem, CL_FALSE, 0, 64, host_buf, 0, NULL, + NULL); + end_time = l_get_timestamp(); + times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + + status = clFinish(cq); + assert(status == CL_SUCCESS); + } + + std::cout << "Measuring clEnqueueReadBuffer..." << std::endl; + times = &results["clEnqueueReadBuffer"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + char host_buf[64] = {0}; + start_time = l_get_timestamp(); + status = clEnqueueReadBuffer(cq, mem, CL_FALSE, 0, 64, host_buf, 0, NULL, + NULL); + end_time = l_get_timestamp(); + times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + + status = clFinish(cq); + assert(status == CL_SUCCESS); + } + + std::cout << "Measuring clCreateKernel/clReleaseKernel..." << std::endl; + create_times = &results["clCreateKernel"]; + release_times = &results["clReleaseKernel"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + start_time = l_get_timestamp(); + kernel = clCreateKernel(program, "kernel4_task_double", &status); + end_time = l_get_timestamp(); + create_times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + assert(kernel); + + start_time = l_get_timestamp(); + status = clReleaseKernel(kernel); + end_time = l_get_timestamp(); + release_times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + } + + kernel = clCreateKernel(program, "kernel4_task_double", &status); + assert(status == CL_SUCCESS); + assert(kernel); + + std::cout << "Measuring clSetKernelArg..." << std::endl; + times = &results["clSetKernelArg"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + start_time = l_get_timestamp(); + status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &mem); + end_time = l_get_timestamp(); + times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + } + + std::cout << "Measuring clEnqueueTask/clReleaseEvent..." << std::endl; + create_times = &results["clEnqueueTask"]; + release_times = &results["clReleaseEvent"]; + for (int inner_rep = 0; inner_rep < INNER_REPS; ++inner_rep) { + cl_event kernel_event; + start_time = l_get_timestamp(); + status = clEnqueueTask(cq, kernel, 0, NULL, &kernel_event); + end_time = l_get_timestamp(); + create_times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + + // wait for kernel to be submitted + cl_int execution_status = CL_QUEUED; + while (execution_status != CL_SUBMITTED) { + status = clGetEventInfo(kernel_event, CL_EVENT_COMMAND_EXECUTION_STATUS, + sizeof(cl_int), &execution_status, NULL); + assert(status == CL_SUCCESS); + } + + int activation_id = kernel_event->cmd.info.ndrange_kernel + .invocation_wrapper->image->activation_id; + acltest_call_kernel_update_callback(activation_id, CL_RUNNING); + acltest_call_kernel_update_callback(activation_id, CL_COMPLETE); + + status = clWaitForEvents(1, &kernel_event); + assert(status == CL_SUCCESS); + + start_time = l_get_timestamp(); + status = clReleaseEvent(kernel_event); + end_time = l_get_timestamp(); + release_times->push_back(end_time - start_time); + assert(status == CL_SUCCESS); + } + + status = clReleaseKernel(kernel); + assert(status == CL_SUCCESS); + status = clReleaseMemObject(mem); + assert(status == CL_SUCCESS); + status = clReleaseCommandQueue(cq); + assert(status == CL_SUCCESS); + status = clReleaseProgram(program); + assert(status == CL_SUCCESS); + status = clReleaseContext(context); + assert(status == CL_SUCCESS); + + std::cout << std::endl; + l_sleep(10 * 1000); // 10 seconds + } + + acl_test_teardown_generic_system(); + + std::cout << "Results:" << std::endl << std::endl; + + if (debug_mode > 0) { + std::cout << "WARNING! You are running this benchmark in debug mode!" + << std::endl; + } + + std::cout << "Name,Time (ns),Standard Deviation (ns)" << std::endl; + results_t::iterator iter = results.begin(); + results_t::iterator iter_end = results.end(); + for (; iter != iter_end; ++iter) { + const std::string &name = iter->first; + times_t × = iter->second; + assert(times.size() == INNER_REPS * OUTER_REPS); + + // The ULL at the end of the constant 0 is actually essential. The type + // of this argument controls the type of the return value of + // std::accumulate(). Without ULL, a 0 by itself is considered to be a + // 32-bit int, which isn't wide enough to store the result. + double mean = + static_cast(std::accumulate(times.begin(), times.end(), 0ULL)) / + static_cast(times.size()); + + double stddev = 0; + times_t::iterator iter2 = times.begin(); + times_t::iterator iter_end2 = times.end(); + for (; iter2 != iter_end2; ++iter2) { + stddev += pow(static_cast(*iter2) - mean, 2); + } + stddev = sqrt(stddev / static_cast(times.size())); + + std::cout << name << "," << mean << "," << stddev << std::endl; + } + + if (debug_mode > 0) { + std::cout << "WARNING! You are running this benchmark in debug mode!" + << std::endl; + } + + std::cout << std::endl; +} diff --git a/fuzz_testing/test/acl_fuzz_test.h b/fuzz_testing/test/acl_fuzz_test.h new file mode 100755 index 00000000..80afaf51 --- /dev/null +++ b/fuzz_testing/test/acl_fuzz_test.h @@ -0,0 +1,114 @@ +// Copyright (C) 2010-2021 Intel Corporation +// SPDX-License-Identifier: BSD-3-Clause + +#ifndef ACL_TEST_H +#define ACL_TEST_H + +#define MIN(X, Y) (((X) < (Y)) ? (X) : (Y)) + +#include + +#include + +#include + +#include + +// Use the default board for the SDK, but with less memory for unit test +// purposes. Must also update the table of hashes in acl_program_test.cpp +#if ACDS_PRO == 0 +#define ACLTEST_DEFAULT_BOARD "s5_net_small" +#else +#define ACLTEST_DEFAULT_BOARD "a10_ref_small" +#endif + +// Default setup and teardown. +void acl_test_setup_generic_system(); +void acl_test_setup_empty_system(); +void acl_test_setup_sample_default_board_system(void); +void acl_test_teardown_generic_system(void); +void acl_test_teardown_system(void); +void acl_test_teardown_sample_default_board_system(void); +void acl_hal_test_setup_generic_system(void); +void acl_hal_test_teardown_generic_system(void); +void acl_test_run_standard_teardown_checks(); + +void acl_test_unsetenv(const char *var); +void acl_test_setenv(const char *var, const char *value); + +cl_context_properties *acl_test_context_prop_preloaded_binary_only(void); + +const unsigned char *acl_test_get_example_binary(size_t *binary_len); + +SimpleString StringFrom(cl_uint x); +SimpleString StringFrom(cl_ulong x); +SimpleString StringFrom(size_t x); + +#ifdef _WIN32 +SimpleString StringFrom(intptr_t x); +#endif + +// Context error notify callback function. +void CL_CALLBACK acl_test_notify_print(const char *errinfo, + const void *private_info, size_t cb, + void *user_data); + +#define ACL_LOCKED(...) \ + do { \ + std::scoped_lock lock{acl_mutex_wrapper}; \ + { __VA_ARGS__; } \ + } while (0) + +/* CAUTION. These are only used in self-tests. + * The runtime does not use these constants any more. + */ +#define ACL_MAX_EVENT (1024 * 16) +#define ACL_MAX_COMMAND \ + ACL_MAX_EVENT /* each event refers to a command. if same number of them, \ + then there is less worry about running out of commands when \ + creating an event */ + +typedef struct mem_data_s mem_data_t; + +typedef struct mem_data_s { + int mmd_interface; + size_t offset; + size_t size; + void *data; + mem_data_t *next; +} mem_data_t; +typedef struct { + cl_bool is_active; + mem_data_t *mem_data; + aocl_mmd_interrupt_handler_fn kernel_interrupt; + void *interrupt_user_data; + aocl_mmd_status_handler_fn kernel_status; + void *status_user_data; +} acl_hal_device_test; + +// This must match the define in acl_kernel_if.c +#define KERNEL_VERSION_ID (0xa0c00001) + +// These must match the defines in acl_kernel_if.c +#define OFFSET_VERSION_ID ((dev_addr_t)0x0000) +#define OFFSET_KERNEL_CRA_SEGMENT ((dev_addr_t)0x0020) +#define OFFSET_SW_RESET ((dev_addr_t)0x0030) +// Default mem_org address. +// Runtime is now using one loaded from autodiscovery, +// rather than hard coded value. +// For tests, autodiscovery will still have the default value. +#define OFFSET_MEM_ORG ((dev_addr_t)0x0018) +#define OFFSET_KERNEL_CRA ((dev_addr_t)0x1000) +#define OFFSET_CONFIGURATION_ROM ((dev_addr_t)0x2000) + +// These must match the defines in acl_pll.c +#define OFFSET_ROM ((dev_addr_t)0x400) +#define OFFSET_RECONFIG_CTRL ((dev_addr_t)0x200) +#define OFFSET_COUNTER ((dev_addr_t)0x100) +#define OFFSET_RESET ((dev_addr_t)0x110) +#define OFFSET_LOCK ((dev_addr_t)0x120) + +// This must match the define in acl_pll.c +#define MAX_KNOWN_SETTINGS 100 + +#endif diff --git a/fuzz_testing/test/acl_globals_fuzz_test.cpp b/fuzz_testing/test/acl_globals_fuzz_test.cpp new file mode 100644 index 00000000..1164c25c --- /dev/null +++ b/fuzz_testing/test/acl_globals_fuzz_test.cpp @@ -0,0 +1,740 @@ +// Copyright (C) 2010-2021 Intel Corporation +// SPDX-License-Identifier: BSD-3-Clause + +#ifdef _MSC_VER +#pragma warning(push) +#pragma warning(disable : 4100) // unreferenced formal parameter +#endif +#include +#ifdef _MSC_VER +#pragma warning(pop) +#endif + +#include + +#include + +#include +#include +#include +#include + +#include "acl_fuzz_test.h" +#include "acl_globals_fuzz_test.h" +#include "acl_hal_fuzz_test.h" + +// Worst case alignment + +// Make these double to ensure alignment. +static ACL_ALIGNED double acltest_global[1024 * 1024 * 4]; + +static ACL_ALIGNED double + acltest_devicelocal[14][16384 / sizeof(double)]; // min permitted local mem + // size is 16K + +// Assumes device only has 32-bit addresses +#define LOCAL_PTR_SIZE_IN_CRA (4) + +static std::vector acltest_laspace_info = {{4, 2048}, + {5, 16768}}; + +static acl_kernel_interface_t acltest_kernels[] = { + {// interface + "kernel0_copy_vecin_vecout", + {{ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_NONE, ACL_ARG_BY_VALUE, sizeof(cl_uint), 0, 0}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_LOCAL, ACL_ARG_BY_VALUE, LOCAL_PTR_SIZE_IN_CRA, 5, 8192}, + {ACL_ARG_ADDR_CONSTANT, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}}}, + {// interface + "kernel1_vecadd_vecin_vecin_vecout", + {{ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}}}, + {// interface + "kernel2_vecscale_vecin_scalar_vecout", + {{ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_NONE, ACL_ARG_BY_VALUE, sizeof(float), 0, 0}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}}}, + {// interface + "kernel3_locals", + { + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_LOCAL, ACL_ARG_BY_VALUE, LOCAL_PTR_SIZE_IN_CRA, 4, 1024}, + {ACL_ARG_ADDR_LOCAL, ACL_ARG_BY_VALUE, LOCAL_PTR_SIZE_IN_CRA, 4, 2048}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + }}, + {// interface + "kernel4_task_double", + { + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + }}, + {// interface + "kernel5_double", + { + {ACL_ARG_ADDR_NONE, ACL_ARG_BY_VALUE, 4, 0, 0}, + }}, + {// interface + "kernel6_profiletest", + {{ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}}}, + {// interface + "kernel7_emptyprofiletest", + {{ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}}}, + {// interface + "kernel8_svm_args", + { + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(int *), 0, 0}, + }}, + {// interface + "kernel9_image_args", + { + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + {ACL_ARG_ADDR_NONE, ACL_ARG_SAMPLER, sizeof(int), 0, 0}, + }}, + {// interface + "kernel11_task_double", + { + {ACL_ARG_ADDR_NONE, ACL_ARG_BY_VALUE, 4, 0, 0}, + }}, + {// interface + "kernel12_task_double", + { + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + }}, + {// interface + "kernel13_multi_vec_lane", + { + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(cl_mem), 0, 0}, + }}, + {// interface + "kernel14_svm_arg_alignment", + { + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(int *), 0, 0, 1}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(int *), 0, 0, 4}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(int *), 0, 0, 8}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(int *), 0, 0, 16}, + {ACL_ARG_ADDR_GLOBAL, ACL_ARG_MEM_OBJ, sizeof(int *), 0, 0, 1024}, + }}}; + +template +static inline constexpr acl_addr_range_t ACL_RANGE_FROM_ARRAY(T (&a)[N]) { + return {reinterpret_cast(&a[0]), + reinterpret_cast( + (reinterpret_cast(&a[0]) + N * sizeof(T)))}; +} + +static acl_system_def_t acltest_empty_system = { + // Device definitions. + 0, +}; + +// accel for acltest_simple_system -> fpga0 +static std::vector acltest_simple_system_accel = { + {0, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[0]), + acltest_kernels[0], + acltest_laspace_info, + {0, 0, 0}, + 0, + /*is_workitem_invariant*/ 0, + /* num_vector_lanes */ 0, + /*num_profiling_counters*/ 149, + 32768}}; + +static acl_system_def_t acltest_simple_system = { + + // Device definitions. + 1, + {{nullptr, + 0, + 1, + 1, + 1, /* half duplex memory transfers */ + 0, + 0, + 0, /* alloc capabilities */ + 0, /* min_host_mem_alignment */ + {"fpga0", + "sample40byterandomhash000000000000000000", + 0, + acltest_simple_system_accel, /* accel */ + {}, /* hal_info */ + 1, // number of global memory systems + { + /* global mem info array */ + { + /* global mem info for memory 0 */ + /* global mem */ ACL_RANGE_FROM_ARRAY(acltest_global), + /* acl_system_global_mem_type_t */ ACL_GLOBAL_MEM_DEVICE_PRIVATE, + /* num_global_bank */ 2, + /* burst_interleaved */ 1, + }, + }}}}}; + +// accel definition for acltest_complex_system->fpga0 +static std::vector acltest_complex_system_device0_accel = { + {0, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[0]), + acltest_kernels[0], + acltest_laspace_info, + {4, 2, 4}, + 0, + 0, + 4, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {1, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[1]), + acltest_kernels[1], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {2, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[2]), + acltest_kernels[2], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {3, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[11]), + acltest_kernels[4], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {4, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[12]), + acltest_kernels[3], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1} // test ptr-to-local + , + {5, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[13]), + acltest_kernels[5], + acltest_laspace_info, + {0, 0, 0}, + 1, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1} // test invariant_workgroup, + , + {6, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[7]), + acltest_kernels[6], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 149, + 32768, + 2, + {}, + {32768, 0, 0}, + 1} // profiler testing + , + {7, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[4]), + acltest_kernels[7], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 0, + {}, + {32768, 0, 0}, + 1}, + {8, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[3]), + acltest_kernels[8], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 1, + {}, + {32768, 0, 0}, + 1}, + {9, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[5]), + acltest_kernels[9], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {10, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[6]), + acltest_kernels[10], + acltest_laspace_info, + {0, 0, 0}, + 1, + 1, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1} // test invariant_workitem, + , + {11, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[11]), + acltest_kernels[11], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1, + 1} // a task that can be fast relaunched + , + {12, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[12]), + acltest_kernels[13], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 0, + 0} // a task that can be fast relaunched + , + {13, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[11]), + acltest_kernels[12], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 3, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1} // test wrong vector lane size +}; + +// accel definition for acltest_complex_system->fpga1 +static std::vector acltest_complex_system_device1_accel = { + {0, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[3]), + acltest_kernels[0], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {1, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[4]), + acltest_kernels[2], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {2, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[12]), + acltest_kernels[3], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1} // test ptr-to-local +}; + +// accel definition for acltest_complex_system->fpga2 +static std::vector acltest_complex_system_device2_accel = { + {0, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[5]), + acltest_kernels[0], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {1, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[6]), + acltest_kernels[4], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {2, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[7]), + acltest_kernels[2], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {3, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[12]), + acltest_kernels[3], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1} // test ptr-to-local +}; + +// accel definition for acltest_complex_system->fpga3 +static std::vector acltest_complex_system_device3_accel = { + {0, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[9]), + acltest_kernels[4], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {1, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[10]), + acltest_kernels[2], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {2, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[8]), + acltest_kernels[0], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1}, + {3, + ACL_RANGE_FROM_ARRAY(acltest_devicelocal[12]), + acltest_kernels[3], + acltest_laspace_info, + {0, 0, 0}, + 0, + 0, + 1, + 0, + 32768, + 3, + {}, + {32768, 0, 0}, + 1} // test ptr-to-local +}; + +static acl_system_def_t acltest_complex_system = { + + /////////// Device definitions + ////////////////////////////////////////////////////////////////////////////////////////////////////////////// + /* num_devices */ 5, + {// All of these have kernels 0 and 2 in common. This is used for + // testing clCreateKernels + {nullptr, + 0, + 1, + 1, + 1, /* half duplex memory transfers */ + 1, + 1, + 0, /* alloc capabilities */ + 0, /* min_host_mem_alignment */ + {"fpga0", + "sample40byterandomhash000000000000000000", + 0, + acltest_complex_system_device0_accel, /* accel */ + {}, /* hal_info */ + 1, // number of global memory systems + { + /* global mem info array */ + { + /* global mem info for memory 0 */ + /* global mem */ ACL_RANGE_FROM_ARRAY(acltest_global), + /* acl_system_global_mem_type_t */ ACL_GLOBAL_MEM_DEVICE_PRIVATE, + /* num_global_bank */ 2, + /* burst_interleaved */ 1, + }, + }}}, + {nullptr, + 1, + 1, + 1, + 2, /* full duplex memory transfers */ + 0, + 0, + 0, /* alloc capabilities */ + 0, /* min_host_mem_alignment */ + {"fpga1", + "sample40byterandomhash000000000000000001", + 0, + acltest_complex_system_device1_accel, /* accel */ + {}, /* hal_info */ + 1, // number of global memory systems + { + /* global mem info array */ + { + /* global mem info for memory 0 */ + /* global mem */ ACL_RANGE_FROM_ARRAY(acltest_global), + /* acl_system_global_mem_type_t */ ACL_GLOBAL_MEM_DEVICE_PRIVATE, + /* num_global_bank */ 2, + /* burst_interleaved */ 1, + }, + }}}, + {nullptr, + 2, + 1, + 1, + 2, /* full duplex memory transfers */ + 0, + 0, + 0, /* alloc capabilities */ + 0, /* min_host_mem_alignment */ + {"fpga2", + "sample40byterandomhash000000000000000002", + 0, + acltest_complex_system_device2_accel, /* accel */ + {}, /* hal_info */ + 1, // number of global memory systems + { + /* global mem info array */ + { + /* global mem info for memory 0 */ + /* global mem */ ACL_RANGE_FROM_ARRAY(acltest_global), + /* acl_system_global_mem_type_t */ ACL_GLOBAL_MEM_DEVICE_PRIVATE, + /* num_global_bank */ 2, + /* burst_interleaved */ 1, + }, + }}}, + {nullptr, + 3, + 1, + 1, + 1, /* half duplex memory transfers */ + 0, + 0, + 0, /* alloc capabilities */ + 0, /* min_host_mem_alignment */ + {"fpga3", + "sample40byterandomhash000000000000000003", + 0, + acltest_complex_system_device3_accel, /* accel */ + {}, /* hal_info */ + 1, // number of global memory systems + { + /* global mem info array */ + { + /* global mem info for memory 0 */ + /* global mem */ ACL_RANGE_FROM_ARRAY(acltest_global), + /* acl_system_global_mem_type_t */ ACL_GLOBAL_MEM_DEVICE_PRIVATE, + /* num_global_bank */ 2, + /* burst_interleaved */ 1, + }, + }}}, + {nullptr, + 4, + 1, + 1, + 1, /* half duplex memory transfers */ + 0, + 0, + 0, /* alloc capabilities */ + 0, /* min_host_mem_alignment */ + {"fpga4", + "sample40byterandomhash000000000000000004", + 0, + {}, /* accel */ + {}, /* hal_info */ + 1, // number of global memory systems + { + /* global mem info array */ + { + /* global mem info for memory 0 */ + /* global mem */ ACL_RANGE_FROM_ARRAY(acltest_global), + /* acl_system_global_mem_type_t */ ACL_GLOBAL_MEM_DEVICE_PRIVATE, + /* num_global_bank */ 2, + /* burst_interleaved */ 1, + }, + }} + + }}}; + +// For use by other tests +const acl_system_def_t *acl_test_get_complex_system_def() { + return &acltest_complex_system; +} + +const acl_system_def_t *acl_test_get_empty_system_def() { + return &acltest_empty_system; +} + +TEST_GROUP(acl_globals_undef){void setup(){acl_mutex_wrapper.lock(); +CHECK(acl_set_hal(acl_test_get_simple_hal())); +} +void teardown() { + acl_reset_hal(); + acl_mutex_wrapper.unlock(); + acl_test_run_standard_teardown_checks(); +} + +void misalign_ptr(void **ptr) { *ptr = ((char *)*ptr) + 1; } +} +; + +TEST(acl_globals_undef, zero_when_unint) { + CHECK(0 == acl_present_board_def()); + CHECK(0 == acl_present_board_is_valid()); +} + +TEST(acl_globals_undef, valid_init_simple) { + CHECK(1 == acl_init(&acltest_simple_system)); + CHECK(0 != acl_present_board_def()); + CHECK(0 != acl_present_board_is_valid()); + // Teardown + acl_reset(); + CHECK(0 == acl_present_board_def()); + CHECK(0 == acl_present_board_is_valid()); +} + +TEST(acl_globals_undef, valid_init_empty) { + CHECK(1 == acl_init(&acltest_empty_system)); + CHECK(0 != acl_present_board_def()); + CHECK(0 != acl_present_board_is_valid()); + // Teardown + acl_reset(); + CHECK(0 == acl_present_board_def()); + CHECK(0 == acl_present_board_is_valid()); +} + +TEST(acl_globals_undef, valid_init_complex) { + CHECK_EQUAL(1, acl_init(&acltest_complex_system)); + CHECK(0 != acl_present_board_def()); + // Teardown + acl_reset(); + CHECK_EQUAL(0, acl_present_board_def()); +} diff --git a/fuzz_testing/test/acl_globals_fuzz_test.h b/fuzz_testing/test/acl_globals_fuzz_test.h new file mode 100755 index 00000000..6149a13f --- /dev/null +++ b/fuzz_testing/test/acl_globals_fuzz_test.h @@ -0,0 +1,12 @@ +// Copyright (C) 2010-2021 Intel Corporation +// SPDX-License-Identifier: BSD-3-Clause + +#ifndef ACL_GLOBALS_TEST_H +#define ACL_GLOBALS_TEST_H + +#include + +const acl_system_def_t *acl_test_get_complex_system_def(); +const acl_system_def_t *acl_test_get_empty_system_def(); + +#endif diff --git a/fuzz_testing/test/acl_hal_fuzz_test.cpp b/fuzz_testing/test/acl_hal_fuzz_test.cpp new file mode 100644 index 00000000..f49bb6f1 --- /dev/null +++ b/fuzz_testing/test/acl_hal_fuzz_test.cpp @@ -0,0 +1,694 @@ +// Copyright (C) 2010-2021 Intel Corporation +// SPDX-License-Identifier: BSD-3-Clause + +#ifdef _MSC_VER +#pragma warning(push) +#pragma warning(disable : 4100) // unreferenced formal parameter +#endif +#include +#ifdef _MSC_VER +#pragma warning(pop) +#endif + +#include + +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +#include "acl_fuzz_test.h" +#include "acl_hal_fuzz_test.h" + +static int acl_test_svm_memory_support = + (CL_DEVICE_SVM_COARSE_GRAIN_BUFFER | CL_DEVICE_SVM_FINE_GRAIN_BUFFER | + CL_DEVICE_SVM_FINE_GRAIN_SYSTEM); +static bool acl_test_physical_memory_support = true; + +// Parts of a valid HAL. +void acltest_hal_init_device(const acl_system_def_t *def); +void acltest_hal_yield(cl_uint num_devices, const cl_device_id *devices) { + UNREFERENCED_PARAMETER(num_devices); + UNREFERENCED_PARAMETER(devices); +} +cl_ulong acltest_hal_get_timestamp(void); +void acltest_hal_copy_hostmem_to_hostmem(cl_event event, const void *src, + void *dest, size_t size); +void acltest_hal_copy_hostmem_to_globalmem(cl_event event, const void *src, + void *dest, size_t size); +void acltest_hal_copy_globalmem_to_hostmem(cl_event event, const void *src, + void *dest, size_t size); +void acltest_hal_copy_globalmem_to_globalmem(cl_event event, const void *src, + void *dest, size_t size); +void acltest_hal_register_callbacks( + acl_event_update_callback event_update, + acl_kernel_update_callback kernel_update, + acl_profile_callback profile_update, + acl_device_update_callback device_update, + acl_process_printf_buffer_callback process_printf); +void acltest_hal_launch_kernel(unsigned int physical_id, + acl_kernel_invocation_wrapper_t *wrapper); +void acltest_hal_unstall_kernel(unsigned int physical_id, int activation_id); +int acltest_hal_program_device(unsigned int physical_id, + const acl_device_def_t *devdef, + const struct acl_pkg_file *binary, + int acl_program_mode); +cl_bool acltest_hal_query_temperature(unsigned int physical_id, cl_int *temp); +int acltest_hal_get_device_official_name(unsigned int physical_device_id, + char *name, size_t size); +int acltest_hal_get_device_vendor_name(unsigned int physical_device_id, + char *name, size_t size); +int acltest_hal_get_profile_data(unsigned int physical_device_id, + unsigned int accel_id, uint64_t *data, + unsigned int length); +int acltest_hal_reset_profile_counters(unsigned int physical_device_id, + unsigned int accel_id); +int acltest_hal_disable_profile_counters(unsigned int physical_device_id, + unsigned int accel_id); +int acltest_hal_enable_profile_counters(unsigned int physical_device_id, + unsigned int accel_id); +int acltest_hal_set_profile_shared_control(unsigned int physical_device_id, + unsigned int accel_id); +int acltest_hal_set_profile_start_cycle(unsigned int physical_device_id, + unsigned int accel_id, uint64_t value); +int acltest_hal_set_profile_stop_cycle(unsigned int physical_device_id, + unsigned int accel_id, uint64_t value); +int acl_test_hal_has_svm_support(unsigned int physical_device_id, int *value); +int acl_test_hal_has_physical_mem(unsigned int physical_device_id); +int acl_test_hal_pll_reconfigure(unsigned int physical_device_id, + const char *pll_settings_str); +void acl_test_hal_reset_kernels(cl_device_id device); +int acl_test_hal_try_devices(cl_uint num_devices, const cl_device_id *devices, + cl_platform_id platform); +int acl_test_hal_close_devices(cl_uint num_devices, + const cl_device_id *devices); +void *acl_test_hal_shared_alloc(cl_device_id device, size_t size, + size_t alignment, mem_properties_t *properties, + int *error); +void *acl_test_hal_host_alloc(const std::vector devices, + size_t size, size_t alignment, + mem_properties_t *properties, int *error); +int acl_test_hal_free(cl_context context, void *ptr); + +static acl_event_update_callback acltest_hal_event_callback = NULL; +static acl_kernel_update_callback acltest_hal_kernel_callback = NULL; +static acl_profile_callback acltest_hal_profile_callback = NULL; +static acl_device_update_callback acltest_hal_device_callback = NULL; +static acl_process_printf_buffer_callback + acltest_process_printf_buffer_callback = NULL; + +static const acl_hal_t simple_hal = {acltest_hal_init_device, + acltest_hal_yield, + acltest_hal_get_timestamp, + acltest_hal_copy_hostmem_to_hostmem, + acltest_hal_copy_hostmem_to_globalmem, + acltest_hal_copy_globalmem_to_hostmem, + acltest_hal_copy_globalmem_to_globalmem, + acltest_hal_register_callbacks, + acltest_hal_launch_kernel, + acltest_hal_unstall_kernel, + acltest_hal_program_device, + acltest_hal_query_temperature, + acltest_hal_get_device_official_name, + acltest_hal_get_device_vendor_name, + 0, + 0, + acltest_hal_get_profile_data, + acltest_hal_reset_profile_counters, + acltest_hal_disable_profile_counters, + acltest_hal_enable_profile_counters, + acltest_hal_set_profile_shared_control, + acltest_hal_set_profile_start_cycle, + acltest_hal_set_profile_stop_cycle, + acl_test_hal_has_svm_support, + acl_test_hal_has_physical_mem, + 0, + acl_test_hal_pll_reconfigure, + acl_test_hal_reset_kernels, + 0, + 0, + 0, + 0, + 0, + 0, + 0, + 0, + acl_test_hal_try_devices, + acl_test_hal_close_devices, + acl_test_hal_host_alloc, + acl_test_hal_free, + acl_test_hal_shared_alloc}; + +// Emulate device memory. +// +// Device memory is modeled as a contiguous set of addresses, and they +// usually start at 0. +// We can't just copy to address 0 on a regular host because that will +// cause a segfault. +// +// So keep a buffer around to store what *should* be in device memory. +// Address 0 translates into the first location in acltest_hal_device_mem. +// We resize it as needed to accommodate the maximum transfer. +// +// This is only required when the runtime is in regular mode; if it's +// emulating an offline device then the host runtime already does mallocs +// for device buffers. +bool acltest_hal_emulate_device_mem = false; +static void *acltest_hal_device_mem = 0; +static size_t acltest_hal_device_mem_size = 0; + +TEST_GROUP(acl_hal){void setup(){acl_mutex_wrapper.lock(); +} +void teardown() { + acl_mutex_wrapper.unlock(); + acl_assert_unlocked(); +} +} +; + +TEST(acl_hal, zero_when_unint) { CHECK(0 == acl_get_hal()); } + +TEST(acl_hal, new_hal_must_exist) { CHECK(0 == acl_set_hal(0)); } + +TEST(acl_hal, valid_init) { + const acl_hal_t *stored_hal; + CHECK(1 == acl_set_hal(&simple_hal)); + + stored_hal = acl_get_hal(); + CHECK(stored_hal != 0); + // Check all fields + CHECK(stored_hal->init_device == simple_hal.init_device); + CHECK(stored_hal->yield == simple_hal.yield); + CHECK(stored_hal->get_timestamp == simple_hal.get_timestamp); + CHECK(stored_hal->copy_hostmem_to_hostmem == + simple_hal.copy_hostmem_to_hostmem); + CHECK(stored_hal->copy_hostmem_to_globalmem == + simple_hal.copy_hostmem_to_globalmem); + CHECK(stored_hal->copy_globalmem_to_hostmem == + simple_hal.copy_globalmem_to_hostmem); + CHECK(stored_hal->copy_globalmem_to_globalmem == + simple_hal.copy_globalmem_to_globalmem); + CHECK(stored_hal->launch_kernel == simple_hal.launch_kernel); + CHECK(stored_hal->unstall_kernel == simple_hal.unstall_kernel); + CHECK(stored_hal->program_device == simple_hal.program_device); + CHECK(stored_hal->query_temperature == simple_hal.query_temperature); + CHECK(stored_hal->get_device_official_name == + simple_hal.get_device_official_name); + CHECK(stored_hal->get_device_vendor_name == + simple_hal.get_device_vendor_name); + CHECK(stored_hal->get_profile_data == simple_hal.get_profile_data); + CHECK(stored_hal->reset_profile_counters == + simple_hal.reset_profile_counters); + CHECK(stored_hal->disable_profile_counters == + simple_hal.disable_profile_counters); + CHECK(stored_hal->enable_profile_counters == + simple_hal.enable_profile_counters); + CHECK(stored_hal->set_profile_start_cycle == + simple_hal.set_profile_start_cycle); + CHECK(stored_hal->set_profile_stop_cycle == + simple_hal.set_profile_stop_cycle); + CHECK(stored_hal->has_svm_memory_support == + simple_hal.has_svm_memory_support); + CHECK(stored_hal->has_physical_mem == simple_hal.has_physical_mem); + CHECK(stored_hal->pll_reconfigure == simple_hal.pll_reconfigure); + CHECK(stored_hal->reset_kernels == simple_hal.reset_kernels); + + // Reset should invalidate + acl_reset_hal(); + CHECK(0 == acl_get_hal()); +} + +TEST(acl_hal, field_check) { + acl_hal_t bad_hal; + + bad_hal = simple_hal; + bad_hal.init_device = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.yield = 0; + CHECK(0 != acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.get_timestamp = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.copy_hostmem_to_hostmem = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.copy_hostmem_to_globalmem = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.copy_globalmem_to_hostmem = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.copy_globalmem_to_globalmem = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.launch_kernel = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.unstall_kernel = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.program_device = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.query_temperature = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.get_device_official_name = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.get_device_vendor_name = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.get_profile_data = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.reset_profile_counters = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.disable_profile_counters = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.enable_profile_counters = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.set_profile_start_cycle = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.set_profile_stop_cycle = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.has_svm_memory_support = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.has_physical_mem = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.pll_reconfigure = 0; + CHECK(0 == acl_set_hal(&bad_hal)); + bad_hal = simple_hal; + bad_hal.reset_kernels = 0; + CHECK(0 == acl_set_hal(&bad_hal)); +} + +TEST(acl_hal, debugging) { + const char test_str[] = "hal debug printf test\n"; + CHECK(acl_set_hal(&simple_hal)); + const acl_hal_t *hal = acl_get_hal(); + + acl_test_setenv("ACL_DEBUG", "0"); + char *acl_debug_env = getenv("ACL_DEBUG"); + CHECK(acl_debug_env); + debug_mode = atoi(acl_debug_env); + // The simple hal starts out in non-debug mode. + CHECK_EQUAL(0, debug_mode); + // Nothing gets printed when debug is 0. + CHECK_EQUAL(0, acl_print_debug_msg("some text")); + CHECK_EQUAL(0, acl_print_debug_msg("")); + + // Now start debugging. + acl_test_setenv("ACL_DEBUG", "4"); + acl_debug_env = getenv("ACL_DEBUG"); + CHECK(acl_debug_env); + debug_mode = atoi(acl_debug_env); + CHECK_EQUAL(4, debug_mode); + acl_test_setenv("ACL_DEBUG", "3"); + acl_debug_env = getenv("ACL_DEBUG"); + CHECK(acl_debug_env); + debug_mode = atoi(acl_debug_env); + CHECK_EQUAL(3, debug_mode); + acl_test_setenv("ACL_DEBUG", "2"); + acl_debug_env = getenv("ACL_DEBUG"); + CHECK(acl_debug_env); + debug_mode = atoi(acl_debug_env); + CHECK_EQUAL(2, debug_mode); + CHECK_EQUAL( + sizeof(test_str) / sizeof(test_str[0]) - 1, + acl_print_debug_msg(test_str)); // -1 because source array includes a NUL + // char, but we don't print it. + CHECK_EQUAL(0, acl_print_debug_msg("")); + acl_test_setenv("ACL_DEBUG", "0"); + acl_debug_env = getenv("ACL_DEBUG"); + CHECK(acl_debug_env); + debug_mode = atoi(acl_debug_env); + CHECK_EQUAL(0, debug_mode); + CHECK_EQUAL(0, acl_print_debug_msg(test_str)); + acl_test_unsetenv("ACL_DEBUG"); + + cl_int temp; + CHECK_EQUAL(1, hal->query_temperature(0, &temp)); + CHECK_EQUAL(0, temp); + +#define BUF_SIZE 1024 + char name[BUF_SIZE]; + int size_returned; + size_returned = hal->get_device_official_name(0, name, BUF_SIZE); + CHECK(size_returned < BUF_SIZE); + CHECK(strcmp(name, "Test Device") == 0); + size_returned = hal->get_device_vendor_name(0, name, BUF_SIZE); + CHECK(size_returned < BUF_SIZE); + CHECK(strcmp(name, "Intel(R) Corporation") == 0); + + uint64_t temp_u64; + // Return values in following six calls are defined by the test HAL as + // 1-6, simply to verify ordering of the function pointers. + CHECK_EQUAL(1, hal->get_profile_data(0, 1, &temp_u64, 1)); + CHECK_EQUAL(2, hal->reset_profile_counters(0, 1)); + CHECK_EQUAL(3, hal->disable_profile_counters(0, 1)); + CHECK_EQUAL(4, hal->enable_profile_counters(0, 1)); + CHECK_EQUAL(5, hal->set_profile_start_cycle(0, 1, 2)); + CHECK_EQUAL(6, hal->set_profile_stop_cycle(0, 1, 2)); +} + +///////////////////////////// +// A very simple HAL + +void acltest_hal_init_device(const acl_system_def_t *def) { + def = def; + + // Only emulate memory if using an offline device. + // In that case the regular runtime will already do malloc's to cover + // device memory. + const char *env = acl_getenv("CL_CONTEXT_OFFLINE_DEVICE_INTELFPGA"); + acltest_hal_emulate_device_mem = (env == NULL || env[0] == 0); +} + +const acl_hal_t *acl_test_get_simple_hal(void) { return &simple_hal; } + +cl_ulong acltest_hal_get_timestamp(void) { + static cl_ulong now = 0; + return now++; // always advancing is reasonable. +} + +void acltest_hal_copy_hostmem_to_hostmem(cl_event event, const void *src, + void *dest, size_t size) { + acltest_hal_event_callback( + event, CL_RUNNING); // in "real life" this in response to a hw message + size_t i; + acl_print_debug_msg(" Copying %zu bytes from %p to %p event %p\n", size, src, + dest, event); + for (i = 0; i < size; i++) { + ((char *)dest)[i] = ((char *)src)[i]; + } + acltest_hal_event_callback( + event, CL_COMPLETE); // in "real life" this in response to a hw message +} +void acltest_hal_copy_hostmem_to_globalmem(cl_event event, const void *src, + void *dest, size_t size) { + // For testing purposes, the same. + (void)acltest_translate_device_address(dest, size); + void *dev_ptr = acltest_translate_device_address(dest, 0); + acltest_hal_copy_hostmem_to_hostmem(event, src, dev_ptr, size); +} +void acltest_hal_copy_globalmem_to_hostmem(cl_event event, const void *src, + void *dest, size_t size) { + // For testing purposes, the same. + (void)acltest_translate_device_address(src, size); + void *dev_ptr = acltest_translate_device_address(src, 0); + acltest_hal_copy_hostmem_to_hostmem(event, dev_ptr, dest, size); +} +void acltest_hal_copy_globalmem_to_globalmem(cl_event event, const void *src, + void *dest, size_t size) { + // For testing purposes, the same. + (void)acltest_translate_device_address(src, size); + (void)acltest_translate_device_address(dest, size); + void *src_ptr = acltest_translate_device_address(src, 0); + void *dest_ptr = acltest_translate_device_address(dest, 0); + acltest_hal_copy_hostmem_to_hostmem(event, src_ptr, dest_ptr, size); +} + +void acltest_hal_register_callbacks( + acl_event_update_callback event_update, + acl_kernel_update_callback kernel_update, + acl_profile_callback profile_update, + acl_device_update_callback device_update, + acl_process_printf_buffer_callback process_printf) { + acltest_hal_event_callback = event_update; + acltest_hal_kernel_callback = kernel_update; + acltest_hal_profile_callback = profile_update; + acltest_hal_device_callback = device_update; + acltest_process_printf_buffer_callback = process_printf; +} + +void acltest_call_event_update_callback(cl_event event, int new_status) { + acltest_hal_event_callback(event, new_status); +} + +void acltest_call_kernel_update_callback(int activation_id, cl_int status) { + acltest_hal_kernel_callback(activation_id, status); +} + +void acltest_call_device_update_callback(unsigned physical_device_id, + int device_status) { + acltest_hal_device_callback(physical_device_id, + (CL_EXCEPTION_TYPE_INTEL)device_status, NULL, 0); +} + +void acltest_call_printf_buffer_callback(int activation_id, int size, + int stalled) { + acltest_process_printf_buffer_callback(activation_id, size, stalled); +} + +void acltest_hal_launch_kernel( + unsigned int physical_id, + acl_kernel_invocation_wrapper_t *invocation_wrapper) { + // Send a message to the device controller, pointing at + // the global buffer pointed at by invocation_wrapper->mem's global + // memory buffer. + + // For unit testing, just trust the tester is doing the right thing. + invocation_wrapper = invocation_wrapper; // avoid warning on MSVC + physical_id = physical_id; +} + +void acltest_hal_unstall_kernel(unsigned int physical_id, int invocation_id) { + invocation_id = invocation_id; // avoid warning on windows + // For unit testing, just trust the tester is doing the right thing. + physical_id = physical_id; +} + +cl_bool acltest_hal_query_temperature(unsigned int physical_id, cl_int *temp) { + *temp = 0; // Avoid Windows warning + physical_id = physical_id; + return 1; // Fake success +} + +int acltest_hal_get_device_official_name(unsigned int physical_device_id, + char *name, size_t size) { + static const char *the_name = "Test Device"; + physical_device_id = physical_device_id; // Avoid Windows warning + const size_t the_size = strnlen(the_name, size - 1) + 1; + strncpy(name, the_name, the_size - 1); + name[the_size - 1] = '\0'; + return static_cast(the_size); +} + +int acltest_hal_get_device_vendor_name(unsigned int physical_device_id, + char *name, size_t size) { + static const char *the_name = "Intel(R) Corporation"; + physical_device_id = physical_device_id; // Avoid Windows warning + const size_t the_size = strnlen(the_name, size - 1) + 1; + strncpy(name, the_name, the_size - 1); + name[the_size - 1] = '\0'; + return static_cast(the_size); +} + +int acltest_hal_program_device(unsigned int physical_id, + const acl_device_def_t *devdef, + const struct acl_pkg_file *binary, + int acl_program_mode) { + devdef = devdef; + binary = binary; + physical_id = physical_id; + acl_program_mode = acl_program_mode; + + char *str_use_jtag_programming = getenv("ACL_PCIE_USE_JTAG_PROGRAMMING"); + // program the device based on the acl_program_mode and + // str_use_jtag_programming + if (acl_program_mode == ACL_PROGRAM_PRESERVE_MEM && + str_use_jtag_programming) { + return ACL_PROGRAM_CANNOT_PRESERVE_GLOBAL_MEM; + } + + return 0; // signals success +} + +int acltest_hal_get_profile_data(unsigned int physical_device_id, + unsigned int accel_id, uint64_t *data, + unsigned int length) { + physical_device_id = + physical_device_id; // avoid warning and hence build break on Windows. + accel_id = accel_id; + data = data; + length = length; + return 1; +} + +int acltest_hal_reset_profile_counters(unsigned int physical_device_id, + unsigned int accel_id) { + physical_device_id = + physical_device_id; // avoid warning and hence build break on Windows. + accel_id = accel_id; + return 2; +} + +int acltest_hal_disable_profile_counters(unsigned int physical_device_id, + unsigned int accel_id) { + physical_device_id = + physical_device_id; // avoid warning and hence build break on Windows. + accel_id = accel_id; + return 3; +} + +int acltest_hal_enable_profile_counters(unsigned int physical_device_id, + unsigned int accel_id) { + physical_device_id = + physical_device_id; // avoid warning and hence build break on Windows. + accel_id = accel_id; + return 4; +} + +int acltest_hal_set_profile_shared_control(unsigned int physical_device_id, + unsigned int accel_id) { + physical_device_id = + physical_device_id; // avoid warning and hence build break on Windows. + accel_id = accel_id; + return 7; +} + +int acltest_hal_set_profile_start_cycle(unsigned int physical_device_id, + unsigned int accel_id, uint64_t value) { + physical_device_id = + physical_device_id; // avoid warning and hence build break on Windows. + accel_id = accel_id; + value = value; + return 5; +} + +int acltest_hal_set_profile_stop_cycle(unsigned int physical_device_id, + unsigned int accel_id, uint64_t value) { + physical_device_id = + physical_device_id; // avoid warning and hence build break on Windows. + accel_id = accel_id; + value = value; + return 6; +} + +void acl_test_hal_set_svm_memory_support(int value) { + acl_test_svm_memory_support = value; +} + +void acl_test_hal_set_physical_memory_support(bool value) { + acl_test_physical_memory_support = value; +} + +int acl_test_hal_has_svm_support(unsigned int physical_device_id, int *value) { + physical_device_id = physical_device_id; // Avoid Windows warning + *value = acl_test_svm_memory_support; + return *value != 0; +} + +int acl_test_hal_has_physical_mem(unsigned int physical_device_id) { + physical_device_id = physical_device_id; // Avoid Windows warning + return acl_test_physical_memory_support; +} + +int acl_test_hal_pll_reconfigure(unsigned int physical_device_id, + const char *pll_settings_str) { + physical_device_id = physical_device_id; + pll_settings_str = pll_settings_str; + return 0; +} +void acl_test_hal_reset_kernels(cl_device_id device) { device = device; } + +int acl_test_hal_try_devices(cl_uint num_devices, const cl_device_id *devices, + cl_platform_id platform) { + // Windows warnings: + platform = platform; + + for (unsigned i = 0; i < num_devices; i++) { + devices[i]->opened_count++; + } + return 0; +} + +int acl_test_hal_close_devices(cl_uint num_devices, + const cl_device_id *devices) { + for (unsigned i = 0; i < num_devices; i++) { + assert(devices[i]->opened_count > 0); + devices[i]->opened_count--; + } + return 0; +} + +void *acl_test_hal_shared_alloc(cl_device_id device, size_t size, + size_t alignment, mem_properties_t *properties, + int *error) { + device = device; + size = size; + alignment = alignment; + properties = properties; + error = error; + return (void *)0xdeadbeefdeadbeef; +} + +void *acl_test_hal_host_alloc(const std::vector, size_t, size_t, + mem_properties_t *, int *) { + return (void *)0xdeadbeefdeadbeef; +} + +int acl_test_hal_free(cl_context context, void *ptr) { + context = context; + ptr = ptr; + return 0; +} + +/////////////// +// Emulate device memory. +// This is only necessary + +// Translate the device pointer to something we can actually write to in host +// memory. And translate them into a host pointer. It's transient though! +void *acltest_translate_device_address(const void *device_ptr, size_t offset) { + if (!acltest_hal_emulate_device_mem) { + // cast away const and get bottom 48 bits + void *result = (void *)(0xffffffffffff & ((uintptr_t)device_ptr)); + return result; + } + + uintptr_t max_dev_addr = 0xfffffff & ((uintptr_t)device_ptr + offset); + acl_print_debug_msg("maxdevaddr %" PRIuPTR "\n", max_dev_addr); + if (!acltest_hal_device_mem || max_dev_addr >= acltest_hal_device_mem_size) { + if (!acltest_hal_device_mem) { + acltest_hal_device_mem = acl_malloc(max_dev_addr + 1); + acl_print_debug_msg("malloc %p\n", acltest_hal_device_mem); + } else { + acltest_hal_device_mem = + acl_realloc(acltest_hal_device_mem, max_dev_addr + 1); + acl_print_debug_msg("realloc %p\n", acltest_hal_device_mem); + } + assert(acltest_hal_device_mem); + acltest_hal_device_mem_size = max_dev_addr + 1; + } + void *result = ((char *)acltest_hal_device_mem) + max_dev_addr; + acl_print_debug_msg(" dev %p --> fake %p (base %p size %llx)\n", + device_ptr, result, acltest_hal_device_mem, + (unsigned long long)acltest_hal_device_mem_size); + return result; +} + +void acltest_hal_teardown(void) { + if (acltest_hal_device_mem) { + acl_free(acltest_hal_device_mem); + } + acltest_hal_device_mem = 0; + acltest_hal_device_mem_size = 0; + acltest_hal_emulate_device_mem = false; +} diff --git a/fuzz_testing/test/acl_hal_fuzz_test.h b/fuzz_testing/test/acl_hal_fuzz_test.h new file mode 100755 index 00000000..a7e3949a --- /dev/null +++ b/fuzz_testing/test/acl_hal_fuzz_test.h @@ -0,0 +1,27 @@ +// Copyright (C) 2010-2021 Intel Corporation +// SPDX-License-Identifier: BSD-3-Clause + +#ifndef ACL_HAL_TEST_H +#define ACL_HAL_TEST_H + +// A simple HAL for testing. + +const acl_hal_t *acl_test_get_simple_hal(void); + +void acltest_hal_teardown(void); + +// Make sure that all these device addresses have host storage representing +// them. And translate them into a host pointer. It's transient though! +void *acltest_translate_device_address(const void *device_ptr, size_t offset); + +void acl_test_hal_set_svm_memory_support(int value); +void acl_test_hal_set_physical_memory_support(bool value); + +extern bool acltest_hal_emulate_device_mem; + +void acltest_call_event_update_callback(cl_event event, int new_status); +void acltest_call_kernel_update_callback(int activation_id, cl_int status); +void acltest_call_printf_buffer_callback(int activation_id, int size, + int stalled); + +#endif