diff --git a/.github/workflows/format.yml b/.github/workflows/format.yml index 83e4a1bd..e09afbfe 100644 --- a/.github/workflows/format.yml +++ b/.github/workflows/format.yml @@ -10,7 +10,7 @@ permissions: jobs: format: - runs-on: ubuntu-latest + runs-on: ubuntu-24.04 steps: - uses: actions/checkout@v4 @@ -20,14 +20,14 @@ jobs: with: python-version: '3.10' - - name: Install dependencies - run: pip install -r assembler_tools/hec-assembler-tools/requirements.txt + - name: Install Development Dependencies + run: pip install -r requirements-dev.txt - - name: Install pre-commit - run: pip install pre-commit + - name: Installing Component-specific Dependencies + run: pip install -r assembler_tools/hec-assembler-tools/requirements.txt - - name: Install pylint - run: pip install pylint + - name: Install Apt Dependencies + run: sudo apt install -y clang-format-14 - name: Fetch main branch for diff run: git fetch origin main diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index d3467f90..2fd0f8d9 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -62,6 +62,7 @@ repos: - -rn # Only display messages - -sn # Don't display the score - --source-roots=p-isa_tools/kerngen # Working directory + - --source-roots=p-isa_tools/program_mapper # Added directory - --source-roots=assembler_tools/hec-assembler-tools # Added directory - id: clang-format-14 name: clang-format-14 @@ -76,3 +77,4 @@ repos: files: \.(c|cc|cxx|cpp|h|hpp|hxx)$ args: - --recursive + - --filter=-build/c++17 diff --git a/p-isa_tools/CMakeLists.txt b/p-isa_tools/CMakeLists.txt index 61b6b2b3..80f767e2 100644 --- a/p-isa_tools/CMakeLists.txt +++ b/p-isa_tools/CMakeLists.txt @@ -28,7 +28,7 @@ message(ENABLE_DATA_FORMATS="${ENABLE_DATA_FORMATS}") option(ENABLE_FUNCTIONAL_MODELER "Enable building of functional modeler" ON) message(ENABLE_FUNCTIONAL_MODELER="${ENABLE_FUNCTIONAL_MODELER}") -option(ENABLE_PROGRAM_MAPPER "Enable building of program mapper" OFF) +option(ENABLE_PROGRAM_MAPPER "Enable building of program mapper" ON) message(ENABLE_PROGRAM_MAPPER="${ENABLE_PROGRAM_MAPPER}") option(ENABLE_P_ISA_UTILITIES "Enable building of p-isa utilities" OFF) diff --git a/p-isa_tools/program_mapper/CMakeLists.txt b/p-isa_tools/program_mapper/CMakeLists.txt new file mode 100644 index 00000000..9c8a6e06 --- /dev/null +++ b/p-isa_tools/program_mapper/CMakeLists.txt @@ -0,0 +1,23 @@ +############################## +# HERACLES Program Mapper +############################## + +project(program_mapper LANGUAGES CXX) + +set(HERACLES_PGM_SOURCES + "p_isa/pisakernel.cpp" + "p_isa/pisa_graph_optimizer.cpp" + "poly_program/polyprogram.cpp" + "trace_parser/program_trace_helper.cpp" + "p_isa/pisa_test_generator.cpp" + "utility_functions.h" + "main.cpp" +) + +add_executable(program_mapper ${HERACLES_PGM_SOURCES} ${IDE_HEADERS}) +if(ENABLE_DATA_FORMATS) + target_link_libraries(program_mapper PUBLIC HERACLES_DATA_FORMATS::heracles_data_formats) + target_compile_definitions(program_mapper PRIVATE ENABLE_DATA_FORMATS) +endif() +target_link_libraries(program_mapper PUBLIC nlohmann_json::nlohmann_json snap OpenMP::OpenMP_CXX common) +target_include_directories(program_mapper PRIVATE ${INCLUDE_DIRS}) diff --git a/p-isa_tools/program_mapper/README.md b/p-isa_tools/program_mapper/README.md new file mode 100644 index 00000000..5bb3b89e --- /dev/null +++ b/p-isa_tools/program_mapper/README.md @@ -0,0 +1,130 @@ +# HERACLES Program Mapper + +## Table of Contents +1. [Requirements](#requirements) +2. [Build Configuration](#build-configuration) + 1. [Build Type](#build-type) + 1. [Third-Party Components](#third-party-components) +3. [Building](#building) +4. [Running the Program Mapper](#running-the-program-mapper) +5. [Code Formatting](#code-formatting) + +## Requirements + +Current build system uses `CMake`. + +Tested Configuration(s) +- Ubuntu 22.04 (also tested on WSL2) +- C++17 +- GCC == 11.3 +- CMake >= 3.22.1 +- SNAP (used to support graph features) +- JSON for Modern CPP >= 3.11 + +## Build Configuration + +The current build system is minimally configurable but will be improved with +time. The project directory is laid out as follows + +- __program_mapper__ *src directory for the program mapper* +- __common__ *Common code used by p-isa tools* + +**NOTE:** If using an IDE then it is recommended to set the `INC_HEADERS` flag +to include the header files in the project filesystem. This can be done +via `-DINC_HEADERS=TRUE`. + +### Build Type + +If no build type is specified, the build system will build in Debug +mode. Use `-DCMAKE_BUILD_TYPE` configuration variable to set your preferred +build type: + +- `-DCMAKE_BUILD_TYPE=Debug` : debug mode (default if no build type is specified). +- `-DCMAKE_BUILD_TYPE=Release` : release mode. Compiler optimizations for release enabled. +- `-DCMAKE_BUILD_TYPE=RelWithDebInfo` : release mode with debug symbols. +- `-DCMAKE_BUILD_TYPE=MinSizeRel` : release mode optimized for size. + +#### Third-Party Components +This backend requires the following third party components: + +- [SNAP](https://github.com/snap-stanford/snap.git) +- [JSON for modern c++](https://github.com/nlohmann/json) + +These external dependencies are fetched and built at configuration time by +`cmake`, see below how to build the project. + +## Building +Build from the top level of p-isa-tools with Cmake as follows: + +```bash +cmake -S . -B build -DCMAKE_BUILD_TYPE=Release +cmake --build build -j +``` + +Build type can also be changed to `Debug` depending on current needs (Debug +should be used if the tool is being used to actively debug failing kernels). + +## Running the Program Mapper + +Located in `build/bin` is an executable called **program_mapper**. This program +can be used to generate a graph of combined p-isa instructions with a +supplementary memory file, and a combined p-isa kernel from a program trace. +The program accepts a number of commandline options to control its usage. + +A typical run is of the form +```bash +program_mapper +``` + +The standard list of currently supported options are listed below. +```bash +Usage: + program_mapper program_trace kerngen_loc OPTIONS + +POSITIONAL ARGUMENTS: 2 +program_trace + Location of a file containing a list in csv format for p_isa instructions + +kerngen_loc + Location of the kerngen.py file + + +OPTIONS: +--cache_dir, --cache, -c + Sets the name of the kernel cache directory [Default: ./kernel_cache] + +--disable_cache, --no_cache, -dc + Disables the use of a cache for Ninja kernels + +--disable_graphs, --graphs, -g + Disables graph building and features + +--disable_namespace, --nns, -n + Disables applying register name spacing on PISAKernel nodes + +--dot_file_name, -df + Sets the name of the output dot file + +--enable_memory_bank_output, --banks, -b + Will output P-ISA programs with registers that include hard coded memory banks when enabled + +--export_dot, -ed + Export seal trace and p_isa graphs to dot file format + +--out_dir, --out, -o + Sets the location for all output files [Default: ./] + +--remove_cache, --rm_cache, -rc + Remove the kernel cache directory at the end of the program + +--verbose, -v + Enables more verbose execution reporting to std out + +-h, /h, \h, --help, /help, \help + Shows this help. +``` + +## Code Formatting +The repository includes `pre-commit` and `clang-format` hooks to help ensure +code consistency. It is recommended to install `pre-commit` and `pre-commit +hooks` prior to committing to repo. diff --git a/p-isa_tools/program_mapper/main.cpp b/p-isa_tools/program_mapper/main.cpp new file mode 100644 index 00000000..0a1cd72f --- /dev/null +++ b/p-isa_tools/program_mapper/main.cpp @@ -0,0 +1,136 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include +#include +#include +#include +#include + +#include "p_isa/pisa_test_generator.h" +#include "program_mapper.h" +#include +#include +#include +#include +#include +#include +#include + +#include + +using DATA_TYPE = uint32_t; +namespace fs = std::filesystem; + +inline pisa::ProgramMapperArguments parseCommandLineArguments(int argc, char **argv) +{ + pisa::ProgramMapperArguments args; + // clang-format off + argmap::ArgMap() + .separator(argmap::ArgMap::Separator::WHITESPACE) + .required() + .positional() + .arg("program_trace", args.program_trace_location, + "Location of a file containing a list in csv format for p_isa instructions", "") + .arg("kerngen_loc", args.kerngen, "Location of the kerngen.py file", "") + .optional() + .toggle() + .arg({"--verbose", "-v"}, args.verbose,"Enables more verbose execution reporting to std out", "") + .arg({"--export_dot", "-ed"}, args.export_dot, "Export polynomial program and p_isa program graphs to dot file format", "") + .arg({"--remove_cache", "--rm_cache", "-rc"}, args.remove_cache, + "Remove the kernel cache directory at the end of the program", "") + .arg({"--enable_memory_bank_output", "--banks", "-b"}, args.output_memory_bank, + "Will output P-ISA programs with registers that include hard coded memory banks when enabled", "") + .arg({"--export_trace", "-pb"}, args.export_program_trace,"Exports trace to opposite of input format, CSV <-> Pb", "") + .arg({"--enable_intermediates", "-ei"}, args.enable_intermediates,"Enables intermediates by disabling name spacing and other optimizations on intermediate values", "") + .toggle(false) + .arg({"--disable_graphs", "--graphs", "-g"}, args.generate_graphs, + "Disables graph building and features", "") + .arg({"--disable_namespace", "--nns", "-n"}, args.apply_name_spacing, + "Disables applying register name spacing on PISAKernel nodes", "") + .arg({"--disable_cache", "--no_cache", "-dc"}, args.use_kernel_cache, + "Disables the use of a cache for Ninja kernels", "") + .named() + .arg({"--dot_file_name", "-df"}, args.dot_file_name , "Sets the name of the output dot file", "") + .arg({"--cache_dir", "--cache", "-c"}, args.cache_dir, "Sets the name of the kernel cache directory") + .arg({"--out_dir", "--out", "-o"}, args.out_dir, "Sets the location for all output files") + .arg({"--generated_json", "--generate", "-gen"}, args.generated_name, "Enables generation of JSON data file and specifies name") + .arg({"--kernel_library", "--kernlib", "-kl"}, args.kernel_library, "Specifies which kernel library to use.") + .parse(argc, argv); + // clang-format on + + // Post-processing of positional arguments + auto strip_substring = [&path = args.program_trace_location](const std::string &substr) -> fs::path { + auto ret = path.stem(); + auto pos = ret.string().find(substr); + if (pos != std::string::npos) + { + ret = ret.string().erase(pos, substr.size()); + } + return ret; + }; + + args.outfile_prefix = args.out_dir / (strip_substring("_program_trace").string() + "_pisa"); + if (args.dot_file_name.empty()) + { + args.dot_file_name = args.out_dir / args.program_trace_location.stem(); + args.dot_file_name.replace_extension("dot"); + } + + return args; +} + +int main(int argc, char **argv) +{ + try + { + auto arguments = parseCommandLineArguments(argc, argv); + + std::shared_ptr program_trace; + // Parses a polynomial program into a vector of HEOPerations + + if (arguments.program_trace_location.extension() == ".csv") + { + program_trace = PolynomialProgramHelper::parse(arguments.program_trace_location, POLYNOMIAL_PROGRAM_FORMAT::CSV); +#ifdef ENABLE_DATA_FORMATS + if (arguments.export_program_trace) + { + PolynomialProgramHelper::writeTraceToProtoBuff(program_trace, arguments.program_trace_location.filename().string() + ".bin"); + } +#endif + } +#ifdef ENABLE_DATA_FORMATS + else if (arguments.program_trace_location.extension() == ".bin") + { + program_trace = PolynomialProgramHelper::parse(arguments.program_trace_location, POLYNOMIAL_PROGRAM_FORMAT::PROTOBUFF); + if (arguments.export_program_trace) + { + PolynomialProgramHelper::writeTraceToCSV(program_trace, arguments.program_trace_location.filename().string() + ".csv"); + } + } +#endif + else + { + throw std::runtime_error("Unsupported data format"); + } + + if (arguments.verbose) + std::cout << "Instruction count: " << program_trace->operations().size() << std::endl; + + pisa::ProgramMapper program_mapper; + program_mapper.setArguments(arguments); + program_mapper.generatePisaProgramFromHEProgram(program_trace); + + return 0; + } + catch (const std::runtime_error &err) + { + std::cerr << "Caught std::runtime_error in main: " << err.what() << std::endl; + return 1; + } + catch (...) + { + std::cerr << "Unknown exception caught in main" << std::endl; + return 1; + } +} diff --git a/p-isa_tools/program_mapper/p_isa/pisa_graph_optimizer.cpp b/p-isa_tools/program_mapper/p_isa/pisa_graph_optimizer.cpp new file mode 100644 index 00000000..715279cf --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/pisa_graph_optimizer.cpp @@ -0,0 +1,318 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include "pisa_graph_optimizer.h" + +PISAGraphOptimizer::PISAGraphOptimizer() +{ +} + +std::vector PISAGraphOptimizer::generateInstructionStreamFromGraph(graph::Graph &p_isa_graph, bool fixed_order, std::vector instr_order) +{ + if (fixed_order == false) + { + auto instruction_graph = p_isa_graph.clone(); + + auto all_nodes = instruction_graph.getNodes(); + for (auto node : all_nodes) + { + if (node.GetDat().type != graph::OPERATION) + { + instruction_graph.removeNodeMaintainConnections(node); + } + } + + auto instruction_graph_consumable = instruction_graph.clone(); + + std::vector>> input_layers; + //Layer peel + while (instruction_graph_consumable.getNodeCount() > 0) + { + auto inputs = instruction_graph_consumable.getInputNodes(); + //input_layers.push_back(inputs); + std::vector> layer; + for (auto &node : inputs) + { + + layer.push_back(instruction_graph.getNode(node.GetId())); + //std::cout << *node.GetDat().instruction << std::endl; + instruction_graph_consumable.removeNode(node); + //std::cout << *node.GetDat().instruction << std::endl; + } + input_layers.push_back(layer); + } + + if (perform_variable_isolation) + { + isolateGraphVariables(p_isa_graph, input_layers); + } + + std::vector instructions; + + for (auto &layer : input_layers) + { + for (auto &node : layer) + { + //std::cout << *node.GetDat().instruction << std::endl; + instructions.push_back(node.GetDat().instruction); + } + } + return instructions; + } + else + { + return instr_order; + } +} + +void PISAGraphOptimizer::isolateGraphVariables(graph::Graph &p_isa_graph, std::vector>> &input_layers) +{ + //Generate rename black list + for (auto &layer : input_layers) + { + for (auto &node : layer) + { + nodeLocklist(node, p_isa_graph); + } + } + + //Adjust variables and generate unique node names + for (auto &layer : input_layers) + { + for (auto &node : layer) + { + nodeVariableAdjustment(node, p_isa_graph); + } + } + //Correct registers for accumulate ops + // for(auto layer = input_layers.rbegin(); layer != input_layers.rend(); layer++){ + // //for(auto& layer : input_layers.rbegin()) { + // for(auto& node : *layer) { + // nodeMACVariableAdjustment(node,p_isa_graph); + // } + // } + + for (auto &layer : input_layers) + { + for (auto &node : layer) + { + nodeInstructionAdjustment(node, p_isa_graph); + } + } + + return; +} + +void PISAGraphOptimizer::applyDuplicateInputVariableSeparation(std::vector &instr_order) +{ + std::vector newOrder; + for (auto instr : instr_order) + { + //Check input for matches + bool match = false; + std::pair matching_indices; + + if (instr->numInputOperands() == 2) + { + auto &operand_0 = instr->getInputOperand(0); + auto &operand_1 = instr->getInputOperand(1); + if (operand_0.location() == operand_1.location()) + { + std::cout << "Duplicate input variable detected" << std::endl; + match = true; + } + } + else if (instr->numInputOperands() == 3) + { + auto &operand_0 = instr->getInputOperand(0); + auto &operand_1 = instr->getInputOperand(1); + auto &operand_2 = instr->getInputOperand(2); + + if (operand_0.location() == operand_1.location()) + { + match = true; + matching_indices = std::pair(0, 1); + } + if (operand_0.location() == operand_2.location()) + { + match = true; + matching_indices = std::pair(0, 2); + } + if (operand_1.location() == operand_2.location()) + { + match = true; + matching_indices = std::pair(1, 2); + } + } + if (match == false) + { + newOrder.push_back(instr); + } + else + { + std::cout << "Duplicate input variable detected" << std::endl; + auto copy_instr = pisa::instruction::Copy().create(); + copy_instr->setPMD(instr->PMD()); + copy_instr->setResidual(instr->residual()); + copy_instr->addInputOperand(instr->getInputOperand(matching_indices.second)); + auto output_operand = instr->getInputOperand(matching_indices.second); + output_operand.setLocation("copyA" + output_operand.location()); + copy_instr->addOutputOperand(output_operand); + newOrder.push_back(copy_instr); + instr->getInputOperand(matching_indices.second).setLocation(output_operand.location()); + newOrder.push_back(instr); + } + } + instr_order = newOrder; + return; +} + +void PISAGraphOptimizer::nodeLocklist(graph::NetworkNode &node, graph::Graph &p_isa_graph) +{ + auto p_isa_graph_node = p_isa_graph.getNode(node.GetId()); + for (int x = 0; x < p_isa_graph_node.GetOutDeg(); x++) + { + auto target_register = p_isa_graph.getNode(p_isa_graph_node.GetOutNId(x)); + //If target register is an output, don't touch it, if not, rename it + if (target_register.GetOutDeg() == 0 || node.GetDat().instruction->Name() == pisa::instruction::Mac().baseName) + { + rename_lock_list[target_register.GetDat().label] = true; + //std::cout << "Adjusting " << target_register.GetDat().label << " to " << "uid_" + std::to_string(unique_counter++) + "_" << target_register.GetDat().label << std::endl; + //target_register.GetDat().label = "uid_" + std::to_string(unique_counter++) + "_" + target_register.GetDat().label; + } + } +} + +void PISAGraphOptimizer::nodeVariableAdjustment(graph::NetworkNode &node, graph::Graph &p_isa_graph) +{ + //node.GetInDeg() + //Adjust outputs only + auto p_isa_graph_node = p_isa_graph.getNode(node.GetId()); + for (int x = 0; x < p_isa_graph_node.GetOutDeg(); x++) + { + auto target_register = p_isa_graph.getNode(p_isa_graph_node.GetOutNId(x)); + //If target register is an output, don't touch it, if not, rename it + if (/*target_register.GetOutDeg() > 0*/ rename_lock_list.count(target_register.GetDat().label) == 0) + { + std::cout << "Adjusting " << target_register.GetDat().label << " to " + << "uid_" + std::to_string(unique_counter++) + "_" << target_register.GetDat().label << std::endl; + target_register.GetDat().label = "uid_" + std::to_string(unique_counter++) + "_" + target_register.GetDat().label; + } + } +} + +void PISAGraphOptimizer::nodeMACVariableAdjustment(graph::NetworkNode &node, graph::Graph &p_isa_graph) +{ + + // + auto p_isa_graph_node = p_isa_graph.getNode(node.GetId()); + //for(int x = 0; x < p_isa_graph_node.GetOutDeg(); x++) { + if (p_isa_graph_node.GetDat().instruction->Name() == pisa::instruction::Mac().baseName) + { + auto target_register_input_reg = p_isa_graph.getNode(p_isa_graph_node.GetInNId(0)); + auto target_register_output_reg = p_isa_graph.getNode(p_isa_graph_node.GetOutNId(0)); + //If target register is an output, don't touch it, if not, rename it + //if(target_register.GetOutDeg() > 0) { + std::cout << "Adjusting Mac Variable registers" << target_register_input_reg.GetDat().label << " to " << target_register_output_reg.GetDat().label << std::endl; + + target_register_input_reg.GetDat().label = target_register_output_reg.GetDat().label; + //} + } + //} +} + +void PISAGraphOptimizer::nodeInstructionAdjustment(graph::NetworkNode &node, graph::Graph &p_isa_graph) +{ + // + auto p_isa_graph_node = p_isa_graph.getNode(node.GetId()); + // std::cout << "In degree:" + if (p_isa_graph_node.GetDat().instruction->Name() == pisa::instruction::Muli().baseName) + { + std::cout << "Muli instruction" << std::endl; + auto input_node_0 = p_isa_graph.getNode(p_isa_graph_node.GetInNId(0)).GetDat(); + auto input_node_1 = p_isa_graph.getNode(p_isa_graph_node.GetInNId(1)).GetDat(); + + auto &operand_0 = node.GetDat().instruction->getInputOperand(0); + auto &operand_1 = node.GetDat().instruction->getInputOperand(1); + + std::cout << "Input label" << 0 << ": " << input_node_0.label << " Immediate: " << operand_0.immediate() << std::endl; + std::cout << "Input label" << 1 << ": " << input_node_1.label << " Immediate: " << operand_1.immediate() << std::endl; + if (input_node_0.type == graph::IMMEDIATE) + { + operand_0.setLocation(input_node_1.label); + operand_1.setLocation(input_node_0.label); + } + else + { + operand_0.setLocation(input_node_0.label); + operand_1.setLocation(input_node_1.label); + } + } + else if (p_isa_graph_node.GetDat().instruction->Name() == pisa::instruction::Mac().baseName) + { + std::cout << "Mac instruction" << std::endl; + auto input_node_0 = p_isa_graph.getNode(p_isa_graph_node.GetInNId(0)).GetDat(); + auto input_node_1 = p_isa_graph.getNode(p_isa_graph_node.GetInNId(1)).GetDat(); + auto input_node_2 = p_isa_graph.getNode(p_isa_graph_node.GetInNId(2)).GetDat(); + + auto output_node_0 = p_isa_graph.getNode(p_isa_graph_node.GetOutNId(0)).GetDat(); + + auto &operand_0 = node.GetDat().instruction->getInputOperand(0); + auto &operand_1 = node.GetDat().instruction->getInputOperand(1); + auto &operand_2 = node.GetDat().instruction->getInputOperand(2); + + auto &output_operand_0 = node.GetDat().instruction->getOutputOperand(0); + + output_operand_0.setLocation(output_node_0.label); + if (output_node_0.label == input_node_0.label) + { + operand_0.setLocation(input_node_0.label); + operand_1.setLocation(input_node_1.label); + operand_2.setLocation(input_node_2.label); + } + else if (output_node_0.label == input_node_1.label) + { + operand_0.setLocation(input_node_1.label); + operand_1.setLocation(input_node_0.label); + operand_2.setLocation(input_node_2.label); + } + else if (output_node_0.label == input_node_2.label) + { + operand_0.setLocation(input_node_2.label); + operand_1.setLocation(input_node_0.label); + operand_2.setLocation(input_node_1.label); + } + else + { + throw std::runtime_error("No match between input and output registers, MAC instruction no valid output!"); + } + + // std::cout << "Input label" << 0 << ": " << input_node_0.label << " Immediate: " << operand_0.immediate() <getInputOperand(x).immediate() << std::endl; + node.GetDat().instruction->getInputOperand(x).setLocation(input_label); + } + } + for (int x = 0; x < p_isa_graph_node.GetOutDeg(); x++) + { + auto input_label = p_isa_graph.getNode(p_isa_graph_node.GetOutNId(x)).GetDat().label; + std::cout << "Output label" << x << ": " << input_label << std::endl; + node.GetDat().instruction->getOutputOperand(x).setLocation(input_label); + } + + return; +} diff --git a/p-isa_tools/program_mapper/p_isa/pisa_graph_optimizer.h b/p-isa_tools/program_mapper/p_isa/pisa_graph_optimizer.h new file mode 100644 index 00000000..f6c6a111 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/pisa_graph_optimizer.h @@ -0,0 +1,31 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include + +class PISAGraphOptimizer +{ +public: + PISAGraphOptimizer(); + std::vector generateInstructionStreamFromGraph(graph::Graph &p_isa_graph, bool fixed_order, std::vector instr_order); + + void isolateGraphVariables(graph::Graph &p_isa_graph, std::vector>> &layers); + + void applyDuplicateInputVariableSeparation(std::vector &instr_order); + + void nodeLocklist(graph::NetworkNode &node, graph::Graph &p_isa_graph); + void nodeVariableAdjustment(graph::NetworkNode &node, graph::Graph &p_isa_graph); + void nodeMACVariableAdjustment(graph::NetworkNode &node, graph::Graph &p_isa_graph); + + void nodeInstructionAdjustment(graph::NetworkNode &node, graph::Graph &p_isa_graph); + //void setNodeHeights(graph::Graph& p_isa_graph); + //int getNodeHeight(graph::NetworkNode& node, graph::Graph p_isa_graph) + + int unique_counter = 1; + bool perform_variable_isolation = false; + std::map rename_lock_list; +}; diff --git a/p-isa_tools/program_mapper/p_isa/pisa_test_generator.cpp b/p-isa_tools/program_mapper/p_isa/pisa_test_generator.cpp new file mode 100644 index 00000000..29b3e54a --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/pisa_test_generator.cpp @@ -0,0 +1,231 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include "pisa_test_generator.h" +#include "functional_modeler/data_handlers/json_data_handler.h" +#include "functional_modeler/pisa_runtime/pisaprogramruntime.h" +#include +#include +#include +#include +#include + +PisaTestGenerator::PisaTestGenerator() +{ +} + +void PisaTestGenerator::populateCalculatedOutputResults(std::vector instructions, json &input) +{ + PISAProgramRuntime evaluator; + //Setup data + auto json_data = JSONDataHandler(input); + std::vector modulus_chain = json_data.getModulusChain(); + auto trace_ntt_twiddle_factors = json_data.getNTTTwiddleFactors(); + auto trace_intt_twiddle_factors = json_data.getINTTTwiddleFactors(); + + auto trace_inputs = json_data.getAllInputs(); + auto trace_immediates = json_data.getAllimmediatesAsVec(1); + + evaluator.setModulusChain(modulus_chain); + auto chain = evaluator.getModulusChain(); + + evaluator.setNTTTwiddleFactors(trace_ntt_twiddle_factors); + evaluator.setINTTTwiddleFactors(trace_intt_twiddle_factors); + + evaluator.setParamMemoryToMultiRegisterDeviceMemory(trace_inputs); + evaluator.setImmediatesToMultiRegisterDeviceMemory(trace_immediates); + + evaluator.executeProgram(instructions); + + auto trace_outputs = json_data.getAllOutputs(); + auto outputs = input["output"].items(); + + for (const auto &output : outputs) + { + auto result = evaluator.getParamMemoryFromMultiRegisterDeviceMemory(output.key()); + + for (int x = 0; x < output.value().size(); x++) + { + input["output"][output.key()][x] = result.second[x]; + } + } + return; +} + +json PisaTestGenerator::generateJSONForGraph(graph::Graph p_isa_graph, pisa::testgenerator::InputGenerationMode gen_mode, unsigned int random_seed) +{ + json new_json; + auto inputs = p_isa_graph.getInputNodes(true, false, false); + auto immediates = p_isa_graph.getInputNodes(false, true, false); + auto outputs = p_isa_graph.getOutputNodes(); + + for (auto &input : inputs) + { + const std::string key = input.GetDat().label; + + for (int x = 0; x < block_size; x++) + { + if (gen_mode == pisa::testgenerator::InputGenerationMode::SINGLE_ONE) + { + new_json["input"][key][x] = (x == 0) ? 1 : 0; + } + else if (gen_mode == pisa::testgenerator::InputGenerationMode::ALL_ONES) + { + new_json["input"][key][x] = 1; + } + else if (gen_mode == pisa::testgenerator::InputGenerationMode::ASCENDING_FROM_ZERO) + { + new_json["input"][key][x] = x; + } + else if (gen_mode == pisa::testgenerator::InputGenerationMode::ONE_RANDOM) + { + new_json["input"][key][x] = (x == 0) ? rand() % modulus_value : 0; + } + else if (gen_mode == pisa::testgenerator::InputGenerationMode::ALL_RANDOM) + { + new_json["input"][key][x] = rand_r(&random_seed) % modulus_value; + } + } + } + + for (auto &output : outputs) + { + const std::string key = output.GetDat().label; + + for (int x = 0; x < block_size; x++) + { + new_json["output"][key][x] = 0; + } + } + + addMetaDataInformation(new_json); + + for (auto &immediate : immediates) + { + const std::string key = immediate.GetDat().label; + + int register_size = 1; + + for (int x = 0; x < register_size; x++) + { + new_json["metadata"]["immediate"][key] = 1; + } + } + + convertPolyRnsChunkToPolyRns(new_json); + + return new_json; +} + +int PisaTestGenerator::findMaxRNSNumber(json &input_json) +{ + int max_rns = 0; + + auto inputs = input_json.find("input"); + + for (auto input : inputs->items()) + { + auto label = input.key(); + int block_divider = label.find_last_of('_'); + auto block_removed = label.substr(0, block_divider); + int rns_divider = block_removed.find_last_of('_'); + std::string rns_val_str = block_removed.substr(rns_divider + 1); + int rns_val = std::atoi(rns_val_str.c_str()); + max_rns = std::max(rns_val, max_rns); + } + + return max_rns + 1; +} + +void PisaTestGenerator::addMetaDataInformation(json &input_json, int rns_num) +{ + input_json["metadata"]["scheme"] = "custom"; + for (int x = 0; x < rns_num; x++) + { + input_json["metadata"]["RNS_modulus"][x] = modulus_value; + } + + for (int x = 0; x < rns_num; x++) + { + for (int y = 0; y < block_size; y++) + { + input_json["metadata"]["twiddle"]["ntt"][x][y] = 1; + input_json["metadata"]["twiddle"]["intt"][x][y] = 1; + } + } + + //Add the default immediate values + input_json["metadata"]["immediate"]["iN"] = 1; + input_json["metadata"]["immediate"]["iN_0"] = 1; + input_json["metadata"]["immediate"]["iN_1"] = 1; + input_json["metadata"]["immediate"]["iN_2"] = 1; + input_json["metadata"]["immediate"]["R2_0"] = 1; + input_json["metadata"]["immediate"]["R2_1"] = 1; + input_json["metadata"]["immediate"]["R2_2"] = 1; + input_json["metadata"]["immediate"]["one"] = 1; + input_json["metadata"]["immediate"]["pinv_q_0"] = 1; + input_json["metadata"]["immediate"]["pinv_q_1"] = 1; + input_json["metadata"]["immediate"]["t_inverse_mod_p_0"] = 1; + input_json["metadata"]["immediate"]["t_0"] = 1; + input_json["metadata"]["immediate"]["t_1"] = 1; + input_json["metadata"]["immediate"]["t_2"] = 1; + + return; +} + +void PisaTestGenerator::addMetaDataInformation(json &input_json) +{ + auto rns_num = findMaxRNSNumber(input_json); + addMetaDataInformation(input_json, rns_num); +} + +void PisaTestGenerator::convertPolyRnsChunkToPolyRns(json &input_json) +{ + convertPolyRnsChunkToPolyRnsHelper(input_json["input"]); + convertPolyRnsChunkToPolyRnsHelper(input_json["output"]); +} + +void PisaTestGenerator::convertPolyRnsChunkToPolyRnsHelper(json &input_json) +{ + std::map>> collections; + for (auto input : input_json.items()) + { + auto label = input.key(); + int block_divider = label.find_last_of('_'); + auto block_removed = label.substr(0, block_divider); + int block_number = std::stoi(label.substr(block_divider + 1)); + + // This creates my ordered vectors + //std::cout << "Adding to collection: " << block_removed << " , " << label << " , " << block_number << std::endl; + collections[block_removed].push_back(std::pair(label, block_number)); + } + for (auto &collection : collections) + { + //Sort vector based on block number + std::sort(collection.second.begin(), collection.second.end(), [](const std::pair &a, const std::pair &b) { return a.second < b.second; }); + + for (auto chunk : collection.second) + { + for (auto block_value : input_json[chunk.first].items()) + { + //std::cout << "Block V" + input_json[collection.first][input_json[collection.first].size()] = block_value.value(); + }; + } + //Remove block item from json once data is transferred + int items_to_remove = collection.second.size(); + for (int x = 0; x < items_to_remove; x++) + { + auto chunk_val = input_json.find(collection.second[x].first); + input_json.erase(chunk_val); + } + } +} + +void PisaTestGenerator::writeJSON(json input_json, std::string file_name) +{ + auto serialized_json = input_json.dump(1, ' ', true); + std::ofstream output; + output.open(file_name); + output << serialized_json; +} diff --git a/p-isa_tools/program_mapper/p_isa/pisa_test_generator.h b/p-isa_tools/program_mapper/p_isa/pisa_test_generator.h new file mode 100644 index 00000000..8e5e289b --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/pisa_test_generator.h @@ -0,0 +1,53 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/tests/pisa_instruction_tests.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests.h" +#include +#include +#include +#include +#include + +using json = nlohmann::json; +namespace pisa::testgenerator { + +enum class InputGenerationMode +{ + SINGLE_ONE, + ALL_ONES, + ASCENDING_FROM_ZERO, + ONE_RANDOM, + ALL_RANDOM +}; + +static std::string AvailableGenerationModesStr() +{ + std::string genModes = "( SINGLE_ONE , ALL_ONES , ASCENDING_FROM_ZERO , ONE_RANDOM , ALL_RANDOM )"; + return genModes; +} +} // namespace pisa::testgenerator + +class PisaTestGenerator +{ +public: + PisaTestGenerator(); + + graph::Graph generateGraphFromProgramTrace(ProgramTrace program_trace); + graph::Graph generateGraphFromHEOperationTrace(std::vector program_trace); + graph::Graph generateGraphFromPISAInstructions(std::vector instructions); + void populateCalculatedOutputResults(std::vector instructions, json &input); + json generateJSONForGraph(graph::Graph p_isa_graph, pisa::testgenerator::InputGenerationMode gen_mode = pisa::testgenerator::InputGenerationMode::SINGLE_ONE, unsigned int random_seed = 0); + int findMaxRNSNumber(json &input_json); + void addMetaDataInformation(json &input_json, int RNS_NUM); + void addMetaDataInformation(json &input_json); + void convertPolyRnsChunkToPolyRns(json &input_json); + void convertPolyRnsChunkToPolyRnsHelper(json &input_json); + void writeJSON(json input_json, std::string file_name); + json trace_file; + + int block_size = 8192; + int modulus_value = 32684; +}; diff --git a/p-isa_tools/program_mapper/p_isa/pisakernel.cpp b/p-isa_tools/program_mapper/p_isa/pisakernel.cpp new file mode 100644 index 00000000..7bae3b54 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/pisakernel.cpp @@ -0,0 +1,509 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include +#include +#include + +#include +#include +#include + +namespace pisa::kernel { + +const std::vector &PISAKernel::getOutput_names() const +{ + return output_names; +} + +const std::vector &PISAKernel::getInput_names() const +{ + return input_names; +} + +static unsigned int global_kernel_id_counter = 0; + +} // namespace pisa::kernel + +// Helper function for generating the kernel generator input given an operation +inline std::string genKernInput(const pisa::poly::PolyOperation &op) +{ + /* + * CONTEXT SCHEME poly_order key_rns current_rns + * DATA symbol num_parts + * OPNAME output input(s) + */ + std::ostringstream input; + // CONTEXT + input << "CONTEXT " << toString(op.parentProgram()->scheme()) + << " " << op.parentProgram()->getPolyModulusDegree() + << " " << op.parentProgram()->getKeyRns() + << " " << op.getInputOperand(0).num_of_rns_terms + << "\n"; + + // DATA + // NOTE: CipherDegree is tied to HEOperation not Operand + auto operation = op.Name(); + std::vector inputs = op.getInputLocations(); + std::vector output = op.getOutputLocations(); + + //#TODO Naming is switched to use generic sequential naming allowing program mapper to control the final naming + //This allows generic kernels to be generated effectively for the cache, while passing names worked well for single + //operations it creates issues for multi ops. A more robust fix with runtime control on which naming control scheme to use may be desirable. + for (int i = 0; i < op.numOutputOperands(); ++i) + { + input << "DATA " << /*output[i].register_name*/ "output" << i + << " " << output[i].num_of_polynomials << "\n"; + } + + for (int i = 0; i < op.numInputOperands(); ++i) + { + input << "DATA " << /*inputs[i].register_name*/ "input" << i + << " " << inputs[i].num_of_polynomials << "\n"; + } + + // OP + input << std::uppercase << op.Name() << std::nouppercase; + for (int i = 0; i < op.numOutputOperands(); ++i) + { + input << " " + << "output" << i /*output[i].register_name*/; + } + for (int i = 0; i < op.numInputOperands(); ++i) + { + input << " " + << "input" << i /*inputs[i].register_name*/; + } + + return input.str(); +} + +// Helper function creating kernel filepath +inline std::filesystem::path createKernelFilepath(const pisa::poly::PolyOperation &op, const pisa::kernel::Cache &kernel_cache) +{ + std::ostringstream kernel_file_name; + kernel_file_name << toString(op.parentProgram()->scheme()) << "_" + << op.Name() << "_" + << op.parentProgram()->getPolyModulusDegree() << "_" + << op.getInputLocations().front().num_of_polynomials << "_" + << op.getInputLocations().front().num_of_rns_terms << ".csv"; + + return std::filesystem::path(kernel_cache.getDirname()) / kernel_file_name.str(); +} + +// Selection function for choosing between legacy and new kernel generators +pisa::kernel::PISAKernel *pisa::kernel::PISAKernel::create(std::string he_op_generator, pisa::poly::PolyOperation *op, const pisa::kernel::Cache &kernel_cache, bool verbose, bool new_kerngen, std::string kern_library) +{ + if (new_kerngen) + { + return create_new(he_op_generator, op, kernel_cache, verbose); + } + if (kern_library == "CSV") + { + return create_legacy(he_op_generator, op, kernel_cache, verbose); + } + else if (kern_library == "HDF") + { + return createHECDataFormats(he_op_generator, op, kernel_cache, verbose); + } + else + throw std::runtime_error("Invalid kernel library"); +} + +pisa::kernel::PISAKernel *pisa::kernel::PISAKernel::create_new(std::string he_op_generator, pisa::poly::PolyOperation *op, const pisa::kernel::Cache &kernel_cache, bool verbose) +{ + PISAKernel *kernel = PISAKernel().createInstance(); + kernel->name = op->Name(); + kernel->kernel_id = pisa::kernel::global_kernel_id_counter++; + + const auto kernel_file_path = createKernelFilepath(*op, kernel_cache); + const std::string command_string = he_op_generator + " -q -l > " + kernel_file_path.c_str() + " <instructions = pisa::PISAParser::parse(kernel_file_path); + kernel->mapped_instructions.resize(kernel->instructions.size()); + for (int x = 0; x < kernel->instructions.size(); x++) + kernel->mapped_instructions[x] = new pisa::PISAInstruction(*kernel->instructions[x]); + + return kernel; +} + +pisa::kernel::PISAKernel *pisa::kernel::PISAKernel::create_legacy(std::string he_op_generator, pisa::poly::PolyOperation *op, const pisa::kernel::Cache &kernel_cache, bool verbose) +{ + PISAKernel *kernel = PISAKernel().createInstance(); + kernel->name = op->Name(); + kernel->kernel_id = pisa::kernel::global_kernel_id_counter++; + + std::ostringstream params; + params << toString(op->parentProgram()->scheme(), true) << " " + << op->Name() << " " + << std::to_string(op->parentProgram()->getPolyModulusDegree()) << " " + << op->getInputOperand(0).num_of_rns_terms; + + auto kernel_file_name = params.str() + "_" + std::to_string(op->getInputOperand(0).num_of_polynomials) + ".csv"; + std::replace(kernel_file_name.begin(), kernel_file_name.end(), ' ', '_'); + + // Add to correct calling params. + params << " " << (op->getInputOperand(0).num_of_rns_terms + 1); + + if (op->Name() == "add") + { + params << " " << op->getInputOperand(0).num_of_polynomials; + } + + if (op->Name() == "relin" || op->Name() == "rotate") + { + // If passing dnum (number of digits), need to also pass alpha (digit size) and k (size of the extended prime) + // dnum == rns and alpha/k are 1 as we're using rns-prime decomposition for key switching + int dnum = op->getInputOperand(0).num_of_rns_terms; + int alpha = op->parentProgram()->getAlpha() == 0 ? 1 : op->parentProgram()->getAlpha(); + int k = alpha; + params << " " << dnum << " " << alpha << " " << k; + } + + if (verbose) + { + std::cout << he_op_generator << " " << params.str() << std::endl; + } + + auto kernel_file_path = std::filesystem::path(kernel_cache.getDirname()) / kernel_file_name; + if (!kernel_cache.use_cache() || !std::filesystem::exists(kernel_file_path)) + { + std::string command_string = he_op_generator + " " + params.str() + " > " + kernel_file_path.c_str(); + std::cout << command_string << std::endl; + [[maybe_unused]] auto rc = system(command_string.c_str()); + } + + kernel->instructions = pisa::PISAParser::parse(kernel_file_path); + kernel->mapped_instructions.resize(kernel->instructions.size()); + for (int x = 0; x < kernel->instructions.size(); x++) + kernel->mapped_instructions[x] = new pisa::PISAInstruction(*kernel->instructions[x]); + + return kernel; +} + +pisa::kernel::PISAKernel *pisa::kernel::PISAKernel::createHECDataFormats(std::string he_op_generator, pisa::poly::PolyOperation *op, const pisa::kernel::Cache &kernel_cache, bool verbose) +{ + PISAKernel *kernel = PISAKernel().createInstance(); + kernel->name = op->Name(); + kernel->kernel_id = pisa::kernel::global_kernel_id_counter++; + + std::ostringstream params; + params << toString(op->parentProgram()->scheme(), true) << " " + << op->Name() << " " + << op->parentProgram()->getPolyModulusDegree() << " " + << op->getRnsTerms(); + + int key_rns_num = op->parentProgram()->getKeyRns(); + uint32_t q_size = op->parentProgram()->getQSize(); + int dnum = op->parentProgram()->getDNum(); + uint32_t alpha = op->parentProgram()->getAlpha(); + // k = alpha + int k = alpha; + + params << " " << key_rns_num; + + if (op->Name() == "relin") + { + // If passing dnum (number of digits), need to also pass alpha (digit size) and k (size of the extended prime) + // dnum == rns and alpha/k are 1 as we're using rns-prime decomposition for key switching + params << " " << dnum << " " << alpha << " " << k << " " << q_size; + } + else if (op->Name() == "add") + { + params << " " << op->getCipherDegree(); + } + else if (op->Name() == "rotate") + { + // If passing dnum (number of digits), need to also pass alpha (digit size) and k (size of the extended prime) + // dnum == rns and alpha/k are 1 as we're using rns-prime decomposition for key switching + params << " " << dnum << " " << alpha << " " << k; + params << " " << q_size << " " << op->getGaloisElt(); + } + else if (op->Name() == "rescale") + { + // qsize is required for Dataformats CKKS + params << " " << q_size; + } + if (verbose) + { + std::cout << he_op_generator << " " << params.str() << std::endl; + } + + // make kernel_file_name use full params + auto kernel_file_name = params.str() + "_" + std::to_string(op->getCipherDegree()) + ".csv"; + std::replace(kernel_file_name.begin(), kernel_file_name.end(), ' ', '_'); + + auto kernel_file_path = std::filesystem::path(kernel_cache.getDirname()) / kernel_file_name; + // debugging + if (!kernel_cache.use_cache() || !std::filesystem::exists(kernel_file_path)) + { + std::string command_string = he_op_generator + " " + params.str() + " > " + kernel_file_path.c_str(); + std::cout << command_string << std::endl; + [[maybe_unused]] auto rc = system(command_string.c_str()); + } + + kernel->instructions = pisa::PISAParser::parse(kernel_file_path); + kernel->mapped_instructions.resize(kernel->instructions.size()); + for (int x = 0; x < kernel->instructions.size(); x++) + kernel->mapped_instructions[x] = new pisa::PISAInstruction(*kernel->instructions[x]); + + return kernel; +} + +bool pisa::kernel::PISAKernel::enableNamespace() const +{ + return enable_namespace; +} + +void pisa::kernel::PISAKernel::setEnableNamespace(bool newNamespace_internal_variables) +{ + enable_namespace = newNamespace_internal_variables; +} + +std::string pisa::kernel::PISAKernel::registerNameRoot(const std::string ®_name) +{ + int size = reg_name.find('_', 0); + return reg_name.substr(0, size); +} + +std::vector pisa::kernel::PISAKernel::nonRepeatingRootsNode(std::vector> &xputs) +{ + std::unordered_set non_repeat_roots; + std::vector non_repeat_roots_insertion_ordered; + for (auto xput : xputs) + { + auto root = registerNameRoot(xput.GetDat().label); + if (non_repeat_roots.count(root) == 0) + { + non_repeat_roots_insertion_ordered.push_back(root); + } + non_repeat_roots.insert(root); + } + return non_repeat_roots_insertion_ordered; +} + +void pisa::kernel::PISAKernel::updateInput(int index, std::string new_name) +{ + std::string old_name = input_names[index]; + + for (auto &i : instructions) + { + for (int x = 0; x < i->numInputOperands(); x++) + { + auto name_root = i->getInputOperand(x).locationRoot(); + if (name_root == old_name) + { + i->getInputOperand(x).setLocationRoot(new_name); + } + } + } + input_names[index] = new_name; +} + +void pisa::kernel::PISAKernel::updateOutput(int index, std::string new_name) +{ + std::string old_name = output_names[index]; + + for (auto &i : instructions) + { + for (int x = 0; x < i->numOutputOperands(); x++) + { + auto name_root = i->getOutputOperand(x).locationRoot(); + if (name_root == old_name) + { + i->getOutputOperand(x).setLocationRoot(new_name); + } + } + } + output_names[index] = new_name; +} + +void pisa::kernel::PISAKernel::updateSymbols(bool verbose) +{ + try + { + if (internal_map.size() == 0) + createInternalVariableMap(); + + for (int i = 0; i < instructions.size(); i++) + { + for (int x = 0; x < instructions[i]->numInputOperands(); x++) + { + auto name_root = instructions[i]->getInputOperand(x).locationRoot(); + if (naming_map.count(name_root) == 1) + { + auto value = naming_map[name_root]; + mapped_instructions[i]->getInputOperand(x).setLocationRoot(value); + if (verbose) + std::cout << "Mapped: " << name_root << "->" << value << std::endl; + } + else if (enable_namespace) + { + auto name_loc = internal_map[instructions[i]->getInputOperand(x).location()]; + mapped_instructions[i]->getInputOperand(x).setLocation(name_loc); + if (verbose) + std::cout << "Mapped: " << instructions[i]->getInputOperand(x).location() << "->" << name_loc << std::endl; + } + } + + for (int x = 0; x < instructions[i]->numOutputOperands(); x++) + { + auto name_root = instructions[i]->getOutputOperand(x).locationRoot(); + if (naming_map.count(name_root) == 1) + { + auto value = naming_map[name_root]; + mapped_instructions[i]->getOutputOperand(x).setLocationRoot(value); + if (verbose) + std::cout << "Mapped: " << name_root << "->" << value << std::endl; + } + else if (enable_namespace) + { + auto name_loc = internal_map[instructions[i]->getOutputOperand(x).location()]; + mapped_instructions[i]->getOutputOperand(x).setLocation(name_loc); + if (verbose) + std::cout << "Mapped: " << instructions[i]->getOutputOperand(x).location() << "->" << name_loc << std::endl; + } + } + map_dirty = false; + } + } + catch (...) + { + throw; + } +} + +void pisa::kernel::PISAKernel::setImmediate(const std::string &key, const std::string &value) +{ + immediate_map[key] = value; + map_dirty = true; +} + +void pisa::kernel::PISAKernel::mapInput(int index, const std::string &new_name) +{ + naming_map[input_names[index]] = new_name; + map_dirty = true; +} + +void pisa::kernel::PISAKernel::mapImmediate(int index, const std::string &new_name) +{ + naming_map[immediate_names[index]] = new_name; + map_dirty = true; +} + +void pisa::kernel::PISAKernel::mapOutput(int index, const std::string &new_name) +{ + naming_map[output_names[index]] = new_name; + map_dirty = true; +} + +const std::vector &pisa::kernel::PISAKernel::getMappedInstructions() +{ + if (map_dirty) + updateSymbols(false); + + return mapped_instructions; +} + +void pisa::kernel::PISAKernel::createInternalVariableMap() +{ + + for (const auto &meta : immediate_names) + { + naming_map[meta] = meta; + internal_map[meta] = meta; + } + for (auto meta : immediate_map) + { + naming_map[meta.first] = meta.second; + internal_map[meta.first] = meta.second; + } + + for (int i = 0; i < instructions.size(); i++) + { + for (int x = 0; x < instructions[i]->numInputOperands(); x++) + { + auto name_root = instructions[i]->getInputOperand(x).locationRoot(); + auto name_loc = instructions[i]->getInputOperand(x).location(); + //if(name_root.size() > 0) { + if (naming_map.count(name_root) == 0 && immediate_map.count(name_loc) == 0 && internal_map.count(name_loc) == 0) + { + internal_map[name_loc] = "internal" + name + std::to_string(kernel_id) + "NS_" + name_loc; + } + } + + for (int x = 0; x < instructions[i]->numOutputOperands(); x++) + { + auto name_loc = instructions[i]->getOutputOperand(x).location(); + auto name_root = instructions[i]->getOutputOperand(x).locationRoot(); + //if(name_root.size() > 0) { + if (naming_map.count(name_root) == 0 && immediate_map.count(name_loc) == 0 && internal_map.count(name_loc) == 0) + { + internal_map[name_loc] = "internal" + name + std::to_string(kernel_id) + "NS_" + name_loc; + } + } + } +} + +bool containsString(const std::string &value, const std::string &substring) +{ + return value.find(substring) != std::string::npos; +} + +void pisa::kernel::PISAKernel::determineVariableNamingViaGraph() +{ + auto instruction_graph = graph::Graph::createGraph(instructions); + auto inputs = instruction_graph.getInputNodes(true, false, false); + auto outputs = instruction_graph.getOutputNodes(); + auto immediates = instruction_graph.getInputNodes(false, true, false); + + auto non_repeat_inputs = nonRepeatingRootsNode(inputs); + auto non_repeat_outputs = nonRepeatingRootsNode(outputs); + for (auto &input : non_repeat_inputs) + { + input_names.push_back(input); + naming_map[input] = input; + } + // temporary solution for two corner cases: + // Two inputs are needed but "d" is parsed before "c" causing input1 and input2 being swapped + // Trace has TWO inputs but only "d" is used, it will cause mapInput function to crash + // NOTE: This could still lead to an issue for second corner case where only "c" is used but trace has two inputs + //if (input_names.size() == 1 && input_names.front() == "d") + // input_names.push_back("c"); + // std::sort(input_names.begin(), input_names.end(), comp); + // Attempt #2, just sort inputs if the label contains "input" + // Leave the rest of the list alone + // TODO: Generalize Function + auto comp = [](const std::string &a, const std::string &b) { + if (containsString(a, "input") && containsString(b, "input")) + return a < b; + else + return false; + }; + std::sort(input_names.begin(), input_names.end(), comp); + + for (auto &output : non_repeat_outputs) + { + output_names.push_back(output); + naming_map[output] = output; + } + std::sort(output_names.begin(), output_names.end()); + + for (auto &immediate : immediates) + { + if (immediate_map.count(immediate.GetDat().label) == 0) + { + immediate_names.push_back(immediate.GetDat().label); + } + immediate_map[immediate.GetDat().label] = immediate.GetDat().label; + } +} diff --git a/p-isa_tools/program_mapper/p_isa/pisakernel.h b/p-isa_tools/program_mapper/p_isa/pisakernel.h new file mode 100644 index 00000000..d094789c --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/pisakernel.h @@ -0,0 +1,112 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include +#include +#include + +#include +#include +#include + +namespace pisa::kernel { + +class Cache +{ +public: + Cache() = delete; + Cache(const std::string &dirname, bool use_cache = true, bool remove_cache = false) : + _dirname(dirname), _use_cache(use_cache), _remove_cache(remove_cache) + { + std::filesystem::create_directory(_dirname); + } + + std::string getDirname() const { return _dirname; } + bool use_cache() const { return _use_cache; } + + ~Cache() + { + if (_remove_cache) + { + std::filesystem::remove_all(_dirname); + } + } + +private: + std::string _dirname; + bool _use_cache; + bool _remove_cache; +}; + +class PISAKernel +{ +public: + PISAKernel() = default; + PISAKernel(const std::vector &input_names_, + const std::vector &output_names_, + const std::vector &immediate_names_ = {}) : + input_names(input_names_), + output_names(output_names_), + immediate_names(immediate_names_) + { + } + + static PISAKernel *create(std::string he_op_generator, pisa::poly::PolyOperation *op, const Cache &kernel_cache, bool verbose = false, bool new_kerngen = false, std::string kern_library = "CSV"); + + static PISAKernel *create_legacy(std::string he_op_generator, pisa::poly::PolyOperation *op, const Cache &kernel_cache, bool verbose = false); + static PISAKernel *create_new(std::string he_op_generator, pisa::poly::PolyOperation *op, const Cache &kernel_cache, bool verbose = false); + + static PISAKernel *createHECDataFormats(std::string he_op_generator, pisa::poly::PolyOperation *op, const Cache &kernel_cache, bool verbose = false); + + virtual PISAKernel *createInstance() { return new PISAKernel(); } + std::vector instructions; + std::vector mapped_instructions; + + std::vector input_names; + std::vector output_names; + std::vector immediate_names; + + bool map_dirty = true; + + void updateInput(int index, std::string new_name); + void updateOutput(int index, std::string new_name); + void updateSymbols(bool verbose = false); + + std::map naming_map; + std::map immediate_map; + + void setImmediate(const std::string &, const std::string &); + void mapInput(int index, const std::string &new_name); + void mapImmediate(int index, const std::string &new_name); + void mapOutput(int index, const std::string &new_name); + const std::vector &getMappedInstructions(); + + void createInternalVariableMap(); + void determineVariableNamingViaGraph(); + std::map internal_map; + + std::string name; + unsigned int kernel_id; + + bool enable_namespace = true; + bool enableNamespace() const; + void setEnableNamespace(bool newNamespace_internal_variables); + + /** + * @brief registerNameRoot Attempts to split a register name removing the RNS and block terms. + * @todo Need to account for outlier cases when naming doesn't match + * @param location + * @return + */ + std::string registerNameRoot(const std::string ®_name); + + std::vector nonRepeatingRootsNode(std::vector> &xputs); + const std::vector &getInput_names() const; + const std::vector &getOutput_names() const; +}; + +} // namespace pisa::kernel diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests.h new file mode 100644 index 00000000..01b4f4a5 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests.h @@ -0,0 +1,28 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "pisa_instruction_tests/add_instruction_test.h" +#include "pisa_instruction_tests/copy_instruction_test.h" +#include "pisa_instruction_tests/intt_instruction_test.h" +#include "pisa_instruction_tests/mac_instruction_test.h" +#include "pisa_instruction_tests/maci_instruction_test.h" +#include "pisa_instruction_tests/mul_instruction_test.h" +#include "pisa_instruction_tests/muli_instruction_test.h" +#include "pisa_instruction_tests/ntt_instruction_test.h" +#include "pisa_instruction_tests/pisa_instruction_test.h" +#include "pisa_instruction_tests/random_instr_stream_instruction_test.h" +#include "pisa_instruction_tests/sub_instruction_test.h" +#include + +static std::map pisa_instruction_tests = { { addInstructionTest::operationName(), new addInstructionTest() }, + { mulInstructionTest::operationName(), new mulInstructionTest() }, + { muliInstructionTest::operationName(), new muliInstructionTest() }, + { macInstructionTest::operationName(), new macInstructionTest() }, + { maciInstructionTest::operationName(), new maciInstructionTest() }, + { nttInstructionTest::operationName(), new nttInstructionTest() }, + { inttInstructionTest::operationName(), new inttInstructionTest() }, + { subInstructionTest::operationName(), new subInstructionTest() }, + { copyInstructionTest::operationName(), new copyInstructionTest() }, + { RandomStreamInstructionTest::operationName(), new RandomStreamInstructionTest() } }; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/add_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/add_instruction_test.h new file mode 100644 index 00000000..e7c4b096 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/add_instruction_test.h @@ -0,0 +1,50 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include + +class addInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName() { return "add_instruction"; } + + addInstructionTest() : + PisaInstructionTest() + { + configuration["Name"] = this->operationName(); + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (const auto &config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { "output0" }; + std::vector input_regs = { { "input0" }, { "input1" } }; + + auto add_instr = new pisa::instruction::Add(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); + + instruction_trace.push_back(add_instr); + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/copy_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/copy_instruction_test.h new file mode 100644 index 00000000..a9785a62 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/copy_instruction_test.h @@ -0,0 +1,48 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include + +class copyInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName() { return "copy_instruction"; } + + copyInstructionTest() : + PisaInstructionTest() + { + configuration["Name"] = this->operationName(); + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { "output0" }; + std::vector input_regs = { { "input0" } }; + + auto instr = new pisa::instruction::Copy(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"])); + + instruction_trace.push_back(instr); + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/intt_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/intt_instruction_test.h new file mode 100644 index 00000000..f486920a --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/intt_instruction_test.h @@ -0,0 +1,53 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include + +class inttInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName() { return "intt_instruction"; } + + inttInstructionTest() : + PisaInstructionTest() + { + configuration["Name"] = this->operationName(); + configuration["WParam"] = "w_0_0_0"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { { "output0" }, { "output1" } }; + std::vector input_regs = { { "input0" }, { "input1" } }; + + auto instr = new pisa::instruction::Intt(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(output_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::WParam(configuration["WParam"]), + std::stoi(configuration["RNS_INDEX"])); + + instruction_trace.push_back(instr); + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/mac_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/mac_instruction_test.h new file mode 100644 index 00000000..b94864b4 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/mac_instruction_test.h @@ -0,0 +1,50 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include + +class macInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName() { return "mac_instruction"; } + + macInstructionTest() : + PisaInstructionTest() + { + configuration["Name"] = this->operationName(); + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { "output0" }; + std::vector input_regs = { { "input1" }, { "input2" } }; + + auto mac_instr = new pisa::instruction::Mac(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); + + instruction_trace.push_back(mac_instr); + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/maci_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/maci_instruction_test.h new file mode 100644 index 00000000..8d10cdaf --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/maci_instruction_test.h @@ -0,0 +1,50 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include + +class maciInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName() { return "maci_instruction"; } + + maciInstructionTest() : + PisaInstructionTest() + { + configuration["Name"] = this->operationName(); + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { "output0" }; + std::vector input_regs = { { "input0" }, { "input1" } }; + + auto instr = new pisa::instruction::Maci(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); + + instruction_trace.push_back(instr); + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/mul_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/mul_instruction_test.h new file mode 100644 index 00000000..5e955bbe --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/mul_instruction_test.h @@ -0,0 +1,50 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include + +class mulInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName() { return "mul_instruction"; } + + mulInstructionTest() : + PisaInstructionTest() + { + configuration["Name"] = this->operationName(); + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { "output0" }; + std::vector input_regs = { { "input0" }, { "input1" } }; + + auto instr = new pisa::instruction::Mul(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); + + instruction_trace.push_back(instr); + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/muli_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/muli_instruction_test.h new file mode 100644 index 00000000..c2e65283 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/muli_instruction_test.h @@ -0,0 +1,50 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include + +class muliInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName() { return "muli_instruction"; } + + muliInstructionTest() : + PisaInstructionTest() + { + configuration["Name"] = this->operationName(); + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { "output0" }; + std::vector input_regs = { { "input0" }, { "input1" } }; + + auto instr = new pisa::instruction::Muli(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); + + instruction_trace.push_back(instr); + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/ntt_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/ntt_instruction_test.h new file mode 100644 index 00000000..7ffd8d2e --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/ntt_instruction_test.h @@ -0,0 +1,53 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include + +class nttInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName() { return "ntt_instruction"; } + + nttInstructionTest() : + PisaInstructionTest() + { + configuration["Name"] = this->operationName(); + configuration["WParam"] = "w_0_0_0"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { { "output0" }, { "output1" } }; + std::vector input_regs = { { "input0" }, { "input1" } }; + + auto instr = new pisa::instruction::Ntt(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(output_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::WParam(configuration["WParam"]), + std::stoi(configuration["RNS_INDEX"])); + + instruction_trace.push_back(instr); + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h new file mode 100644 index 00000000..9435296c --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h @@ -0,0 +1,52 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include +#include + +class PisaInstructionTest +{ + +public: + PisaInstructionTest() + { + configuration["Name"] = "Default"; + configuration["RNS_INDEX"] = "0"; + configuration["Poly_mod_log2"] = "14"; + configuration["Chunk_INDEX"] = "0"; + } + virtual void constructTest() + { + } + + const std::vector &getInstructionTrace() + { + try + { + if (created == true) + { + return instruction_trace; + } + else + throw std::runtime_error("Tried to use without creating first"); + } + catch (...) + { + throw; + } + } + + std::map &getConfiguration() + { + return configuration; + } + +protected: + std::vector instruction_trace; + bool created = false; + std::map configuration; +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/random_instr_graph_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/random_instr_graph_instruction_test.h new file mode 100644 index 00000000..9a57aa89 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/random_instr_graph_instruction_test.h @@ -0,0 +1,50 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include + +class mulInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName() { return "mul_instruction"; } + + mulInstructionTest() : + PisaInstructionTest() + { + configuration["Name"] = this->operationName(); + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { "output0" }; + std::vector input_regs = { { "input0" }, { "input1" } }; + + auto instr = new pisa::instruction::Mul(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); + + instruction_trace.push_back(instr); + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/random_instr_stream_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/random_instr_stream_instruction_test.h new file mode 100644 index 00000000..4ebbe4ef --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/random_instr_stream_instruction_test.h @@ -0,0 +1,253 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include +#include +#include + +class RandomStreamInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName(); + + RandomStreamInstructionTest(); + ; + + // PisaKernelTest interface +public: + void constructTest() override; + + pisa::PISAInstruction *createInstr(std::string op, std::string output, std::string input0, std::string input1); + pisa::PISAInstruction *createAddInstr(std::string output, std::string input0, std::string input1); + pisa::PISAInstruction *createMulInstr(std::string output, std::string input0, std::string input1); + pisa::PISAInstruction *createMacInstr(std::string output, std::string input0, std::string input1); + pisa::PISAInstruction *createMaciInstr(std::string output, std::string input0, std::string input1); + pisa::PISAInstruction *createMuliInstr(std::string output, std::string input0, std::string input1); + pisa::PISAInstruction *createSubInstr(std::string output, std::string input0, std::string input1); + pisa::PISAInstruction *createCopyInstr(std::string output, std::string input0); +}; + +inline std::string RandomStreamInstructionTest::operationName() +{ + return "random_stream_instruction"; +} + +inline RandomStreamInstructionTest::RandomStreamInstructionTest() : + PisaInstructionTest() +{ + configuration["Name"] = this->operationName(); + configuration["Intermediate_registers"] = "10"; + configuration["Add_ops"] = "5"; + configuration["Mul_ops"] = "0"; + configuration["Copy_ops"] = "0"; + configuration["Mac_ops"] = "0"; + configuration["Maci_ops"] = "0"; + configuration["Muli_ops"] = "0"; + configuration["Sub_ops"] = "0"; + configuration["Random_seed"] = "0"; +} + +inline void RandomStreamInstructionTest::constructTest() +{ + + unsigned int random_seed = std::stoi(configuration["Random_seed"]); + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { "output0" }; + std::vector intermediate_regs; + int intermediate_register_count = std::stoi(configuration["Intermediate_registers"]); + for (unsigned int x = 0; x < intermediate_register_count; x++) + { + intermediate_regs.push_back("intermediate" + std::to_string(x)); + } + std::vector input_regs = { { "input0" }, { "input1" } }; + + std::vector op_tokens; + + //Initialize all intermediates with 1 value + instruction_trace.push_back(createInstr("copy", intermediate_regs[0], input_regs[0], input_regs[0])); + for (unsigned int x = 1; x < intermediate_regs.size(); x++) + { + instruction_trace.push_back(createInstr("copy", intermediate_regs[x], intermediate_regs[x - 1], intermediate_regs[x - 1])); + } + + //Create a bucket of test tokens + //Add ops + int add_ops = std::stoi(configuration["Add_ops"]); + for (unsigned int x = 0; x < add_ops; x++) + { + op_tokens.push_back("add"); + } + int mul_ops = std::stoi(configuration["Mul_ops"]); + for (unsigned int x = 0; x < mul_ops; x++) + { + op_tokens.push_back("mul"); + } + int copy_ops = std::stoi(configuration["Copy_ops"]); + for (unsigned int x = 0; x < copy_ops; x++) + { + op_tokens.push_back("copy"); + } + int mac_ops = std::stoi(configuration["Mac_ops"]); + for (unsigned int x = 0; x < mac_ops; x++) + { + op_tokens.push_back("mac"); + } + int maci_ops = std::stoi(configuration["Maci_ops"]); + for (unsigned int x = 0; x < maci_ops; x++) + { + op_tokens.push_back("maci"); + } + int muli_ops = std::stoi(configuration["Muli_ops"]); + for (unsigned int x = 0; x < muli_ops; x++) + { + op_tokens.push_back("muli"); + } + int sub_ops = std::stoi(configuration["Sub_ops"]); + for (unsigned int x = 0; x < sub_ops; x++) + { + op_tokens.push_back("sub"); + } + + //Shuffle the bucket + std::random_shuffle(op_tokens.begin() + intermediate_regs.size(), op_tokens.end()); + + //Create first instruction + auto first_token = op_tokens.front(); + int output_reg = rand_r(&random_seed) % intermediate_regs.size(); + if (op_tokens.size() > 1) + { + instruction_trace.push_back(createInstr(first_token, intermediate_regs[output_reg], intermediate_regs.back(), input_regs[1])); + } + else + { + instruction_trace.push_back(createInstr(first_token, output_regs[0], input_regs[0], input_regs[1])); + } + //create intermediate instructions + int next_reg = rand_r(&random_seed) % intermediate_regs.size(); + for (int x = 1; x < op_tokens.size() - 1; x++) + { + instruction_trace.push_back(createInstr(op_tokens[x], intermediate_regs[next_reg], intermediate_regs[output_reg], intermediate_regs[rand() % intermediate_regs.size()])); + output_reg = next_reg; + next_reg = rand_r(&random_seed) % intermediate_regs.size(); + } + + //create final instruction + auto final_token = op_tokens.back(); + if (op_tokens.size() > 1) + { + instruction_trace.push_back(createInstr(final_token, output_regs[0], intermediate_regs[output_reg], intermediate_regs[rand() % intermediate_regs.size()])); + } + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; +} + +inline pisa::PISAInstruction *RandomStreamInstructionTest::createInstr(std::string op, std::string output, std::string input0, std::string input1) +{ + pisa::PISAInstruction *instr; + if (op == "add") + { + instr = createAddInstr(output, input0, input1); + } + else if (op == "mul") + { + instr = createMulInstr(output, input0, input1); + } + else if (op == "copy") + { + instr = createCopyInstr(output, input0); + } + else if (op == "mac") + { + instr = createMacInstr(output, input0, input1); + } + else if (op == "maci") + { + instr = createMaciInstr(output, input0, input1); + } + else if (op == "muli") + { + instr = createMuliInstr(output, input0, input1); + } + else if (op == "sub") + { + instr = createSubInstr(output, input0, input1); + } + + return instr; +} + +inline pisa::PISAInstruction *RandomStreamInstructionTest::createAddInstr(std::string output, std::string input0, std::string input1) +{ + return new pisa::instruction::Add(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input0 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input1 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); +} + +inline pisa::PISAInstruction *RandomStreamInstructionTest::createMulInstr(std::string output, std::string input0, std::string input1) +{ + return new pisa::instruction::Mul(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input0 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input1 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); +} + +inline pisa::PISAInstruction *RandomStreamInstructionTest::createMacInstr(std::string output, std::string input0, std::string input1) +{ + return new pisa::instruction::Mac(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input0 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input1 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); +} + +inline pisa::PISAInstruction *RandomStreamInstructionTest::createMaciInstr(std::string output, std::string input0, std::string input1) +{ + return new pisa::instruction::Maci(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input0 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input1 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); +} + +inline pisa::PISAInstruction *RandomStreamInstructionTest::createMuliInstr(std::string output, std::string input0, std::string input1) +{ + return new pisa::instruction::Muli(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input0 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input1 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); +} + +inline pisa::PISAInstruction *RandomStreamInstructionTest::createSubInstr(std::string output, std::string input0, std::string input1) +{ + return new pisa::instruction::Sub(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input0 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input1 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); +} + +inline pisa::PISAInstruction *RandomStreamInstructionTest::createCopyInstr(std::string output, std::string input0) +{ + return new pisa::instruction::Copy(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input0 + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"])); +} diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/sub_instruction_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/sub_instruction_test.h new file mode 100644 index 00000000..547ddc2c --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_instruction_tests/sub_instruction_test.h @@ -0,0 +1,50 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_instruction_tests/pisa_instruction_test.h" +#include +#include +#include +#include + +class subInstructionTest : public PisaInstructionTest +{ +public: + static std::string operationName() { return "sub_instruction"; } + + subInstructionTest() : + PisaInstructionTest() + { + configuration["Name"] = this->operationName(); + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + std::vector output_regs = { "output0" }; + std::vector input_regs = { { "input0" }, { "input1" } }; + + auto instr = new pisa::instruction::Sub(std::stoi(configuration["Poly_mod_log2"]), + pisa::Operand(output_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[0] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + pisa::Operand(input_regs[1] + "_" + configuration["RNS_INDEX"] + "_" + configuration["Chunk_INDEX"]), + std::stoi(configuration["RNS_INDEX"])); + + instruction_trace.push_back(instr); + + // } + created = true; + // std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests.h new file mode 100644 index 00000000..01c6b054 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests.h @@ -0,0 +1,44 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "pisa_kernel_tests/add.h" +#include "pisa_kernel_tests/add_corrected.h" +#include "program_mapper/poly_program/poly_operation_library.h" +//#include "pisa_kernel_tests/add_plain.h" +//#include "pisa_kernel_tests/base_operation.h" +//#include "pisa_kernel_tests/chained_adds.h" +//#include "pisa_kernel_tests/intt.h" +//#include "pisa_kernel_tests/mod_switch.h" +//#include "pisa_kernel_tests/mul.h" +//#include "pisa_kernel_tests/mul_plain.h" +//#include "pisa_kernel_tests/multiply_constant_inplace.h" +//#include "pisa_kernel_tests/ntt.h" +//#include "pisa_kernel_tests/relin.h" +//#include "pisa_kernel_tests/rescale.h" +//#include "pisa_kernel_tests/rotate.h" +//#include "pisa_kernel_tests/square.h" +//#include "pisa_kernel_tests/wide_add.h" +#include + +//#TODO These items are commented out for now as they will need to be updated to new programtrace format but +//the trace is still undergoing some additional improvements/reworks so to avoid lots of unnecessary code updates +//just a few test classes are enabled until full refactor done then all tests will be reworked to upgraded style. +static std::map pisa_kernel_tests = { + // { ChainedAdds::operationName(), new ChainedAdds() }, + { AddOperation::operationName(), new AddOperation() }, + { AddCorrected::operationName(), new AddCorrected() }, + // { AddPlain::operationName(), new AddPlain() }, + // { Intt::operationName(), new Intt() }, + // { ModSwitchOperation::operationName(), new ModSwitchOperation() }, + // { MulOperation::operationName(), new MulOperation() }, + // { MulPlainOperation::operationName(), new MulPlainOperation() }, + // { MultiplyConstantInplaceOperation::operationName(), new MultiplyConstantInplaceOperation() }, + // { NttOperation::operationName(), new NttOperation() }, + // { RelinOperation::operationName(), new RelinOperation() }, + // { RescaleOperation::operationName(), new RescaleOperation() }, + // { RotateOperation::operationName(), new RotateOperation() }, + // { SquareOperation::operationName(), new SquareOperation() }, + // { WideAdd::operationName(), new WideAdd() } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/add.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/add.h new file mode 100644 index 00000000..422f3e99 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/add.h @@ -0,0 +1,52 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" +#include + +class AddOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "add_operation"; } + + AddOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + //FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, RNS_TERM, CIPHER_DEGREE, OP_NAME, OUTPUT_ARGUMENT, INPUT_ARGUMENT, INPUT_ARGUMENT + + auto add_operation = pisa::poly::library::createPolyOperation("add", + { configuration["Scheme"], + configuration["Poly_mod_log2"], + configuration["Key_RNS"], + configuration["RNS"], + configuration["CipherDegree"], + "add", + "output0", + "input0", + "input1" }, + program_trace); + + program_trace->addOperation(add_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace->operations().size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/add_corrected.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/add_corrected.h new file mode 100644 index 00000000..66a70016 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/add_corrected.h @@ -0,0 +1,45 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class AddCorrected : public PisaKernelTest +{ +public: + static std::string operationName() { return "AddCorrected_operation"; } + + AddCorrected() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto add_corrected_operation = pisa::poly::library::createPolyOperation("add_corrected"); + program_trace->addOperation(add_corrected_operation); + program_trace->setScheme(fromString(configuration["Scheme"])); + + add_corrected_operation->setOperationName(add_corrected_operation->Name()); + + add_corrected_operation->addInput("input0", std::stoi(configuration["RNS"]), std::stoi(configuration["CipherDegree"])); + add_corrected_operation->addInput("input1", std::stoi(configuration["RNS"]), std::stoi(configuration["CipherDegree"])); + add_corrected_operation->addOutput("output0", std::stoi(configuration["RNS"]), std::stoi(configuration["CipherDegree"])); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace->operations().size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/add_plain.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/add_plain.h new file mode 100644 index 00000000..77ecec7b --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/add_plain.h @@ -0,0 +1,58 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class AddPlain : public PisaKernelTest +{ +public: + static std::string operationName() { return "add_plain_operation"; } + + AddPlain() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto add_plain_operation = pisa::poly::add_plain().create(); + add_plain_operation->setOperationName(add_plain_operation->Name()); + add_plain_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + add_plain_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + add_plain_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + add_plain_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + add_plain_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + add_plain_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + add_plain_operation->addInput("input0"); + add_plain_operation->addInput("input1"); + add_plain_operation->addOutput("output0"); + program_trace.push_back(add_plain_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/base_operation.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/base_operation.h new file mode 100644 index 00000000..db47630f --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/base_operation.h @@ -0,0 +1,58 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class BaseOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "base_operation"; } + + BaseOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto base_operation = pisa::poly::Add().create(); + base_operation->setOperationName(base_operation->baseName()); + base_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + base_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + base_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + base_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + base_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + base_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + base_operation->addInput("input0"); + base_operation->addInput("input1"); + base_operation->addOutput("output0"); + program_trace.push_back(base_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/chained_adds.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/chained_adds.h new file mode 100644 index 00000000..be7a2035 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/chained_adds.h @@ -0,0 +1,57 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" +#include +#include +#include + +class ChainedAdds : public PisaKernelTest +{ +public: + static std::string operationName() { return "chained_add"; } + + ChainedAdds() : + PisaKernelTest() + { + configuration["Name"] = this->operationName(); + configuration["Number_of_adds"] = "200"; + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + number_of_adds = std::stoi(configuration["Number_of_adds"]); + + for (int x = 0; x < number_of_adds; x++) + { + auto add_operation = pisa::poly::Add().create(); + add_operation->setOperationName("add"); + add_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + add_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + add_operation->setRnsTerms(std::stoi(configuration["RNS"])); + add_operation->setScheme(SCHEME::BGV); + + add_operation->addInput("a"); + add_operation->addInput("b"); + add_operation->addOutput("a"); + + program_trace.push_back(add_operation); + } + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } + + int number_of_adds = 2000; +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/intt.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/intt.h new file mode 100644 index 00000000..d6a3207d --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/intt.h @@ -0,0 +1,57 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class Intt : public PisaKernelTest +{ +public: + static std::string operationName() { return "intt_operation"; } + + Intt() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto intt_operation = pisa::poly::intt().create(); + intt_operation->setOperationName(intt_operation->Name()); + intt_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + intt_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + intt_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + intt_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + intt_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + intt_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + intt_operation->addInput("input0"); + intt_operation->addOutput("output0"); + program_trace.push_back(intt_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/mod_switch.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/mod_switch.h new file mode 100644 index 00000000..b35e9b85 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/mod_switch.h @@ -0,0 +1,57 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class ModSwitchOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "mod_switch_operation"; } + + ModSwitchOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto mod_switch_operation = pisa::poly::mod_switch().create(); + mod_switch_operation->setOperationName(mod_switch_operation->Name()); + mod_switch_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + mod_switch_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + mod_switch_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + mod_switch_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + mod_switch_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + mod_switch_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + mod_switch_operation->addInput("input0"); + mod_switch_operation->addOutput("output0"); + program_trace.push_back(mod_switch_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/mul.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/mul.h new file mode 100644 index 00000000..5ab5cb37 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/mul.h @@ -0,0 +1,58 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class MulOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "mul_operation"; } + + MulOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto mul_operation = pisa::poly::Mul().create(); + mul_operation->setOperationName(mul_operation->Name()); + mul_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + mul_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + mul_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + mul_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + mul_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + mul_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + mul_operation->addInput("input0"); + mul_operation->addInput("input1"); + mul_operation->addOutput("output0"); + program_trace.push_back(mul_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/mul_plain.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/mul_plain.h new file mode 100644 index 00000000..6dd50057 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/mul_plain.h @@ -0,0 +1,58 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class MulPlainOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "mul_plain_operation"; } + + MulPlainOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto mul_plain_operation = pisa::poly::mul_plain().create(); + mul_plain_operation->setOperationName(mul_plain_operation->Name()); + mul_plain_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + mul_plain_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + mul_plain_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + mul_plain_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + mul_plain_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + mul_plain_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + mul_plain_operation->addInput("input0"); + mul_plain_operation->addInput("input1"); + mul_plain_operation->addOutput("output0"); + program_trace.push_back(mul_plain_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/multiply_constant_inplace.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/multiply_constant_inplace.h new file mode 100644 index 00000000..e40e98ca --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/multiply_constant_inplace.h @@ -0,0 +1,57 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class MultiplyConstantInplaceOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "multiply_constant_inplace_operation"; } + + MultiplyConstantInplaceOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto multiply_constant_inplace_operation = pisa::poly::multiply_constant_inplace().create(); + multiply_constant_inplace_operation->setOperationName(multiply_constant_inplace_operation->Name()); + multiply_constant_inplace_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + multiply_constant_inplace_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + multiply_constant_inplace_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + multiply_constant_inplace_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + multiply_constant_inplace_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + multiply_constant_inplace_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + multiply_constant_inplace_operation->addInput("inputoutput0"); + multiply_constant_inplace_operation->addOutput("inputoutput0"); + program_trace.push_back(multiply_constant_inplace_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/ntt.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/ntt.h new file mode 100644 index 00000000..e0f097f1 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/ntt.h @@ -0,0 +1,57 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class NttOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "ntt"; } + + NttOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto ntt_operation = pisa::poly::ntt().create(); + ntt_operation->setOperationName(ntt_operation->Name()); + ntt_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + ntt_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + ntt_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + ntt_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + ntt_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + ntt_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + ntt_operation->addInput("input0"); + ntt_operation->addOutput("output0"); + program_trace.push_back(ntt_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h new file mode 100644 index 00000000..514842f2 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h @@ -0,0 +1,53 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/poly_program/poly_operation_library.h" +#include +#include + +class PisaKernelTest +{ + +public: + PisaKernelTest() + { + configuration["Name"] = "Default"; + configuration["RNS"] = "8"; + configuration["Key_RNS"] = "9"; + configuration["Poly_mod_log2"] = "14"; + configuration["Scheme"] = "BGV"; + program_trace = pisa::poly::PolyProgram::create(); + } + virtual void constructTest() + { + } + + const std::shared_ptr getProgramTrace() + { + try + { + if (created == true) + { + return program_trace; + } + else + throw std::runtime_error("Tried to use without creating first"); + } + catch (...) + { + throw; + } + } + + std::map &getConfiguration() + { + return configuration; + } + +protected: + std::shared_ptr program_trace; + bool created = false; + std::map configuration; +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/relin.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/relin.h new file mode 100644 index 00000000..3d44e472 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/relin.h @@ -0,0 +1,57 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class RelinOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "relin_operation"; } + + RelinOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto relin_operation = pisa::poly::relin().create(); + relin_operation->setOperationName(relin_operation->Name()); + relin_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + relin_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + relin_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + relin_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + relin_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + relin_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + relin_operation->addInput("input0"); + relin_operation->addOutput("output0"); + program_trace.push_back(relin_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/rescale.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/rescale.h new file mode 100644 index 00000000..2da4dda0 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/rescale.h @@ -0,0 +1,57 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class RescaleOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "rescale_operation"; } + + RescaleOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto rescale_operation = pisa::poly::rescale().create(); + rescale_operation->setOperationName(rescale_operation->Name()); + rescale_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + rescale_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + rescale_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + rescale_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + rescale_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + rescale_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + rescale_operation->addInput("input0"); + rescale_operation->addOutput("output0"); + program_trace.push_back(rescale_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/rotate.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/rotate.h new file mode 100644 index 00000000..dab35fcf --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/rotate.h @@ -0,0 +1,57 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class RotateOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "rotate_operation"; } + + RotateOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto rotate_operation = pisa::poly::rotate().create(); + rotate_operation->setOperationName(rotate_operation->Name()); + rotate_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + rotate_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + rotate_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + rotate_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + rotate_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + rotate_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + rotate_operation->addInput("input0"); + rotate_operation->addOutput("output0"); + program_trace.push_back(rotate_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/square.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/square.h new file mode 100644 index 00000000..808f6fe0 --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/square.h @@ -0,0 +1,57 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class SquareOperation : public PisaKernelTest +{ +public: + static std::string operationName() { return "square_operation"; } + + SquareOperation() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + auto square_operation = pisa::poly::square().create(); + square_operation->setOperationName(square_operation->Name()); + square_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + square_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + square_operation->setRnsTerms(std::stoi(configuration["RNS"])); + if (configuration["Scheme"] == "BGV") + square_operation->setScheme(SCHEME::BGV); + else if (configuration["Scheme"] == "CKKS") + { + square_operation->setScheme(SCHEME::CKKS); + } + else if (configuration["Scheme"] == "BGV") + { + square_operation->setScheme(SCHEME::BFV); + } + else + throw std::runtime_error("Invalid Scheme"); + + square_operation->addInput("input0"); + square_operation->addOutput("output0"); + program_trace.push_back(square_operation); + + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } +}; diff --git a/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/wide_add.h b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/wide_add.h new file mode 100644 index 00000000..7edb63ea --- /dev/null +++ b/p-isa_tools/program_mapper/p_isa/tests/pisa_kernel_tests/wide_add.h @@ -0,0 +1,53 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "program_mapper/p_isa/pisa_test_generator.h" +#include "program_mapper/p_isa/tests/pisa_kernel_tests/pisa_kernel_test.h" + +class WideAdd : public PisaKernelTest +{ +public: + static std::string operationName() { return "wide_add"; } + + WideAdd() : + PisaKernelTest() + { + configuration["Name"] = operationName(); + configuration["Number_of_adds"] = "10"; + configuration["CipherDegree"] = "2"; + }; + + // PisaKernelTest interface +public: + void constructTest() override + { + std::cout << "Entered Construct Test" << std::endl; + std::cout << "Configuration:" << std::endl; + for (auto config : configuration) + { + std::cout << config.first << " : " << config.second << std::endl; + } + + number_of_adds = std::stoi(configuration["Number_of_adds"]); + for (int x = 0; x < number_of_adds; x++) + { + auto add_operation = pisa::poly::Add().create(); + add_operation->setOperationName("add"); + add_operation->setCipherDegree(std::stoi(configuration["CipherDegree"])); + add_operation->setPolyModulusDegree(pow(2, std::stoi(configuration["Poly_mod_log2"]))); + add_operation->setRnsTerms(std::stoi(configuration["RNS"])); + add_operation->setScheme(SCHEME::BGV); + add_operation->addInput("a" + std::to_string(x)); + add_operation->addInput("b" + std::to_string(x)); + add_operation->addOutput("a" + std::to_string(x)); + + program_trace.push_back(add_operation); + } + created = true; + std::cout << "Created program trace, program trace size: " << program_trace.size() << std::endl; + } + + int number_of_adds = 20; +}; diff --git a/p-isa_tools/program_mapper/poly_program/operations/core.h b/p-isa_tools/program_mapper/poly_program/operations/core.h new file mode 100644 index 00000000..e221565a --- /dev/null +++ b/p-isa_tools/program_mapper/poly_program/operations/core.h @@ -0,0 +1,153 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "poly_description.h" + +using namespace pisa::poly; + +namespace pisa::poly::library::core { +/** \brief PolyOperation addition PolyOperationDesc + * Adds two polynomials with specified specified polymodulus degree, RNS terms, and polynomial parts and writes the result to the specified output. + * | Param | description | + * | ----- | ----------- | + * | FHE_SCHEME | Not Used | + * | POLYMOD_DEG_LOG2 | Specifies the modulus degree of the input polynomials | + * | KEY_RNS | Specifies number of RNS key values. Not Used | + * | RNS_TERM | specifies the number of RNS terms of the input polynomials | + * | CIPHER_DEGREE | Specifies the number of polynomial parts | + * | OP_NAME | add | + * | OUTPUT_ARGUMENT | Destination ciphertext | + * | INPUT_ARGUMENT | Input1 ciphertext label | + * | INPUT_ARGUMENT | Input2 ciphertext label | + * */ +static const PolyOperationDesc Add("add", { OP_NAME, FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, OUTPUT_ARGUMENT, INPUT_ARGUMENT, INPUT_ARGUMENT }); + +/** \brief PolyOperation subtraction PolyOperationDesc + * Op name: add + * | Param | description | + * | ----- | ----------- | + * | FHE_SCHEME | specifies the FHE_SCHEME of the poly operation | + * */ +static const PolyOperationDesc Sub("sub", { OP_NAME, FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, OUTPUT_ARGUMENT, INPUT_ARGUMENT, INPUT_ARGUMENT }); + +/** \brief PolyOperation multiplication PolyOperationDesc + * Op name: add + * | Param | description | + * | ----- | ----------- | + * | FHE_SCHEME | specifies the FHE_SCHEME of the poly operation | + * */ +static const PolyOperationDesc Mul("mul", { OP_NAME, FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, OUTPUT_ARGUMENT, INPUT_ARGUMENT, INPUT_ARGUMENT }); + +/** \brief PolyOperation Square PolyOperationDesc + * Adds two polynomials with specified specified polymodulus degree, RNS terms, and polynomial parts and writes the result to the specified output. + * | Param | description | + * | ----- | ----------- | + * | FHE_SCHEME | Not Used | + * | POLYMOD_DEG_LOG2 | Specifies the modulus degree of the input polynomials | + * | KEY_RNS | Specifies number of RNS key values. Not Used | + * | RNS_TERM | specifies the number of RNS terms of the input polynomials | + * | CIPHER_DEGREE | Specifies the number of polynomial parts | + * | OP_NAME | add | + * | OUTPUT_ARGUMENT | Destination ciphertext | + * | INPUT_ARGUMENT | Input1 ciphertext label | + * */ +static const PolyOperationDesc Square("square", { OP_NAME, FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, OUTPUT_ARGUMENT, INPUT_ARGUMENT }); + +/** \brief PolyOperation Relin PolyOperationDesc + * Adds two polynomials with specified specified polymodulus degree, RNS terms, and polynomial parts and writes the result to the specified output. + * | Param | description | + * | ----- | ----------- | + * | FHE_SCHEME | Not Used | + * | POLYMOD_DEG_LOG2 | Specifies the modulus degree of the input polynomials | + * | KEY_RNS | Specifies number of RNS key values. Not Used | + * | RNS_TERM | specifies the number of RNS terms of the input polynomials | + * | CIPHER_DEGREE | Specifies the number of polynomial parts | + * | OP_NAME | add | + * | OUTPUT_ARGUMENT | Destination ciphertext | + * | INPUT_ARGUMENT | Input1 ciphertext label | + * | INPUT_ARGUMENT | Input2 ciphertext label | + * */ +static const PolyOperationDesc Relin("relin", { OP_NAME, FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, OUTPUT_ARGUMENT, INPUT_ARGUMENT, ALPHA, QSIZE, DNUM }); + +/** \brief PolyOperation ModSwitch PolyOperationDesc + * Adds two polynomials with specified specified polymodulus degree, RNS terms, and polynomial parts and writes the result to the specified output. + * | Param | description | + * | ----- | ----------- | + * | FHE_SCHEME | Not Used | + * | POLYMOD_DEG_LOG2 | Specifies the modulus degree of the input polynomials | + * | KEY_RNS | Specifies number of RNS key values. Not Used | + * | RNS_TERM | specifies the number of RNS terms of the input polynomials | + * | CIPHER_DEGREE | Specifies the number of polynomial parts | + * | OP_NAME | add | + * | OUTPUT_ARGUMENT | Destination ciphertext | + * | INPUT_ARGUMENT | Input1 ciphertext label | + * | INPUT_ARGUMENT | Input2 ciphertext label | + * */ +static const PolyOperationDesc ModSwitch("mod", { OP_NAME, FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, OUTPUT_ARGUMENT, INPUT_ARGUMENT }); + +/** \brief PolyOperation Ntt PolyOperationDesc + * Adds two polynomials with specified specified polymodulus degree, RNS terms, and polynomial parts and writes the result to the specified output. + * | Param | description | + * | ----- | ----------- | + * | FHE_SCHEME | Not Used | + * | POLYMOD_DEG_LOG2 | Specifies the modulus degree of the input polynomials | + * | KEY_RNS | Specifies number of RNS key values. Not Used | + * | RNS_TERM | specifies the number of RNS terms of the input polynomials | + * | CIPHER_DEGREE | Specifies the number of polynomial parts | + * | OP_NAME | add | + * | OUTPUT_ARGUMENT | Destination ciphertext | + * | INPUT_ARGUMENT | Input1 ciphertext label | + * | INPUT_ARGUMENT | Input2 ciphertext label | + * */ +static const PolyOperationDesc Ntt("ntt", { OP_NAME, FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, OUTPUT_ARGUMENT, INPUT_ARGUMENT }); + +/** \brief PolyOperation Intt PolyOperationDesc + * Adds two polynomials with specified specified polymodulus degree, RNS terms, and polynomial parts and writes the result to the specified output. + * | Param | description | + * | ----- | ----------- | + * | FHE_SCHEME | Not Used | + * | POLYMOD_DEG_LOG2 | Specifies the modulus degree of the input polynomials | + * | KEY_RNS | Specifies number of RNS key values. Not Used | + * | RNS_TERM | specifies the number of RNS terms of the input polynomials | + * | CIPHER_DEGREE | Specifies the number of polynomial parts | + * | OP_NAME | add | + * | OUTPUT_ARGUMENT | Destination ciphertext | + * | INPUT_ARGUMENT | Input1 ciphertext label | + * | INPUT_ARGUMENT | Input2 ciphertext label | + * */ +static const PolyOperationDesc Intt("intt", { OP_NAME, FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, OUTPUT_ARGUMENT, INPUT_ARGUMENT }); + +/** \brief PolyOperation Rescale PolyOperationDesc + * Adds two polynomials with specified specified polymodulus degree, RNS terms, and polynomial parts and writes the result to the specified output. + * | Param | description | + * | ----- | ----------- | + * | FHE_SCHEME | Not Used | + * | POLYMOD_DEG_LOG2 | Specifies the modulus degree of the input polynomials | + * | KEY_RNS | Specifies number of RNS key values. Not Used | + * | RNS_TERM | specifies the number of RNS terms of the input polynomials | + * | CIPHER_DEGREE | Specifies the number of polynomial parts | + * | OP_NAME | add | + * | OUTPUT_ARGUMENT | Destination ciphertext | + * | INPUT_ARGUMENT | Input1 ciphertext label | + * | INPUT_ARGUMENT | Input2 ciphertext label | + * */ +static const PolyOperationDesc Rescale("rescale", { OP_NAME, FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, OUTPUT_ARGUMENT, INPUT_ARGUMENT, QSIZE }); + +/** \brief PolyOperation Rotate PolyOperationDesc + * Adds two polynomials with specified specified polymodulus degree, RNS terms, and polynomial parts and writes the result to the specified output. + * | Param | description | + * | ----- | ----------- | + * | FHE_SCHEME | Not Used | + * | POLYMOD_DEG_LOG2 | Specifies the modulus degree of the input polynomials | + * | KEY_RNS | Specifies number of RNS key values. Not Used | + * | RNS_TERM | specifies the number of RNS terms of the input polynomials | + * | CIPHER_DEGREE | Specifies the number of polynomial parts | + * | OP_NAME | add | + * | OUTPUT_ARGUMENT | Destination ciphertext | + * | INPUT_ARGUMENT | Input1 ciphertext label | + * | INPUT_ARGUMENT | Input2 ciphertext label | + * */ +static const PolyOperationDesc Rotate("rotate", { OP_NAME, FHE_SCHEME, POLYMOD_DEG_LOG2, KEY_RNS, OUTPUT_ARGUMENT, INPUT_ARGUMENT, GALOIS_ELT, ALPHA, QSIZE, DNUM }); +} // namespace pisa::poly::library::core diff --git a/p-isa_tools/program_mapper/poly_program/operations/poly_description.h b/p-isa_tools/program_mapper/poly_program/operations/poly_description.h new file mode 100644 index 00000000..bca412b8 --- /dev/null +++ b/p-isa_tools/program_mapper/poly_program/operations/poly_description.h @@ -0,0 +1,51 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include + +namespace pisa::poly { + +enum PARAM_TYPE +{ + OP_NAME, + INPUT_ARGUMENT, + OUTPUT_ARGUMENT, + INPUT_OUTPUT_ARGUMENT, + POLYMOD_DEG_LOG2, + CIPHER_DEGREE, + RNS_TERM, + FHE_SCHEME, + PARAM, + GALOIS_ELT, + FACTOR, + KEY_RNS, + ALPHA, + QSIZE, + DNUM, +}; + +/** + * @brief The InstructionDesc struct stores a vector of param type objects used to describe the type of parameter in each location of an instruction + */ +struct OperationDesc +{ + OperationDesc() = default; + OperationDesc(const std::initializer_list &_params) : + params(_params) + { + } + std::vector params; +}; +struct PolyOperationDesc +{ + PolyOperationDesc() {} + PolyOperationDesc(std::string _name, OperationDesc _desc) : + name(_name), desc(_desc) {} + std::string name; + OperationDesc desc; + bool force_desc_op_name = true; +}; +} // namespace pisa::poly diff --git a/p-isa_tools/program_mapper/poly_program/poly_operation_library.h b/p-isa_tools/program_mapper/poly_program/poly_operation_library.h new file mode 100644 index 00000000..e53636f8 --- /dev/null +++ b/p-isa_tools/program_mapper/poly_program/poly_operation_library.h @@ -0,0 +1,84 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "polyprogram.h" + +namespace pisa::poly::library { +/** Map used to specify the collection of kernels */ +static std::map core_operation_library = { + { "add", library::core::Add }, + { "add_plain", library::core::Add }, + { "sub", library::core::Sub }, + { "mul", library::core::Mul }, + { "mul_plain", library::core::Mul }, + { "square", library::core::Square }, + { "ntt", library::core::Ntt }, + { "intt", library::core::Intt }, + { "relin", library::core::Relin }, + { "mod_switch", library::core::ModSwitch }, + { "rescale", library::core::Rescale }, + { "rotate", library::core::Rotate }, +}; + +static std::map extended_operation_set; +static std::map active_polynomial_operation_library; + +static void activatePolynomialLibrary() +{ + active_polynomial_operation_library.merge(core_operation_library); + + for (auto item : extended_operation_set) + { + active_polynomial_operation_library[item.first] = item.second; + } +} + +static bool library_ready = false; + +static PolyOperationDesc getPolyOperationDesc(std::string operation) +{ + if (library_ready == false) + { + activatePolynomialLibrary(); + } + + if (active_polynomial_operation_library.count(operation) == 0) + { + throw std::runtime_error("Operation: " + operation + " requested during parseInstruction but no instruction description found"); + } + auto new_instance = active_polynomial_operation_library[operation]; + return new_instance; +} + +static std::shared_ptr createPolyOperation(std::string operation) +{ + auto new_instance = getPolyOperationDesc(operation); + auto new_instruction = pisa::poly::PolyOperation::create(new_instance); + return new_instruction; +} + +static std::shared_ptr createPolyOperation(std::string operation, std::initializer_list args, std::shared_ptr parent) +{ + auto new_instance = getPolyOperationDesc(operation); + auto new_instruction = pisa::poly::PolyOperation::create(new_instance, args, parent); + return new_instruction; +} +static std::shared_ptr createPolyOperation(std::string operation, std::initializer_list args) +{ + return createPolyOperation(operation, args, std::make_shared(pisa::poly::default_global_poly_program)); +} + +static std::shared_ptr createPolyOperation(std::string operation, std::vector args, std::shared_ptr parent) +{ + auto new_instance = getPolyOperationDesc(operation); + auto new_instruction = pisa::poly::PolyOperation::create(new_instance, args, parent); + return new_instruction; +} +static std::shared_ptr createPolyOperation(std::string operation, std::vector args) +{ + return createPolyOperation(operation, args, std::make_shared(pisa::poly::default_global_poly_program)); +} + +} // namespace pisa::poly::library diff --git a/p-isa_tools/program_mapper/poly_program/polynomial.h b/p-isa_tools/program_mapper/poly_program/polynomial.h new file mode 100644 index 00000000..7773a16d --- /dev/null +++ b/p-isa_tools/program_mapper/poly_program/polynomial.h @@ -0,0 +1,58 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include + +namespace pisa::poly { + +enum class OPERAND_TYPE +{ + POLYNOMIAL, + IMMEDIATE +}; + +class Polynomial +{ +public: + Polynomial() {} + Polynomial(std::string _name) + { + register_name = _name; + } + Polynomial(std::string _name, int _num_of_rns_terms, int _num_of_polynomials) + { + register_name = _name; + num_of_polynomials = _num_of_polynomials; + num_of_rns_terms = _num_of_rns_terms; + } + static std::tuple decomposePolyStringForm(std::string polyString) + { + int first_delim_pos = polyString.find('-', 0); + int second_delim_pos = polyString.find('-', first_delim_pos + 1); + + std::string label = polyString.substr(0, first_delim_pos); + int poly_parts = std::stoi(polyString.substr(first_delim_pos + 1, second_delim_pos - first_delim_pos)); + int rns_num = std::stoi(polyString.substr(second_delim_pos + 1, polyString.npos - second_delim_pos)); + + return { label, poly_parts, rns_num }; + } + + std::string location() { return register_name; } + bool immediate() { return operand_type == OPERAND_TYPE::IMMEDIATE; } + std::string register_name; + bool in_ntt_form = false; + bool in_montgomery_form = true; + int num_of_polynomials = 2; + int num_of_rns_terms = 1; + int num_of_coefficients = 8192; + OPERAND_TYPE operand_type = OPERAND_TYPE::POLYNOMIAL; + +private: + std::vector>> data_poly_rns_coefficient; +}; + +} // namespace pisa::poly diff --git a/p-isa_tools/program_mapper/poly_program/polyprogram.cpp b/p-isa_tools/program_mapper/poly_program/polyprogram.cpp new file mode 100644 index 00000000..53ceec57 --- /dev/null +++ b/p-isa_tools/program_mapper/poly_program/polyprogram.cpp @@ -0,0 +1,413 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include "polyprogram.h" +#include + +std::string toString(SCHEME scheme, bool lowercase) +{ + switch (scheme) + { + case SCHEME::BGV: + if (!lowercase) + return "BGV"; + else + return "bgv"; + case SCHEME::CKKS: + if (!lowercase) + return "CKKS"; + else + return "ckks"; + case SCHEME::BFV: + if (!lowercase) + return "BFV"; + else + return "bfv"; + default: + throw std::runtime_error("Error converting enum scheme to string"); + } +} + +namespace pisa::poly { +const std::string &PolyOperation::Name() const +{ + return params.at("operation_name").first; +} + +void PolyOperation::setOperationName(const std::string &newOperation_name) +{ + params["operation_name"] = std::pair(newOperation_name, ValueType::STRING); +} + +int PolyOperation::getCipherDegree() const +{ + return input_locations.front().num_of_polynomials; +} + +void PolyOperation::setCipherDegree(int newCipher_degree) +{ + params["cipher_degree"] = std::pair(std::to_string(newCipher_degree), ValueType::UINT32); +} + +int PolyOperation::getRnsTerms() const +{ + return this->input_locations.front().num_of_rns_terms; +} + +void PolyOperation::setRnsTerms(int newRns_terms) +{ + params["rns_terms"] = std::pair(std::to_string(newRns_terms), ValueType::UINT32); +} + +int PolyProgram::getPolyModulusDegree() const +{ + return m_N; +} + +void PolyProgram::setPolyModulusDegree(int newPoly_modulus_degree) +{ + m_N = newPoly_modulus_degree; +} + +int PolyOperation::getGaloisElt() const +{ + try + { + return std::stoi(params.at("galois_elt").first); + } + catch (...) + { + return 0; + } +} + +void PolyOperation::setGaloisElt(int newGalois_elt) +{ + params["galois_elt"] = std::pair(std::to_string(newGalois_elt), ValueType::UINT32); +} + +int PolyProgram::getAlpha() const +{ + return alpha; +} +void PolyProgram::setAlpha(int newAlpha) +{ + alpha = newAlpha; +} + +const std::vector> &PolyProgram::operations() const +{ + return m_operations; +} + +const std::vector PolyProgram::operationsRaw() const +{ + std::vector operations; + for (auto i : m_operations) + { + operations.push_back(i.get()); + } + return operations; +} + +void PolyProgram::addOperation(std::shared_ptr operation) +{ + operation->setParentProgram(std::make_shared(*this)); + m_operations.push_back(operation); +} + +void PolyProgram::setOperations(const std::vector> &newOperations) +{ + m_operations = newOperations; +} +int PolyProgram::getDNum() const +{ + return dnum; +} +void PolyProgram::setDNum(int newDNum) +{ + dnum = newDNum; +} +int PolyProgram::getQSize() const +{ + return q_size; +} +void PolyProgram::setQSize(int newQ_size) +{ + q_size = newQ_size; +} + +int PolyProgram::getKeyRns() const +{ + return key_rns; +} + +void PolyProgram::setKeyRns(int newKey_rns) +{ + key_rns = newKey_rns; +} + +uint32_t PolyOperation::getFactor() const +{ + try + { + return std::stoul(params.at("factor").first); + } + catch (...) + { + return 0; + } +} + +void PolyOperation::setFactor(uint32_t newFactor) +{ + params["factor"] = std::pair(std::to_string(newFactor), ValueType::UINT32); +} + +SCHEME PolyProgram::scheme() const +{ + return m_scheme; +} + +void PolyProgram::setScheme(SCHEME newScheme) +{ + m_scheme = newScheme; +} + +void PolyOperation::addInput(std::string arg, int poly_num, int rns_num) +{ + + input_locations.push_back(Polynomial(arg, rns_num, poly_num)); +} + +void PolyOperation::addOutput(std::string arg, int poly_num, int rns_num) +{ + output_locations.push_back(Polynomial(arg, rns_num, poly_num)); +} + +const OperationDesc &PolyOperation::description() const +{ + return m_description; +} + +void PolyOperation::setDescription(const OperationDesc &newDescription) +{ + m_description = newDescription; +} +#ifdef ENABLE_DATA_FORMATS +void PolyOperation::setComponents(const heracles::fhe_trace::Instruction &instr_pb) +{ + for (const auto &src : instr_pb.args().dests()) + addOutput(src.symbol_name(), src.num_rns(), src.order()); + for (const auto &src : instr_pb.args().srcs()) + addInput(src.symbol_name(), src.num_rns(), src.order()); + setRnsTerms(instr_pb.args().srcs(0).num_rns()); + setCipherDegree(instr_pb.args().srcs(0).order()); + for (const auto &[k, v] : instr_pb.args().params()) + { + if (k == "galois_elt") + setGaloisElt(stoi(v.value())); + if (k == "factor") + setFactor(stoi(v.value())); + } +} +heracles::fhe_trace::Instruction *PolyOperation::getProtobuffFHETraceInstruction() +{ + heracles::fhe_trace::Instruction *instruction = new heracles::fhe_trace::Instruction(); + + instruction->set_op(this->Name()); + for (auto &output : this->output_locations) + { + auto dest = instruction->mutable_args()->add_dests(); + dest->set_symbol_name(output.register_name); + dest->set_order(output.num_of_polynomials); + dest->set_num_rns(output.num_of_rns_terms); + } + for (auto &input : this->input_locations) + { + auto src = instruction->mutable_args()->add_srcs(); + src->set_symbol_name(input.register_name); + src->set_order(input.num_of_polynomials); + src->set_num_rns(input.num_of_rns_terms); + } + + // Parse through description and process any special arguments + auto desc = this->description(); + //TODO: additional arguments not currently supported. + for (auto param : desc.params) + { + if (param == GALOIS_ELT) + { + *((*instruction->mutable_args()->mutable_params())["galois_elt"].mutable_value()) = std::to_string(this->galois_elt); + } + if (param == FACTOR) + { + *((*instruction->mutable_args()->mutable_params())["factor"].mutable_value()) = std::to_string(this->factor); + } + } + + return instruction; +} +#endif + +const std::pair PolyOperation::getParam(std::string key) const +{ + return params.at(key); +} + +const std::pair PolyOperation::getParam(int component_index) const +{ + try + { + auto key = param_index_lookup.at(component_index); + return params.at(key); + } + catch (const std::runtime_error &err) + { + std::cout << "Runtime error during" << __FUNCTION__ << ", err: " << err.what() << std::endl; + throw err; + } + catch (...) + { + std::cout << "Unknown exception caught in " << __FUNCTION__ << " in file " << __FILE__ << std::endl; + throw; + } +} + +const std::string PolyOperation::getParamKey(int component_index) const +{ + try + { + return param_index_lookup.at(component_index); + } + catch (const std::runtime_error &err) + { + std::cout << "Runtime error during" << __FUNCTION__ << ", err: " << err.what() << std::endl; + throw err; + } + catch (...) + { + std::cout << "Unknown exception caught in " << __FUNCTION__ << " in file " << __FILE__ << std::endl; + throw; + } +} + +void PolyOperation::setParam(std::pair> param) +{ + int contains = params.count(param.first); + + params[param.first] = param.second; + if (contains == 0) + { + param_index_lookup[param_index_lookup.size()] = param.first; + } +} + +std::map> &PolyOperation::getParams() +{ + return params; +} + +void PolyOperation::setParams(const std::map> &newParams) +{ + params = newParams; +} + +const std::vector &PolyOperation::getOutputLocations() const +{ + return output_locations; +} + +const std::vector &PolyOperation::getInputLocations() const +{ + return input_locations; +} + +std::shared_ptr PolyOperation::parentProgram() const +{ + return m_parent_program; +} + +void PolyOperation::setParentProgram(std::shared_ptr newParent_program) +{ + m_parent_program = newParent_program; +} + +} // namespace pisa::poly + +#ifdef ENABLE_DATA_FORMATS +heracles::common::Scheme toFHETrace(SCHEME scheme) +{ + try + { + if (scheme == SCHEME::BGV) + { + return heracles::common::SCHEME_BGV; + } + else if (scheme == SCHEME::BFV) + { + return heracles::common::SCHEME_BFV; + } + else if (scheme == SCHEME::CKKS) + { + return heracles::common::SCHEME_CKKS; + } + else + throw std::runtime_error("unknown scheme conversion request"); + } + catch (...) + { + throw std::runtime_error("Error encountered during toFHETrace"); + } +} + +SCHEME toPolyProgram(heracles::common::Scheme scheme) +{ + try + { + if (scheme == heracles::common::SCHEME_BGV) + { + return SCHEME::BGV; + } + else if (scheme == heracles::common::SCHEME_BFV) + { + return SCHEME::BFV; + } + else if (scheme == heracles::common::SCHEME_CKKS) + { + return SCHEME::CKKS; + } + else + throw std::runtime_error("unknown scheme conversion request"); + } + catch (...) + { + throw std::runtime_error("Error encountered during toHEProgram"); + } +} +#endif + +SCHEME fromString(std::string scheme) +{ + try + { + if (scheme == "bgv" || scheme == "BGV") + { + return SCHEME::BGV; + } + else if (scheme == "bfv" || scheme == "BFV") + { + return SCHEME::BFV; + } + else if (scheme == "ckks" || scheme == "CKKS") + { + return SCHEME::CKKS; + } + else + throw std::runtime_error("unknown scheme conversion request"); + } + catch (...) + { + throw std::runtime_error("Error encountered during toHEProgram"); + } +} diff --git a/p-isa_tools/program_mapper/poly_program/polyprogram.h b/p-isa_tools/program_mapper/poly_program/polyprogram.h new file mode 100644 index 00000000..196b1a4e --- /dev/null +++ b/p-isa_tools/program_mapper/poly_program/polyprogram.h @@ -0,0 +1,279 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include +#include +#include +#include + +#include "operations/core.h" +#include "polynomial.h" + +#ifdef ENABLE_DATA_FORMATS +#include +#endif + +enum class SCHEME +{ + BGV, + CKKS, + BFV +}; + +std::string toString(SCHEME scheme, bool lowercase = false); +SCHEME fromString(std::string scheme); +#ifdef ENABLE_DATA_FORMATS +heracles::common::Scheme toFHETrace(SCHEME scheme); +SCHEME toPolyProgram(heracles::common::Scheme scheme); +std::string getOpName(const heracles::fhe_trace::Instruction &instruction_pb); +#endif + +namespace pisa::poly { + +class PolyOperation; + +/** + * @brief The PolyProgram class holds a collection of Polynomial Operations and context information in which to perform those operations. + */ +class PolyProgram +{ +public: + PolyProgram() {} + static std::shared_ptr create() { return std::make_shared(); } + SCHEME scheme() const; + void setScheme(SCHEME newScheme); + int getPolyModulusDegree() const; + void setPolyModulusDegree(int newPoly_modulus_degree); + int getKeyRns() const; + void setKeyRns(int newKey_rns); + int getDNum() const; + void setDNum(int newDNum); + int getQSize() const; + void setQSize(int newQ_size); + int getAlpha() const; + void setAlpha(int newAlpha); + const std::vector> &operations() const; + const std::vector operationsRaw() const; + void addOperation(std::shared_ptr operation); + void setOperations(const std::vector> &newOperations); + +private: + int m_N = 14; + int key_rns = 4; + int alpha = 0; + int dnum = 0; + int q_size = 1; + SCHEME m_scheme = SCHEME::BGV; + + std::vector> m_operations; +}; + +static PolyProgram default_global_poly_program; + +//ValueType enum for POLY_PROGRAM +enum ValueType +{ + UINT32 = 0, + UINT64 = 1, + INT32 = 2, + INT64 = 3, + FLOAT = 4, + DOUBLE = 5, + STRING = 6 +}; + +struct Operand +{ + Operand(const std::string &location) : + location_(location) {} + std::string location() const { return location_; } + + const std::string location_; + /* Added temporarily as a compatibility fix */ + bool immediate() const { return false; } +}; + +/** + * @brief The PolyOperation class represents a polynomial operation involving 1 or more input polynomial objects. + */ +class PolyOperation +{ +public: + PolyOperation() {} + PolyOperation(const PolyOperationDesc &desc) + { + m_description = desc.desc; + operation_name = desc.name; + } + template + PolyOperation(const PolyOperationDesc &desc, Iterator start, Iterator end, std::shared_ptr parent) + { + m_description = desc.desc; + operation_name = desc.name; + setParentProgram(parent); + setOperationName(desc.name); + + std::tuple decomposed; + int index = 0; + for (auto iter = start; iter != end; ++iter) + { + PARAM_TYPE operation = m_description.params[index]; + std::string value = *iter; + switch (operation) + { + case OP_NAME: + if (!desc.force_desc_op_name) + setOperationName(value); + break; + case INPUT_ARGUMENT: + decomposed = Polynomial::decomposePolyStringForm(value); + addInput(std::get<0>(decomposed), std::get<1>(decomposed), std::get<2>(decomposed)); + break; + case OUTPUT_ARGUMENT: + decomposed = Polynomial::decomposePolyStringForm(value); + addOutput(std::get<0>(decomposed), std::get<1>(decomposed), std::get<2>(decomposed)); + break; + case INPUT_OUTPUT_ARGUMENT: + decomposed = Polynomial::decomposePolyStringForm(value); + addInput(std::get<0>(decomposed), std::get<1>(decomposed), std::get<2>(decomposed)); + addOutput(std::get<0>(decomposed), std::get<1>(decomposed), std::get<2>(decomposed)); + break; + case POLYMOD_DEG_LOG2: + m_parent_program->setPolyModulusDegree(stoi(value)); + break; + case CIPHER_DEGREE: + setCipherDegree(stoi(value)); + break; + case RNS_TERM: + setRnsTerms(stoi(value)); + break; + case FHE_SCHEME: + m_parent_program->setScheme(fromString(value)); + break; + case PARAM: + throw std::runtime_error("Params not current supported by poly_program initializer list"); + break; + case GALOIS_ELT: + setGaloisElt(stoi(value)); + break; + case FACTOR: + setFactor(stoi(value)); + break; + case KEY_RNS: + m_parent_program->setKeyRns(stoi(value)); + break; + case ALPHA: + m_parent_program->setAlpha(stoi(value)); + break; + case QSIZE: + m_parent_program->setQSize(stoi(value)); + break; + case DNUM: + m_parent_program->setDNum(stoi(value)); + break; + } + index++; + } + } + + static std::string baseName() { return std::string("Default OP"); } + + virtual std::shared_ptr create() { return std::make_shared(); } + static std::shared_ptr create(PolyOperationDesc desc) { return std::make_shared(desc); } + + static std::shared_ptr create(PolyOperationDesc desc, std::initializer_list args, std::shared_ptr parent) + { + if (args.size() > desc.desc.params.size()) + { + throw std::runtime_error("Number of arguments does not match requested polynomial operation description"); + } + else if ((args.size() < desc.desc.params.size())) + { + std::cerr << "WARNING: " << desc.name << " specifies " << desc.desc.params.size() << " but only " << args.size() << " were provided"; + } + + return std::make_shared(desc, args.begin(), args.end(), parent); + } + static std::shared_ptr create(PolyOperationDesc desc, std::vector args, std::shared_ptr parent) + { + if (args.size() > desc.desc.params.size()) + { + throw std::runtime_error("Number of arguments does not match requested polynomial operation description"); + } + else if ((args.size() < desc.desc.params.size())) + { + std::cerr << "WARNING: " << desc.name << " specifies " << desc.desc.params.size() << " but only " << args.size() << " were provided"; + } + return std::make_shared(desc, args.begin(), args.end(), parent); + } + + std::string outLabel() { return ""; } + + void addInput(std::string arg, int poly_num, int rns_num); + void addOutput(std::string arg, int poly_num, int rns_num); + int numInputOperands() const { return input_locations.size(); } + int numOutputOperands() const { return output_locations.size(); } + Polynomial getInputOperand(int x) const { return input_locations[x]; } + Polynomial getOutputOperand(int x) const { return output_locations[x]; } + + const OperationDesc &description() const; + void setDescription(const OperationDesc &newDescription); + const std::string &Name() const; + void setOperationName(const std::string &newOperation_name); + int getCipherDegree() const; + void setCipherDegree(int newCipher_degree); + int getRnsTerms() const; + void setRnsTerms(int newRns_terms); + int getGaloisElt() const; + void setGaloisElt(int newGalois_elt); + uint32_t getFactor() const; + void setFactor(uint32_t newFactor); + +#ifdef ENABLE_DATA_FORMATS + void setComponents(const heracles::fhe_trace::Instruction &instr_pb); + heracles::fhe_trace::Instruction *getProtobuffFHETraceInstruction(); +#endif + + // Param System + const std::pair getParam(std::string key) const; + const std::pair getParam(int component_index) const; + const std::string getParamKey(int component_index) const; + void setParam(std::pair> param); + + // PolyProgram level operations if present + + std::shared_ptr parentProgram() const; + void setParentProgram(std::shared_ptr newParent_program); + + const std::vector &getInputLocations() const; + const std::vector &getOutputLocations() const; + +protected: + std::map> &getParams(); + void setParams(const std::map> &newParams); + + int rns_terms; + int cipher_degree; + int galois_elt; + uint32_t factor; + + std::string operation_name; + std::vector input_locations; + std::vector output_locations; + std::map> params; + std::map param_index_lookup; + OperationDesc m_description; + std::shared_ptr m_parent_program = std::make_shared(default_global_poly_program); +}; + +} // namespace pisa::poly + +class ProgramTrace +{ +public: + ProgramTrace() = default; +}; diff --git a/p-isa_tools/program_mapper/program_mapper.cpp b/p-isa_tools/program_mapper/program_mapper.cpp new file mode 100644 index 00000000..313b9610 --- /dev/null +++ b/p-isa_tools/program_mapper/program_mapper.cpp @@ -0,0 +1,252 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include "program_mapper.h" + +namespace pisa { + +template +void ProgramMapper::generatePisaProgramFromHEProgram(std::shared_ptr program_trace) +{ + // Generates a map from raw seal trace to V0 trace to allow for + // input/output variable name alignment with trace file to support + // validation. + std::map register_map; + + auto program_graph = graph::Graph::createGraph(program_trace->operationsRaw()); + auto program_inputs = program_graph.getInputNodes(true, true, false); + auto program_outputs = program_graph.getOutputNodes(); + for (auto input : program_inputs) + { + register_map[input.GetDat().label] = input.GetDat().label; + } + for (auto output : program_outputs) + { + register_map[output.GetDat().label] = output.GetDat().label; + } + + if (arguments.enable_intermediates == true) + { + auto all_program_data_nodes = program_graph.getDataGraph(); + for (auto data : all_program_data_nodes.getNodes()) + { + register_map[data.GetDat().label] = data.GetDat().label; + } + } + + // Generates map mapping all input and output locations in the program + // trace not part of the program_trace input/output set to include + // operation namespace. + // This is used for efficiently linking HE_Operations together. + // Currently input/output operations are determined via the trace @TODO + // better to switch this to graph based approach + // and simply raise warnings when there is a mismatch between graph and + // trace. A middle term solution maybe to create option to use a handinput_parser_v0 + // aligned map. + register_map = mapProgramTraceOperationsIntoRegisterMap(program_trace->operations(), register_map); + + // Generate PISAKernels(ninja kernels) as needed by the program trace. + // Checks for each kernel the HE operation and parameters and calls + // kerngen.py to generate the ninja kernel if it does not already + // exist in the kernel_cache. + + int max_rns_term = 0; + auto p_isa_kernels = generatePISAKernelsFromHEOperationVector(program_trace->operations(), arguments.kerngen, max_rns_term); + + /* Remaps kernel input/output names based upon the trace mapping There + * exists a 1 : 1 mapping between HEOperation -> PISAKernel, this + * function aligns the PISA input/output names to match HEOperation + */ + mapKernelInputOutputToRegisterMap(p_isa_kernels, program_trace->operations(), register_map); + + // Generates a vector containing all of the remapped instructions for + // all of the p_isa_kernels resulting in a single set of all P-ISA + // instructions necessary to execute + // the operations specified in the seal trace. + auto combined_instructions = outputCombinedPisaInstructions(p_isa_kernels, arguments.apply_name_spacing); + + /* Apply instruction graph rewritter and instruction hardware fixes and optimization */ + PISAGraphOptimizer graph_optimizer; + graph_optimizer.applyDuplicateInputVariableSeparation(combined_instructions); + + /* generates graph from combined p_isa instructions */ + auto p_isa_graph = graph::Graph::createGraph(combined_instructions); + combined_instructions = graph_optimizer.generateInstructionStreamFromGraph(p_isa_graph, true, combined_instructions); + + /* Re-generate combined instructions as needed */ + if (arguments.generated_name != "") + { + PisaTestGenerator test_gen; + auto generated_json = test_gen.generateJSONForGraph(p_isa_graph); + test_gen.populateCalculatedOutputResults(combined_instructions, generated_json); + test_gen.writeJSON(generated_json, arguments.generated_name); + } + + /* renders instructions from graph, Current renders two graphs, the + * seal program trace graph at the HEOperation level and the P-ISA + * instruction level graph*/ + if (arguments.generate_graphs) + { + + if (arguments.export_dot) + { + std::cout << "Writing graph to dot file: " << arguments.dot_file_name << std::endl; + std::string dot_file_name = arguments.outfile_prefix.replace_extension("dot"); +#pragma omp parallel sections + { +#pragma omp section + program_graph.writeDotFile(arguments.dot_file_name, graph::NAME); +#pragma omp section + p_isa_graph.writeDotFile(dot_file_name, graph::NAME); + } // End of parallel region + + auto inputs = p_isa_graph.getInputNodes(); + std::cout << "P_ISA Graph Input Nodes\n" + << graph::with_delimiter(inputs, "\n"); + } + } + + /* Outputs the combined p_isa_kernel to an instruction stream, + * PISAInstructions are implemented so that they print out functionally + * with stream operator + */ + const auto final_kernel_filename = arguments.outfile_prefix.replace_extension("csv"); + writeToFileBy(final_kernel_filename.string(), combined_instructions, + [membank = arguments.output_memory_bank](auto *instruction) { + instruction->setOutputBlock(membank); + return *instruction; + }); + + if (arguments.verbose) + { + for (const auto &instruction : combined_instructions) + std::cout << *instruction << std::endl; + } + + /* generates memory file for p_isa_graph */ + std::vector mem_file = generateMemFile(p_isa_graph, max_rns_term); + const auto memory_filename = arguments.outfile_prefix.replace_extension("tw.mem"); + writeToFileBy(memory_filename.string(), mem_file, [](const auto &line) { return line; }); + + return; +} + +template +std::vector ProgramMapper::generatePISAKernelsFromHEOperationVector(std::vector> program_trace, + std::string kerngen_loc, int &max_rns_term) +{ + try + { + auto kernel_cache = pisa::kernel::Cache(arguments.cache_dir, arguments.use_kernel_cache, arguments.remove_cache); + std::vector p_isa_kernels(program_trace.size()); + + std::transform(program_trace.begin(), program_trace.end(), p_isa_kernels.begin(), + [&kernel_cache, &max_rns_term, &kerngen_loc, this](auto s) { + // max_rns_term = std::max(max_rns_term, s->getRnsTerms()); + auto kernel = pisa::kernel::PISAKernel::create(kerngen_loc, s.get(), kernel_cache, arguments.verbose, arguments.new_kerngen, arguments.kernel_library); + //Graph variable extraction + kernel->determineVariableNamingViaGraph(); + return kernel; + }); + + return p_isa_kernels; + } + catch (const std::runtime_error &err) + { + std::cout << "Runtime error during kernel generation, err: " << err.what() << std::endl; + throw err; + } + catch (...) + { + std::cout << "Unknown exception caught in " << __FUNCTION__ << "in file" << __FILE__ << std::endl; + throw; + } +} + +template +void ProgramMapper::mapKernelInputOutputToRegisterMap(std::vector &p_isa_kernels, const std::vector> &program_trace, + std::map ®ister_map, std::vector>> immediates) +{ + + for (int y = 0; y < program_trace.size(); y++) + { + std::shared_ptr i = program_trace[y]; + for (int x = 0; x < i->numInputOperands(); x++) + { + std::string name = register_map[i->getInputOperand(x).location()]; + p_isa_kernels[y]->mapInput(x, name); + } + for (int x = 0; x < i->numOutputOperands(); x++) + { + std::string name = register_map[i->getOutputOperand(x).location()]; + p_isa_kernels[y]->mapOutput(x, name); + } + } +} + +// Generates map mapping all input and output locations in the program +// trace not part of the program_trace input/output set to include +// operation namespace. +// This is used for efficiently linking HE_Operations together. +// Currently input/output operations are determined via the trace @TODO +// better to switch this to graph based approach +// and simply raise warnings when there is a mismatch between graph and +// trace. A middle term solution maybe to create option to use a hand +// aligned map +template +std::map ProgramMapper::mapProgramTraceOperationsIntoRegisterMap(std::vector> program_trace, + std::map register_map) +{ + + for (int x = 0; x < program_trace.size(); x++) + { + std::string location; + for (int y = 0; y < program_trace[x]->numInputOperands(); y++) + { + location = program_trace[x]->getInputOperand(y).location(); + std::string value = register_map[location]; + if (value.size() == 0) + { + register_map[location] = program_trace[x]->Name() + std::to_string(x) + "input" + std::to_string(y); + } + } + for (int y = 0; y < program_trace[x]->numOutputOperands(); y++) + { + location = program_trace[x]->getOutputOperand(y).location(); + std::string value = register_map[location]; + if (value.size() == 0) + { + register_map[location] = program_trace[x]->Name() + std::to_string(x) + "output" + std::to_string(y); + } + } + } + return register_map; +} + +template +std::vector ProgramMapper::outputCombinedPisaInstructions(std::vector p_isa_kernels, bool apply_namespacing) +{ + std::vector combined_instructions; + for (const auto &kernel : p_isa_kernels) + { + kernel->setEnableNamespace(apply_namespacing); + const auto &instructions = kernel->getMappedInstructions(); + combined_instructions.insert(combined_instructions.end(), instructions.begin(), instructions.end()); + } + + return combined_instructions; +} + +template +const ProgramMapperArguments &ProgramMapper::getArguments() const +{ + return arguments; +} + +template +void ProgramMapper::setArguments(const ProgramMapperArguments &newArguments) +{ + arguments = newArguments; +} + +} // namespace pisa diff --git a/p-isa_tools/program_mapper/program_mapper.h b/p-isa_tools/program_mapper/program_mapper.h new file mode 100644 index 00000000..4dc099cb --- /dev/null +++ b/p-isa_tools/program_mapper/program_mapper.h @@ -0,0 +1,78 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include +#include +#include +#include + +#include "p_isa/pisa_test_generator.h" +#include "utility_functions.h" +#include +#include +#include +#include +#include +#include +#include + +namespace pisa { + +using DATA_TYPE = uint32_t; +namespace fs = std::filesystem; + +struct ProgramMapperArguments +{ + fs::path program_trace_location; + fs::path outfile_prefix; + fs::path kerngen; + fs::path dot_file_name; + fs::path cache_dir = "./kernel_cache"; + fs::path out_dir = "./"; + fs::path generated_name = ""; + bool verbose = false; + bool export_dot = false; + bool output_memory_bank = false; + bool remove_cache = false; + bool new_kerngen = true; + bool generate_graphs = true; + bool apply_name_spacing = true; + bool use_kernel_cache = true; + std::string kernel_library = "HDF"; + bool export_program_trace = false; + bool enable_intermediates = false; +}; + +template +class ProgramMapper +{ +public: + ProgramMapper() + { + } + + void generatePisaProgramFromHEProgram(std::shared_ptr program_trace); + + // Generate PISAKernels(ninja kernels) as needed by the program trace. + // Checks for each kernel the HE operation and parameters and calls + // kerngen.py to generate the ninja kernel if it does not already + // exist in the kernel_cache. + std::vector generatePISAKernelsFromHEOperationVector(std::vector> operations, std::string kerngen_loc, int &max_rns_terms); + + void mapKernelInputOutputToRegisterMap(std::vector &p_isa_kernels, const std::vector> &program_trace, + std::map ®ister_map, std::vector>> immediates = std::vector>>()); + std::map mapProgramTraceOperationsIntoRegisterMap(std::vector> program_trace, std::map register_map = std::map()); + std::vector outputCombinedPisaInstructions(std::vector p_isa_kernels, bool apply_namespacing); + + const ProgramMapperArguments &getArguments() const; + void setArguments(const ProgramMapperArguments &newArguments); + + ProgramMapperArguments arguments; +}; + +} // namespace pisa +#include "program_mapper.cpp" diff --git a/p-isa_tools/program_mapper/trace_parser/program_trace_helper.cpp b/p-isa_tools/program_mapper/trace_parser/program_trace_helper.cpp new file mode 100644 index 00000000..e0ea9822 --- /dev/null +++ b/p-isa_tools/program_mapper/trace_parser/program_trace_helper.cpp @@ -0,0 +1,591 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include +#include +#include +#include + +#ifdef ENABLE_DATA_FORMATS +#include "google/protobuf/util/json_util.h" +#endif +#include + +#include "program_mapper/poly_program/poly_operation_library.h" +#include "program_trace_helper.h" + +#ifdef ENABLE_DATA_FORMATS +#include +#endif + +#ifdef ENABLE_DATA_FORMATS +std::shared_ptr PolynomialProgramHelper::parse(const heracles::fhe_trace::Trace &trace_pb, bool verbose) +{ + try + { + if (verbose) + heracles::util::fhe_trace::print_trace(trace_pb); + auto program = pisa::poly::PolyProgram::create(); + + std::string scheme = heracles::common::Scheme_descriptor() + ->FindValueByNumber(trace_pb.scheme()) + ->options() + .GetExtension(heracles::common::string_name); + + program->setScheme(toPolyProgram(trace_pb.scheme())); + program->setPolyModulusDegree(trace_pb.n()); + program->setKeyRns(trace_pb.key_rns_num()); + program->setAlpha(trace_pb.alpha()); + program->setQSize(trace_pb.q_size()); + program->setDNum(trace_pb.dnum()); + + for (const auto &inst_pb : trace_pb.instructions()) + { + std::string op = inst_pb.op(); + if (op.substr(0, 3) == "bk_") + continue; + + auto instr = pisa::poly::library::createPolyOperation(op); + instr->setOperationName(op); + instr->setComponents(inst_pb); + program->addOperation(instr); + } + return program; + } + catch (const std::runtime_error &err) + { + std::cout << "Runtime error during ProgramTraceParser::Parse, err: " << err.what() << std::endl; + throw err; + } + catch (...) + { + std::cout << "Unknown exception caught in " << __FUNCTION__ << "in file" << __FILE__ << std::endl; + throw; + } +} +#endif + +std::shared_ptr PolynomialProgramHelper::parse(const std::string &filename_or_prefix, POLYNOMIAL_PROGRAM_FORMAT format, bool ignore_header) +{ + try + { + + if (format == POLYNOMIAL_PROGRAM_FORMAT::CSV) + { + return parseCSV(filename_or_prefix, ignore_header); + } +#ifdef ENABLE_DATA_FORMATS + else if (format == POLYNOMIAL_PROGRAM_FORMAT::PROTOBUFF) + { + return parseProtoBuff(filename_or_prefix); + } +#endif + else + { + throw std::runtime_error("UNSUPPORTED TRACE FORMAT"); + } + } + catch (const std::runtime_error &err) + { + std::cout << "Runtime error during parse, err: " << err.what() << std::endl; + throw err; + } + catch (...) + { + std::cout << "Unknown exception caught in " << __FUNCTION__ << "in file" << __FILE__ << std::endl; + throw; + } +} + +std::shared_ptr PolynomialProgramHelper::parseCSV(const std::string &filename, bool ignore_header) +{ + try + { + auto new_poly_program = pisa::poly::PolyProgram::create(); + + std::ifstream file(filename); + if (!file.is_open()) + { + throw std::runtime_error("File not found: " + filename); + } + + std::string current_line; + if (ignore_header) + std::getline(file, current_line); + + while (std::getline(file, current_line)) + { + std::vector components; + std::stringstream current_line_ss(current_line); + + std::string component; + while (std::getline(current_line_ss, component, ',')) + { + components.push_back(trim(component)); + } + new_poly_program->addOperation(parseInstruction(components, new_poly_program)); + } + return new_poly_program; + } + catch (const std::runtime_error &err) + { + std::cout << "Runtime error during parse, err: " << err.what() << std::endl; + throw err; + } + catch (...) + { + std::cout << "Unknown exception caught in " << __FUNCTION__ << "in file" << __FILE__ << std::endl; + throw; + } +} + +#ifdef ENABLE_DATA_FORMATS +std::shared_ptr PolynomialProgramHelper::parseProtoBuff(const std::string &filename, bool verbose) +{ + try + { + heracles::fhe_trace::Trace trace_pb = heracles::fhe_trace::load_trace(filename); + return parse(trace_pb); + } + catch (const std::runtime_error &err) + { + std::cout << "Runtime error during parse, err: " << err.what() << std::endl; + throw err; + } + catch (...) + { + std::cout << "Unknown exception caught in " << __FUNCTION__ << "in file" << __FILE__ << std::endl; + throw; + } +} +#endif + +void PolynomialProgramHelper::writeTraceToCSV(const std::shared_ptr trace, std::string file_name) +{ + try + { + std::vector instructions; + std::ofstream file(file_name); + if (!file.is_open()) + { + throw std::runtime_error("File not found: " + file_name); + } + + std::string header = "scheme,poly_modulus_degree,rns_terms,cipher_degree,instruction,arg0,arg1,arg2,arg3,arg4,arg5,arg6,arg7,arg8,arg9"; + file << header << "\n"; + for (auto &HE_op : trace->operations()) + { + + auto instruction_components = writeToASCIIComponents(HE_op); + + std::string he_op_string; + //Write string + for (auto component : instruction_components) + { + he_op_string += component + ","; + } + he_op_string.pop_back(); + + file << he_op_string << "\n"; + } + } + catch (const std::runtime_error &err) + { + std::cout << "Runtime error during writeTraceToCSV, err: " << err.what() << std::endl; + throw err; + } + catch (...) + { + std::cout << "Unknown exception caught in " << __FUNCTION__ << "in file" << __FILE__ << std::endl; + throw; + } +} + +#ifdef ENABLE_DATA_FORMATS +void PolynomialProgramHelper::writeTraceToProtoBuff(const std::shared_ptr trace, std::string file_name) +{ + try + { + // Create sample trace + heracles::fhe_trace::Trace protobuff_trace; + protobuff_trace.set_n(trace->getPolyModulusDegree()); + // - context + auto HE_scheme = trace->scheme(); + protobuff_trace.set_scheme(toFHETrace(HE_scheme)); + auto key_rns = trace->getKeyRns(); + protobuff_trace.set_key_rns_num(key_rns); + + for (auto &instr : trace->operations()) + { + auto proto_instr = instr->getProtobuffFHETraceInstruction(); + protobuff_trace.add_instructions()->CopyFrom(*proto_instr); + } + + std::cout << "debug string: " << protobuff_trace.DebugString() << std::endl; + std::string json; + auto rc = google::protobuf::util::MessageToJsonString(protobuff_trace, &json); + std::cout << "json: " << json << std::endl; + + // // accessing enums as default strings and as our own version .. + auto scheme = protobuff_trace.scheme(); + std::cout << "scheme: as-num=" << scheme + << " / as-default-string=" << heracles::common::Scheme_descriptor()->FindValueByNumber(scheme)->name() + << " / as-friendly-string=" + << heracles::common::Scheme_descriptor()->value(scheme)->options().GetExtension( + heracles::common::string_name) + << std::endl; + + // serialize it to file + if (!heracles::fhe_trace::store_trace(file_name, protobuff_trace)) + { + std::cerr << "Could not serialize" << std::endl; + exit(1); + } + } + catch (const std::runtime_error &err) + { + std::cout << "Runtime error during writeTraceToCSV, err: " << err.what() << std::endl; + throw err; + } + catch (...) + { + std::cout << "Unknown exception caught in " << __FUNCTION__ << "in file" << __FILE__ << std::endl; + throw; + } +} +#endif + +std::shared_ptr PolynomialProgramHelper::parseInstruction(const std::vector &components, std::shared_ptr program) +{ + try + { + std::string operation = components[OP_CODE_LOCATION]; + auto trimmed = std::remove(operation.begin(), operation.end(), ' '); + operation.erase(trimmed, operation.end()); + + auto new_instruction = pisa::poly::library::createPolyOperation(operation, components, program); + + return new_instruction; + } + catch (const std::out_of_range &err) + { + throw std::runtime_error("No Instruction Desc found for operation in InstructionMap map. Operation: " + components[OP_CODE_LOCATION]); + } + catch (...) + { + std::cout << "Invalid instruction detected during parsing."; + throw; + } +} + +std::vector PolynomialProgramHelper::writeToASCIIComponents(const std::shared_ptr operation) +{ + try + { + std::vector components; + auto op_desc = operation->description(); + + int count = 0; + for (int component_index = 0; component_index < op_desc.params.size(); component_index++) + { + std::string op_value; + extractComponent(operation, op_value, component_index); + components.push_back(op_value); + } + return components; + } + catch (...) + { + std::cout << "Error during parswriteToASCIIComponentsing."; + throw; + } +} + +#ifdef ENABLE_DATA_FORMATS +std::shared_ptr PolynomialProgramHelper::parseInstruction(const heracles::fhe_trace::Instruction &instruction_pb) +{ + // get op name + std::string op = instruction_pb.op(); + try + { + auto instr = pisa::poly::library::createPolyOperation(op); + + instr->setOperationName(op); + instr->setComponents(instruction_pb); + + return instr; + } + catch (const std::out_of_range &err) + { + throw std::runtime_error("No Instruction Desc found for operation in InstructionMap map. Operation: " + op); + } + catch (...) + { + std::cout << "Invalid instruction detected during parsing."; + throw; + } +} +#endif + +void PolynomialProgramHelper::parseComponent(const std::string &component, PARAM_TYPE type, std::shared_ptr instr) +{ + switch (type) + { + case CIPHER_DEGREE: + parse_CIPHER_DEGREE(component, instr); + break; + case OP_NAME: + parse_OP_NAME(component, instr); + break; + case INPUT_ARGUMENT: + parse_INPUT_ARGUMENT(component, instr); + break; + case OUTPUT_ARGUMENT: + parse_OUTPUT_ARGUMENT(component, instr); + break; + case INPUT_OUTPUT_ARGUMENT: + parse_INPUT_OUTPUT_ARGUMENT(component, instr); + break; + case POLYMOD_DEG_LOG2: + parse_POLYMOD_DEG_LOG2(component, instr); + break; + case RNS_TERM: + parse_RNS_TERM(component, instr); + break; + case FHE_SCHEME: + parse_FHE_SCHEME(component, instr); + break; + case GALOIS_ELT: + parse_GALOIS_ELT(component, instr); + break; + case FACTOR: + parse_FACTOR(component, instr); + break; + case KEY_RNS: + parse_KEY_RNS(component, instr); + break; + case PARAM: + case ALPHA: + case DNUM: + case QSIZE: + throw std::runtime_error("Not implemented"); + default: + throw std::runtime_error("Unhandled component during parsing"); + } +} + +void PolynomialProgramHelper::extractComponent(const std::shared_ptr instr, std::string &component, int component_index) +{ + auto type = instr->description().params[component_index]; + + switch (type) + { + case CIPHER_DEGREE: + extract_CIPHER_DEGREE(component, instr, component_index); + break; + case OP_NAME: + extract_OP_NAME(component, instr, component_index); + break; + case INPUT_ARGUMENT: + extract_INPUT_ARGUMENT(component, instr, component_index); + break; + case OUTPUT_ARGUMENT: + extract_OUTPUT_ARGUMENT(component, instr, component_index); + break; + case INPUT_OUTPUT_ARGUMENT: + extract_INPUT_OUTPUT_ARGUMENT(component, instr, component_index); + break; + case POLYMOD_DEG_LOG2: + extract_POLYMOD_DEG_LOG2(component, instr, component_index); + break; + case RNS_TERM: + extract_RNS_TERM(component, instr, component_index); + break; + case FHE_SCHEME: + extract_FHE_SCHEME(component, instr, component_index); + break; + case GALOIS_ELT: + extract_GALOIS_ELT(component, instr, component_index); + break; + case FACTOR: + extract_FACTOR(component, instr, component_index); + break; + case KEY_RNS: + extract_KEY_RNS(component, instr, component_index); + break; + case PARAM: + case ALPHA: + case DNUM: + case QSIZE: + throw std::runtime_error("Not implemented"); + default: + throw std::runtime_error("Unhandled component during parsing"); + } +} + +#ifdef ENABLE_DATA_FORMATS +heracles::fhe_trace::Instruction *PolynomialProgramHelper::getProtobuffInstruction(std::shared_ptr instr) +{ + return instr->getProtobuffFHETraceInstruction(); +} +#endif + +void PolynomialProgramHelper::parse_OP_NAME(const std::string &component, std::shared_ptr instr) +{ + instr->setOperationName(whiteSpaceRemoved(component)); +} + +void PolynomialProgramHelper::parse_CIPHER_DEGREE(const std::string &component, std::shared_ptr instr) +{ + instr->setCipherDegree(stoi(component)); +} + +void PolynomialProgramHelper::parse_INPUT_ARGUMENT(const std::string &component, std::shared_ptr instr) +{ + // instr->addInput(component); +} + +void PolynomialProgramHelper::parse_OUTPUT_ARGUMENT(const std::string &component, std::shared_ptr instr) +{ + // instr->addOutput(component); +} + +void PolynomialProgramHelper::parse_INPUT_OUTPUT_ARGUMENT(const std::string &component, std::shared_ptr instr) +{ + // instr->addInput(component); + // instr->addOutput(component); +} + +void PolynomialProgramHelper::parse_POLYMOD_DEG_LOG2(const std::string &component, std::shared_ptr instr) +{ + instr->parentProgram()->setPolyModulusDegree(std::stoi(component)); +} + +void PolynomialProgramHelper::parse_RNS_TERM(const std::string &component, std::shared_ptr instr) +{ + instr->setRnsTerms(std::stoi(component)); +} + +void PolynomialProgramHelper::parse_FHE_SCHEME(const std::string &component, std::shared_ptr instr) +{ + instr->parentProgram()->setScheme(fromString(component)); +} + +void PolynomialProgramHelper::parse_GALOIS_ELT(const std::string &component, std::shared_ptr instr) +{ + instr->setGaloisElt(std::stoi(component)); +} + +void PolynomialProgramHelper::parse_FACTOR(const std::string &component, std::shared_ptr instr) +{ + instr->setFactor(std::stoi(component)); +} + +void PolynomialProgramHelper::parse_KEY_RNS(const std::string &component, std::shared_ptr instr) +{ + instr->parentProgram()->setKeyRns(stoi(component)); +} + +void PolynomialProgramHelper::parse_PARAM(const std::pair> component, std::shared_ptr instr) +{ + instr->setParam(component); +} + +void PolynomialProgramHelper::extract_OP_NAME(std::string &component, const std::shared_ptr instr, int component_index) +{ + component = instr->Name(); +} + +void PolynomialProgramHelper::extract_CIPHER_DEGREE(std::string &component, const std::shared_ptr instr, int component_index) +{ + component = std::to_string(instr->getCipherDegree()); +} + +void PolynomialProgramHelper::extract_INPUT_ARGUMENT(std::string &component, const std::shared_ptr instr, int component_index) +{ + int desired_input_index = 0; + auto desc = instr->description(); + for (int x = 0; x < component_index; x++) + { + if (desc.params[x] == PARAM_TYPE::INPUT_ARGUMENT || desc.params[x] == PARAM_TYPE::INPUT_OUTPUT_ARGUMENT) + { + desired_input_index++; + } + } + component = instr->getInputOperand(desired_input_index).register_name; +} + +void PolynomialProgramHelper::extract_OUTPUT_ARGUMENT(std::string &component, const std::shared_ptr instr, int component_index) +{ + int desired_output_index = 0; + auto desc = instr->description(); + for (int x = 0; x < component_index; x++) + { + if (desc.params[x] == PARAM_TYPE::OUTPUT_ARGUMENT || desc.params[x] == PARAM_TYPE::INPUT_OUTPUT_ARGUMENT) + { + desired_output_index++; + } + } + component = instr->getOutputOperand(desired_output_index).register_name; +} + +void PolynomialProgramHelper::extract_INPUT_OUTPUT_ARGUMENT(std::string &component, const std::shared_ptr instr, int component_index) +{ + int desired_output_index = 0; + auto desc = instr->description(); + for (int x = 0; x < component_index; x++) + { + if (desc.params[x] == PARAM_TYPE::OUTPUT_ARGUMENT || desc.params[x] == PARAM_TYPE::INPUT_OUTPUT_ARGUMENT) + { + desired_output_index++; + } + } + component = instr->getOutputOperand(desired_output_index).register_name; +} + +void PolynomialProgramHelper::extract_POLYMOD_DEG_LOG2(std::string &component, const std::shared_ptr instr, int component_index) +{ + component = std::to_string(instr->parentProgram()->getPolyModulusDegree()); +} + +void PolynomialProgramHelper::extract_RNS_TERM(std::string &component, const std::shared_ptr instr, int component_index) +{ + component = std::to_string(instr->getRnsTerms()); +} + +void PolynomialProgramHelper::extract_FHE_SCHEME(std::string &component, const std::shared_ptr instr, int component_index) +{ + component = toString(instr->parentProgram()->scheme()); +} + +void PolynomialProgramHelper::extract_GALOIS_ELT(std::string &component, const std::shared_ptr instr, int component_index) +{ + component = std::to_string(instr->getGaloisElt()); +} + +void PolynomialProgramHelper::extract_FACTOR(std::string &component, const std::shared_ptr instr, int component_index) +{ + component = std::to_string(instr->getFactor()); +} + +void PolynomialProgramHelper::extract_KEY_RNS(std::string &component, const std::shared_ptr instr, int component_index) +{ + component = std::to_string(instr->parentProgram()->getKeyRns()); +} + +void PolynomialProgramHelper::extract_PARAM(std::pair> &component, const std::shared_ptr instr, int component_index) +{ + int desired_param_index = 0; + auto desc = instr->description(); + for (int x = 0; x < component_index; x++) + { + if (desc.params[x] == PARAM_TYPE::PARAM) + { + desired_param_index++; + } + } + + auto key = instr->getParamKey(component_index); + component.first = key; + component.second = instr->getParam(key); +} diff --git a/p-isa_tools/program_mapper/trace_parser/program_trace_helper.h b/p-isa_tools/program_mapper/trace_parser/program_trace_helper.h new file mode 100644 index 00000000..aacacaa9 --- /dev/null +++ b/p-isa_tools/program_mapper/trace_parser/program_trace_helper.h @@ -0,0 +1,79 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#ifdef ENABLE_DATA_FORMATS +#include +#endif + +enum POLYNOMIAL_PROGRAM_FORMAT +{ + PROTOBUFF, + CSV +}; + +inline std::string trim(const std::string &component) +{ + std::string res; + std::copy_if(component.cbegin(), component.cend(), std::back_inserter(res), + [](const char c) { return c != '\n' && c != '\r'; }); + return res; +} + +struct PolynomialProgramHelper +{ + PolynomialProgramHelper() = default; + + static std::shared_ptr parse(const std::string &filename_or_prefix, POLYNOMIAL_PROGRAM_FORMAT format = CSV, bool ignore_header = true); + static std::shared_ptr parseCSV(const std::string &filename, bool ignore_header = true); + static void writeTraceToCSV(const std::shared_ptr trace, std::string file_name); + +#ifdef ENABLE_DATA_FORMATS + static std::shared_ptr parse(const heracles::fhe_trace::Trace &trace_pb, bool verbose = false); + static std::shared_ptr parseProtoBuff(const std::string &filename, bool verbose = false); + static void writeTraceToProtoBuff(const std::shared_ptr trace, std::string file_name); +#endif + + static std::shared_ptr parseInstruction(const std::vector &components, std::shared_ptr program); + static std::vector writeToASCIIComponents(const std::shared_ptr operation); + +#ifdef ENABLE_DATA_FORMATS + static std::shared_ptr parseInstruction(const heracles::fhe_trace::Instruction &instruction_pb); +#endif + static void parseComponent(const std::string &component, PARAM_TYPE type, std::shared_ptr instr); + static void extractComponent(const std::shared_ptr instr, std::string &component, int component_index); +#ifdef ENABLE_DATA_FORMATS + static heracles::fhe_trace::Instruction *getProtobuffInstruction(std::shared_ptr instr); +#endif + //CSV Parsing functions + static void parse_OP_NAME(const std::string &component, std::shared_ptr instr); + static void parse_CIPHER_DEGREE(const std::string &component, std::shared_ptr instr); + static void parse_INPUT_ARGUMENT(const std::string &component, std::shared_ptr instr); + static void parse_OUTPUT_ARGUMENT(const std::string &component, std::shared_ptr instr); + static void parse_INPUT_OUTPUT_ARGUMENT(const std::string &component, std::shared_ptr instr); + static void parse_POLYMOD_DEG_LOG2(const std::string &component, std::shared_ptr instr); + static void parse_RNS_TERM(const std::string &component, std::shared_ptr instr); + static void parse_FHE_SCHEME(const std::string &component, std::shared_ptr instr); + static void parse_GALOIS_ELT(const std::string &component, std::shared_ptr instr); + static void parse_FACTOR(const std::string &component, std::shared_ptr instr); + static void parse_KEY_RNS(const std::string &component, std::shared_ptr instr); + static void parse_PARAM(const std::pair> component, std::shared_ptr instr); + + //CSV Writing functions + static void extract_OP_NAME(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_CIPHER_DEGREE(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_INPUT_ARGUMENT(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_OUTPUT_ARGUMENT(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_INPUT_OUTPUT_ARGUMENT(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_POLYMOD_DEG_LOG2(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_RNS_TERM(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_FHE_SCHEME(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_GALOIS_ELT(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_FACTOR(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_KEY_RNS(std::string &component, const std::shared_ptr instr, int component_index); + static void extract_PARAM(std::pair> &component, const std::shared_ptr instr, int component_index); + + constexpr static int OP_CODE_LOCATION = 0; +}; diff --git a/p-isa_tools/program_mapper/utility_functions.h b/p-isa_tools/program_mapper/utility_functions.h new file mode 100644 index 00000000..e24afdaa --- /dev/null +++ b/p-isa_tools/program_mapper/utility_functions.h @@ -0,0 +1,176 @@ +// Copyright (C) 2025 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include +#include +#include + +#include "p_isa/pisa_test_generator.h" +#include "program_mapper.h" +#include +#include +#include +#include +#include +#include +#include + +using DATA_TYPE = uint32_t; +namespace fs = std::filesystem; + +std::vector generateMemFile(const graph::Graph &graph, int max_rns_terms) +{ + const auto inputs = graph.getInputNodes(true, false); + const auto outputs = graph.getOutputNodes(); + + int counter = 0; + std::vector memory_file = { + "dload, ntt_auxiliary_table, " + std::to_string(counter++), + "dload, ntt_routing_table, " + std::to_string(counter++), + "dload, intt_auxiliary_table, " + std::to_string(counter++), + "dload, intt_routing_table, " + std::to_string(counter++) + }; + + // Get twid/ones iterations + int high_rns_iters = 1 + ((max_rns_terms - 1) / 64); + + //Add preamble strings + std::generate_n(std::back_inserter(memory_file), 8 * high_rns_iters, [&counter]() { + return "dload, twid, " + std::to_string(counter++); + }); + std::generate_n(std::back_inserter(memory_file), high_rns_iters, [&counter]() { + return "dload, ones, " + std::to_string(counter++); + }); + + //Add inputs + std::map hbm_address_map; + std::transform(inputs.begin(), inputs.end(), std::back_inserter(memory_file), + [&counter, &hbm_address_map](const auto &x) { + if (hbm_address_map.count(x.GetDat().label) == 0) + { + hbm_address_map[x.GetDat().label] = counter++; + } + return "dload, poly, " + std::to_string(hbm_address_map[x.GetDat().label]) + ", " + x.GetDat().label; + }); + + //Add outputs + std::transform(outputs.begin(), outputs.end(), std::back_inserter(memory_file), + [&counter, &hbm_address_map](const auto &x) { + //dstore, output_0_0_0, 73 + if (hbm_address_map.count(x.GetDat().label) == 0) + { + hbm_address_map[x.GetDat().label] = counter++; + } + return "dstore, " + x.GetDat().label + ", " + std::to_string(hbm_address_map[x.GetDat().label]); + }); + + return memory_file; +} +/** + * @brief registerNameRoot Attempts to split a register name removing the RNS and block terms. + * @todo Need to account for outlier cases when naming doesn't match + * @param location + * @return + */ +inline std::string registerNameRoot(const std::string ®_name) +{ + int size = reg_name.find('_', 0); + return reg_name.substr(0, size); +} + +template +inline std::vector nonRepeatingRoots(const T &xputs) +{ + std::vector roots; + roots.reserve(xputs.size()); + std::transform(xputs.begin(), xputs.end(), std::back_inserter(roots), + [](const auto &xput) { + const auto &root = registerNameRoot(xput.first); + return root; + }); + // removes consecutive (adjacent) duplicates + auto it = std::unique(roots.begin(), roots.end()); + roots.erase(it, roots.end()); + return roots; +} + +template +inline std::vector nonRepeatingRootsNode(const T &xputs) +{ + std::vector roots; + roots.reserve(xputs.size()); + std::transform(xputs.begin(), xputs.end(), std::back_inserter(roots), + [](const auto &xput) { + const auto &root = registerNameRoot(xput.GetDat().label); + return root; + }); + // removes consecutive (adjacent) duplicates + auto it = std::unique(roots.begin(), roots.end()); + roots.erase(it, roots.end()); + return roots; +} + +/** + * @brief generateRegisterMap Maps all input/output variable name roots in trace to a map structure. Current structure is 1 : 1 + * @param input_parser_v0 + * @return + */ +inline std::map generateRegisterMap(const JSONDataHandler &input_parser_v0) +{ + const auto &inputs_v0 = input_parser_v0.getAllInputs(); + const auto &outputs_v0 = input_parser_v0.getAllOutputs(); + const auto &intermediates_v0 = input_parser_v0.getAllIntermediatess(); + auto input_roots = nonRepeatingRoots(inputs_v0); + auto output_roots = nonRepeatingRoots(outputs_v0); + auto intermediate_roots = nonRepeatingRoots(intermediates_v0); + + for (const auto &input : inputs_v0) + { + std::cout << input.first << '\n'; + } + std::cout << std::flush; + + std::map register_map; + + for (const auto &root : input_roots) + { + register_map[root] = root; + } + + for (const auto &root : output_roots) + { + register_map[root] = root; + } + + for (const auto &root : intermediate_roots) + { + register_map[root] = root; + } + + return register_map; +} + +inline void dumpMapToFile(const std::string &file_name, const std::map &map) +{ + std::ofstream map_file(file_name); + if (!map_file.is_open()) + throw std::runtime_error("Could not open file '" + file_name + "'."); + for (const auto &[k, v] : map) + { + map_file << k << "," << v << std::endl; + } +} + +template +inline void writeToFileBy(const std::string &filename, const std::vector &inputs, FN fn) +{ + std::ofstream ofile(filename); + if (!ofile.is_open()) + throw std::runtime_error("Could not open file '" + filename + "'."); + for (const auto &input : inputs) + ofile << fn(input) << '\n'; +} diff --git a/requirements-dev.txt b/requirements-dev.txt new file mode 100644 index 00000000..79b35d79 --- /dev/null +++ b/requirements-dev.txt @@ -0,0 +1,15 @@ +# Pre-commit framework +pre-commit>=3.5.0 + +# Testing +pytest>=7.4.0 + +# Python development tools +pylint>=3.0.0 + +# Optional: For IDE integration and manual use (pre-commit handles these automatically) +black==25.1.0 +mypy==1.16.0 + +# C++ linting +cpplint>=1.6.0