From 3fd5eb7a6075d05b62771b73cdaeb21553d5d4d0 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:04:18 -0400 Subject: [PATCH 01/38] feat(rocrand): port over utils from rocprim for kernel tuner --- projects/rocrand/tuning/run_tuning.py | 199 ++++++++ projects/rocrand/tuning/tuner/base_tuner.py | 427 ++++++++++++++++++ projects/rocrand/tuning/utils/__init__.py | 23 + .../utils/compare_simulated_strategies.py | 257 +++++++++++ projects/rocrand/tuning/utils/utils.py | 92 ++++ 5 files changed, 998 insertions(+) create mode 100644 projects/rocrand/tuning/run_tuning.py create mode 100644 projects/rocrand/tuning/tuner/base_tuner.py create mode 100644 projects/rocrand/tuning/utils/__init__.py create mode 100644 projects/rocrand/tuning/utils/compare_simulated_strategies.py create mode 100644 projects/rocrand/tuning/utils/utils.py diff --git a/projects/rocrand/tuning/run_tuning.py b/projects/rocrand/tuning/run_tuning.py new file mode 100644 index 000000000000..b45a914287c7 --- /dev/null +++ b/projects/rocrand/tuning/run_tuning.py @@ -0,0 +1,199 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import importlib.util +import re +import sys +import os +from pathlib import Path +from types import ModuleType +from typing import List, Optional +from contextlib import contextmanager +import traceback +from utils import Parser, BASE_DIR +from glob import glob + +from tuner.base_tuner import TunerArgs + +@contextmanager +def working_directory(path: Path): + """Context manager for changing the current working directory.""" + prev_cwd = Path.cwd() + os.chdir(path) + try: + yield + finally: + os.chdir(prev_cwd) + + +def import_module_from_file(file_path: str, module_name: str) -> Optional[ModuleType]: + """Import a module from a file path.""" + try: + spec = importlib.util.spec_from_file_location(module_name, file_path) + if spec is None or spec.loader is None: + print(f"Failed to load spec for {file_path}") + return None + + module = importlib.util.module_from_spec(spec) + spec.loader.exec_module(module) + return module + except Exception as e: + print(f"Error importing {file_path}: {e}") + return None + + +def get_available_algorithms() -> List[str]: + """Return list of supported tuning algorithms.""" + return sorted([re.sub(r'^tuning_|\.py', '', os.path.basename(f)) for f in glob(f'{BASE_DIR}/tuner/tuning_*')]) + +def filter_algorithms(available_algos: List[str], pattern: str) -> List[str]: + """Filter algorithms based on regex pattern.""" + try: + regex = re.compile(pattern) + return [algo for algo in available_algos if regex.search(algo)] + except re.error as e: + print(f"Invalid regex pattern: {e}") + sys.exit(1) + + +def run_tuning( + algorithms: List[str], + size: int | None = None, + output_dir: str | None = None, + exclude_default_config: bool | None = None, + max_fevals: int | None = None, + strategy: str | None = None, + simulation_mode: bool | None = False, + arch_name: str | None = None, + seed: int | None = None, +) -> None: + """Run tuning for specified algorithms.""" + current_dir = Path(__file__).parent + tuning_dir = current_dir / "tuner" + + with working_directory(tuning_dir): + for algo in algorithms: + print(f"\nRunning tuning for {algo}...") + tuning_file = tuning_dir / f"tuning_{algo}.py" + + if not tuning_file.exists(): + print(f"Error: Tuning file not found for {algo}") + continue + + module_name = f"tuning_{algo}" + tuning_module = import_module_from_file(str(tuning_file), module_name) + + if tuning_module is None: + print(f"Error: Failed to import tuning module for {algo}") + continue + + print(f"Working directory: {Path.cwd()}") + try: + if hasattr(tuning_module, "Tuner"): + args: TunerArgs = tuning_module.Tuner._get_default_args() + args.update_with_kwargs( + algo_full_name=algo, + size=size, + output_dir=output_dir, + exclude_default_config=exclude_default_config, + max_fevals=max_fevals, + strategy=strategy, + simulation_mode=simulation_mode, + arch_name=arch_name, + seed=seed, + ) + tuner = tuning_module.Tuner(args) + tuner.tune() + else: + print(f"Warning: No tuner class found in {algo} module") + except Exception: + # Guarantees that earlier print statements get printed first + sys.stdout.flush() + + traceback.print_exc() + + +class RunTuningParser(Parser): + @classmethod + def get_run_tuning_parser(cls): + parser = cls.get_parser() + parser.description = "Run tuning for algorithms matching the specified regex" + parser.add_argument( + "--algo-regex", + type=str, + default=".*", + help='Regex pattern to match algorithm names (default: ".*" matches all)', + ) + parser.add_argument( + "--list", + action="store_true", + default=False, + help="List available algorithms", + ) + return parser + + +def main(): + available_algos = get_available_algorithms() + + parser = RunTuningParser.get_run_tuning_parser() + args = parser.parse_args() + + if args.list: + print("Available algorithms:") + for algo in available_algos: + print(f" - {algo}") + return + + args.size = int(args.size) if args.size else None + + matched_algos = filter_algorithms(available_algos, args.algo_regex) + + if not matched_algos: + print(f"No algorithms matched the algo_regex: {args.algo_regex}") + print("Available algorithms:") + for algo in available_algos: + print(f" - {algo}") + sys.exit(1) + + print(f"Running tuning for algorithms matching '{args.algo_regex}':") + for algo in matched_algos: + print(f" - {algo}") + + output_dir = Path(args.output_dir) + output_dir.mkdir(parents=True, exist_ok=True) + + run_tuning( + matched_algos, + args.size, + args.output_dir, + args.exclude_default_config, + args.max_fevals, + args.strategy, + args.simulation_mode, + args.arch_name, + args.seed, + ) + + +if __name__ == "__main__": + main() \ No newline at end of file diff --git a/projects/rocrand/tuning/tuner/base_tuner.py b/projects/rocrand/tuning/tuner/base_tuner.py new file mode 100644 index 000000000000..00ceaff8882a --- /dev/null +++ b/projects/rocrand/tuning/tuner/base_tuner.py @@ -0,0 +1,427 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +from abc import ABC, abstractmethod +from typing import Any, Callable, List, Optional, OrderedDict, Dict +from kernel_tuner.file_utils import store_metadata_file, store_output_file +import kernel_tuner +import json +from pathlib import Path +import numpy as np +from jinja2 import Environment, FileSystemLoader +from utils import Parser, BASE_DIR +from hip import hip # type: ignore (pyright doesn't detect hip-python correctly) +import warnings +from dataclasses import dataclass +import confgen.parse +import pathlib + +""" +The following base class is used when implementing the tuning for new algorithms +using Kernel Tuner. For a detailed example, check out the tuning of device_merge. + +The Tuner class manages the complete tuning workflow, handling parameter configurations, +device memory, result collection, and creates the necessary C++ +wrapper code that interfaces between Kernel Tuner and rocPRIM templates. Since we are +tuning templated functions in rocPRIM rather than direct kernels, we need to create C +extern wrappers that return a float representing execution time. This wrapping is +necessary because we cannot know the C++ mangled name of the compiled templated function +ahead of time. For each combination of key-value types, the Tuner creates a +unique header file. We then use Kernel Tuner's C compiler functionality to compile and +benchmark these wrappers, rather than using its HIP backend. This is because we're +tuning complete algorithm functions that may internally launch one or more kernels, +rather than tuning individual kernels directly. The wrapper files serve as the +interface between Kernel Tuner and the rocPRIM algorithm implementations, allowing us to +measure and optimize performance across the entire algorithm execution. A base template +is available to use with the different algorithms. This template can be constomized +by implementing the different jinja blocks when extending the base template. + +For each type combination, the tuning workflow in the Tuner class prepares the specific +parameters needed for tune_kernel, including tuning parameters like block sizes and +items per thread, function arguments from the wrapper, problem size specifications, and +parameter space restrictions. It also handles compiler configuration and strategy +options. A key feature integrated into the tuning process is the ability to +incorporate previous best configurations from older tuning runs. This ensures +continuity and improvement - new tuning results will be either equivalent to or better +than previous ones. This is achieved by parsing the existing configuration header file +to extract relevant configs for the current type combination being tuned. All this +happens with the ConfigParser class in config_parser.py. These configurations are added +to the parameter space, and if the chosen strategy doesn't naturally explore them, they +are benchmarked separately to ensure inclusion in the final results. + +The final output consists of JSON files containing the tuning results, organized by +architecture, algorithm name, and data type combinations. These files include time, +performance metrics, parameter configurations, and optional metadata for analysis. One +particularly useful aspect of using Kernel Tuner is its ability to pause and resume +tuning sessions without losing previous progress, thanks to the caching system that +utilizes these JSON output files. + +To adapt this framework for a new algorithm, you'll need to implement the Tuner and a +extend the jinja base template for your specific algorithm requirements. The device_merge +implementation serves as a comprehensive example. + +For detailed information: +Kernel Tuner documentation: https://kerneltuner.github.io/kernel_tuner/stable/contents.html +ROCm HIP Python Wrapper: https://rocm.docs.amd.com/projects/hip-python/en/latest/index.html +""" + +""" +Inclusive range for params tuning, edit these to adjust tuning grid range. +""" +BLOCK_SIZES = [32, 64, 128, 256, 512, 1024, 2048] +GRID_SIZES = [32, 64, 128, 256, 512, 1024] + +@dataclass +class TunerArgs: + algo_full_name: Optional[str] = None + """The full algorithm name, with type and name. + """ + + size: int = 1024 * 1024 * 32 * 4 + """Size in bytes for the problem/data to tune with. + """ + + max_fevals: int = 100 + """Maximum number of function evaluations for the tuning strategy + """ + + strategy: str = "dual_annealing" + """Tuning strategy to use (e.g. "dual_annealing", "brute_force"). + """ + + output_dir: str = "../output" + """Directory path to store tuning results and metadata. + """ + + exclude_default_config: bool = False + """Whether to also test default configurations. + """ + + save_metadata: bool = False + """Whether to save additional tuning metadata. + """ + + simulation_mode: bool = False + seed: int = -1 + arch_name: Optional[str] = None + + def update_with_kwargs(self, **kwargs): + for k, v in kwargs.items(): + if k in self.__dict__.keys() and not v is None: + self.__dict__[k] = v + + +class BaseTuner(ABC): + def __init__(self, args: TunerArgs): + """Initialize the tuner with configuration parameters.""" + self.algo_name = args.algo_full_name + + self.device_id = 0 + self.bytes_size = args.size + self.output_dir = Path(args.output_dir) + self.output_dir.mkdir(parents=True, exist_ok=True) + self.simulation_mode = args.simulation_mode + if not self.simulation_mode: + # TODO: remove or replace this to get rid of hip-python + self.device_properties = hip.hipDeviceProp_t() + hip.hipGetDeviceProperties(self.device_properties, self.device_id) + self.arch_name = self.device_properties.gcnArchName.decode().split(":")[0] + else: + assert ( + args.arch_name + ), "Simulation mode requires --arch-name to be specified" + self.arch_name = args.arch_name + self.exclude_default_config = args.exclude_default_config + self.max_fevals = args.max_fevals + self.save_metadata = args.save_metadata + self.strategy = args.strategy + self.seed = None if args.seed == -1 else args.seed + + if not self.exclude_default_config: + with open( + f"{BASE_DIR}/../library/src/rng/config/{self.algo_name}_config.hpp" + ) as f: + self.existing_config = confgen.parse.parse_lines(f.readlines()) + else: + self.existing_config = None + + @classmethod + @abstractmethod + def _get_default_args(cls) -> TunerArgs: + """Returns the default arguments. Used to generate default values for the CLI.""" + pass + + @classmethod + def cli(cls): + defaults: TunerArgs = cls._get_default_args() + args = dict( + Parser.get_parser( + default_bytes=defaults.size, + default_max_feval=defaults.max_fevals, + default_output_dir=defaults.output_dir, + default_strategy=defaults.strategy, + ) + .parse_args() + ._get_kwargs() + ) + defaults.update_with_kwargs(**args) + cls(defaults).tune() + + def _get_tune_params(self) -> OrderedDict: + params = OrderedDict() + params['block_size_x'] = BLOCK_SIZES + params['grid_size'] = GRID_SIZES + return params + + def _get_restrictions(self) -> Callable[[dict], bool] | List[str]: + """Define constraints for what parameter combinations are valid during tuning. + + Two options: + 1. Return list of string expressions that must all evaluate to True + Example: ["block_size_x * ipt <= size", "block_size_x <= 1024"] + These expressions are evaluated as boolean conditions by the Kernel Tuner. + Each string must be a valid Python expression using the parameter names. + + 2. Return a callable (function/lambda) that takes a params dictionary as input + and returns True/False to validate the configuration. + + See Kernel Tuner's documentation for more details + """ + min_total_threads = 32768 + def validate(params): + threads = params['block_size_x'] + blocks = params['grid_size'] + + return threads * blocks >= min_total_threads + + return validate + + def _get_grid_div_x(self): + """Return the grid_div_x parameter for kernel_tuner.tune_kernel()""" + return ["block_size_x"] + + def tune(self) -> None: + """ + Run tuning + """ + + print(f"\nTuning {self.algo_name}") + strategy_print_message = ( + f"Using strategy: {self.strategy if self.strategy else 'brute_force'}" + ) + strategy_print_message += ( + f", max fevals: {self.max_fevals}" if self.strategy != "brute_force" else "" + ) + print(strategy_print_message) + + try: + tune_kernel_args = self._get_base_tune_kernel_args() + # Run main tuning + results, _ = kernel_tuner.tune_kernel(**tune_kernel_args) + # Run default config if enabled + self._run_default_config(tune_kernel_args) + + cache_file_path = self._get_cache_file_path() + self._save_output(cache_file_path, results) + + except Exception as e: + print(f"Failed tuning for {self.algo_name}") + print(f"Error: {str(e)}") + raise + + def _run_default_config( + self, tune_kernel_args: Dict) -> None: + """Runs the default configuration if enabled.""" + if self.exclude_default_config or self.simulation_mode: + return + + target_arch = self.arch_name + if self.arch_name is None: + warnings.warn(f"Could not detect current architecture!") + return + + if self.existing_config is None: + return + + targets = [ + self.existing_config[k]['arch_name'] for k in self.existing_config if self.existing_config[k]['arch_name'] == target_arch + ] + target = None + if len(targets) > 0: + target = targets[0] + if len(targets) > 1: + warnings.warn( + f"More than one default config for current target {target_arch} found, picking {targets[0]}" + ) + if target is None: + warnings.warn( + f"No appropiate default config for current target {target_arch}!" + ) + return + + arch_config = self.existing_config[target] + + # Get first matching config + config = arch_config['tune_params'] + + if config is None: + warnings.warn( + f"No existing configuration found for {self.algo_name}'" + ) + return + + default_tune_params = { k: [v] for k, v in config.items()} + + # Get the base tuning archs and force set the range of the tune parameters + # to the single-element lists 'default_tune_params'. We also change the + # strategy to bruteforce and clear any set strategy options. + tune_kernel_args = self._get_base_tune_kernel_args().copy() + tune_kernel_args.update( + { + "tune_params": default_tune_params, + "strategy": "brute_force", + "strategy_options": {}, + } + ) + + print(f"Found existing configuration: {default_tune_params}") + kernel_tuner.tune_kernel(**tune_kernel_args) + + def generate_wrapper( + self, + config: OrderedDict, + ) -> str: + """Generate wrapper code using Jinja2 template inheritance.""" + template_dir = pathlib.Path(f"{BASE_DIR}/tuner/templates") + env = Environment( + loader=FileSystemLoader(template_dir), trim_blocks=True, lstrip_blocks=True + ) + + template_name = f"{self.algo_name}.cpp.jinja2" + template = env.get_template(template_name) + # The parameter values are defined using #define by kernel tuner. + # We only need the right symbol name, so 'config.ipt' becomes 'ipt'. + # In the templates, the full '{{ config.ipt }}'-syntax is preferred, + # since this makes the templates more consistent with confgen templates. + context = { + "algo_name": self.algo_name, + "config": config, + } + + content = template.render(**context) + return content + + def _get_base_tune_kernel_args( + self, + ) -> Dict: + """Returns base arguments for kernel_tuner.tune_kernel().""" + np.random.seed(self.seed) + + wrapper_string = lambda config: self.generate_wrapper( + config=config, + ) + + tune_kernel_args = { + "defines": {}, + "kernel_name": f"{self.algo_name}_wrapper", + "kernel_source": wrapper_string, + "arguments": [np.uint64(self.bytes_size)], + "tune_params": self._get_tune_params(), + "strategy": self.strategy, + "grid_div_x": self._get_grid_div_x(), + "cache": str(self._get_cache_file_path()), + "problem_size": 10, + "lang": "C", + "compiler": "hipcc", + "compiler_options": self._get_compiler_options(), + "restrictions": self._get_restrictions(), + "verbose": False, + "iterations": 1, + "log": False, + "device": self.device_id, + "simulation_mode": self.simulation_mode, + "objective_higher_is_better": False, + "quiet": True, + } + + if self.strategy != "brute_force": + tune_kernel_args["strategy_options"] = {"max_fevals": self.max_fevals} + + return tune_kernel_args + + def _get_cache_file_name(self): + """Return the name of the cache file based on algo name, arch name and key value types""" + cache_file_path = f'{self.algo_name}_{self.arch_name}_cache.json' + + return cache_file_path + + def _get_cache_file_path(self): + "Return the path of the cache file" + return self.output_dir / self._get_cache_file_name() + + def _save_output( + self, + cache_file: str | pathlib.Path, + results: List | object | Any | None = None, + ): + """Save tuning results and metadata to files.""" + if self.simulation_mode: + assert results + cache_file = Path( + f"../simulated_output/{self.strategy}_fevals{self.max_fevals}" + ) + cache_file.mkdir(parents=True, exist_ok=True) + cache_file = cache_file / self._get_cache_file_name() + store_output_file(str(cache_file), results, self._get_tune_params()) + + with open(cache_file, "r") as f: + cache_dict = json.load(f) + if "arch_name" not in cache_dict: + with open(cache_file, "r") as f: + content = f.read() + + new_content = "{\n" + new_content += f'"arch_name": "{self.arch_name}",\n' + new_content += f'"algo_name": "{self.algo_name}",\n' + new_content += content.lstrip()[1:] + with open(cache_file, "w") as f: + f.write(new_content) + + if self.save_metadata: + store_metadata_file(str(cache_file).replace("cache", "metadata")) + + def _get_compiler_options(self) -> List[str]: + """Returns a list with all compiler options to pass to Kernel Tuner""" + monorepo_dir = (pathlib.Path(BASE_DIR) / "../../..").resolve() + rocrand_dir = monorepo_dir / "projects/rocrand" + return [ + "-fPIC", + "-std=c++17", + f"-I{rocrand_dir / 'library/include'}", + f"-I{rocrand_dir / 'library/src/'}", + f"-I{rocrand_dir / 'benchmark'}", + f"-I{monorepo_dir / 'shared/primbench'}", + "-Wno-#pragma-messages", + f"--offload-arch={self.arch_name}", + "-lamd_smi", + "-lrocrand", + ] \ No newline at end of file diff --git a/projects/rocrand/tuning/utils/__init__.py b/projects/rocrand/tuning/utils/__init__.py new file mode 100644 index 000000000000..df914b1e5501 --- /dev/null +++ b/projects/rocrand/tuning/utils/__init__.py @@ -0,0 +1,23 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +from .utils import Parser, BASE_DIR diff --git a/projects/rocrand/tuning/utils/compare_simulated_strategies.py b/projects/rocrand/tuning/utils/compare_simulated_strategies.py new file mode 100644 index 000000000000..bcb321042a39 --- /dev/null +++ b/projects/rocrand/tuning/utils/compare_simulated_strategies.py @@ -0,0 +1,257 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import os +import json +import argparse +from typing import Any +import numpy as np +import pandas as pd +import seaborn as sns +import matplotlib.pyplot as plt + +""" +This script performs comparative analysis of different optimization strategies for kernel tuning simulations. +It is designed to work with Kernel Tuner output data to evaluate and visualize the effectiveness of +various tuning strategies. + +Key Features: +- Compares multiple simulated tuning strategies against actual best-known performances +- Calculates relative performance metrics (as percentages) for each strategy +- Generates detailed performance statistics (mean, min, max, median, stddev) +- Creates a visual heatmap representation of strategy performances + +Input Requirements: +- Simulated results should be in '../simulated_output//' directories +- Actual results should be in '../output/' directory +- JSON files must follow naming convention: containing both algorithm and architecture names +- Each JSON file should follow the JSON formats of Kernel Tuner + +Output Formats: +- Console output: Detailed performance statistics for each strategy +- PNG file: Heatmap visualization showing relative performance across all configurations +- Heatmap naming: 'optimization_strategies_heatmap__.png' + +Performance Metrics: +- Values are calculated as: (best_known_time / strategy_time) * 100 +- Higher percentages indicate better performance (closer to optimal solutions) +- 100% means the strategy found the optimal solution +- Values below 100% indicate suboptimal solutions + +Usage: +python3 script.py --arch --algo +Example: python3 script.py --arch gfx90a --algo device_merge +""" + + +def parse_arguments(): + parser = argparse.ArgumentParser( + description="Compare optimization strategies for specific architecture and algorithm." + ) + parser.add_argument( + "--arch", type=str, required=True, help="Architecture name (e.g., gfx90a)" + ) + parser.add_argument( + "--algo", type=str, required=True, help="Algorithm name (e.g., device_merge)" + ) + return parser.parse_args() + + +def is_matching_file(filename, algo, arch): + """ + Check if the JSON file contains both the algorithm name and architecture name. + """ + return algo in filename and arch in filename and filename.endswith("_cache.json") + + +def find_best_time_from_cache(data): + """ + Find the best (lowest) time from results or cache entries. + """ + times = [] + + # For simulated results structure + if "results" in data: + for result in data["results"]: + for measurement in result.get("measurements", []): + if measurement["name"] == "time" and measurement["unit"] == "ms": + time = measurement["value"] + if isinstance(time, (int, float)) and time > 0: + times.append(time) + + # For output results structure with cache + else: + cache = data.get("cache", {}) + for entry_key, entry_data in cache.items(): + time = entry_data.get("time", float("inf")) + if isinstance(time, (int, float)) and time > 0: + times.append(time) + + return min(times) if times else float("inf") + + +def load_json_files(simulated_output_dir, output_dir, algo, arch): + """ + Load and process JSON files from both simulated and actual output directories. + Returns best times for each kernel configuration. + """ + results = {"simulated": {}, "best_times": {}} + + # First, process actual output files to get best times + for json_file in os.listdir(output_dir): + if json_file.endswith(".json") and is_matching_file(json_file, algo, arch): + with open(os.path.join(output_dir, json_file), "r") as f: + data = json.load(f) + best_time = find_best_time_from_cache(data) + results["best_times"][json_file] = best_time + + # Process simulated output directories (different strategies) + for strategy in os.listdir(simulated_output_dir): + strategy_path = os.path.join(simulated_output_dir, strategy) + if os.path.isdir(strategy_path): + results["simulated"][strategy] = {} + for json_file in os.listdir(strategy_path): + if json_file.endswith(".json") and is_matching_file( + json_file, algo, arch + ): + with open(os.path.join(strategy_path, json_file), "r") as f: + data = json.load(f) + best_time = find_best_time_from_cache(data) + results["simulated"][strategy][json_file] = best_time + + if not results["best_times"]: + raise ValueError( + f"No matching files found for algorithm '{algo}' and architecture '{arch}' in {output_dir}" + ) + + if not any(results["simulated"].values()): + raise ValueError( + f"No matching files found in any strategy directory for algorithm '{algo}' and architecture '{arch}'" + ) + + return results + + +def calculate_performance_stats(results: dict[str, Any], algo, arch): + """ + Calculate and print performance statistics for each strategy. + Returns the DataFrame for the heatmap. + """ + strategies = sorted(list(results["simulated"].keys())) + kernels = sorted(list(results["best_times"].keys())) + + data: list[list[float]] = [] + for strategy in strategies: + row: list[float] = [] + for kernel in kernels: + strategy_time: float = results["simulated"][strategy][kernel] + best_time: float = results["best_times"][kernel] + relative_perf = (best_time / strategy_time) * 100 + row.append(relative_perf) + data.append(row) + + df = pd.DataFrame(data, index=np.array(strategies), columns=np.array(kernels)) + + # Calculate summary statistics + summary_stats = pd.DataFrame( + { + "Mean": df.mean(axis=1).round(1), + "Min": df.min(axis=1).round(1), + "Max": df.max(axis=1).round(1), + "Median": df.median(axis=1).round(1), + "Std Dev": df.std(axis=1).round(1), + } + ).sort_values("Mean", ascending=False) + + # Print performance summary + print(f"\nOptimization Strategy Performance Summary for {algo} on {arch}") + print("=" * 80) + print("\nPerformance statistics (%):") + print(summary_stats.to_string()) + print("\nNote: Higher values indicate better performance") + print("=" * 80) + + return df + + +def create_performance_heatmap(df, output_filename, algo, arch): + """ + Create a heatmap comparing different strategies' performance. + """ + # Create figure + plt.figure(figsize=(20, 8)) + + # Sort DataFrame rows by mean performance (descending) + df = df.reindex(df.mean(axis=1).sort_values(ascending=False).index) + + # Normalize kernel names for better display + df.columns = [ + col.replace(f"{algo}_", "").replace("_cache.json", "") for col in df.columns + ] + + min_val = df.min().min() + max_val = df.max().max() + + # Create heatmap with adjusted parameters + sns.heatmap( + df, + annot=True, # Show numbers in cells + fmt=".1f", # Format numbers to 1 decimal place + cmap="RdYlGn", # Red (bad/low) to Green (good/high) + vmin=min_val, # Use actual minimum value + vmax=max_val, # Use actual maximum value + center=None, # Remove center to maximize color spread + robust=True, # Use robust quantiles for color scaling + square=True, # Make cells square + cbar_kws={"label": "Performance relative to best (%)"}, + ) + + plt.title( + f"Optimization Strategies Performance for {algo} on {arch}\nValues range: {min_val:.1f}% - {max_val:.1f}%" + ) + plt.xlabel("Kernel Configurations") + plt.ylabel("Optimization Strategies") + plt.xticks(rotation=45, ha="right") + plt.tight_layout() + plt.savefig(output_filename, dpi=300, bbox_inches="tight") + plt.close() + + +def main(): + args = parse_arguments() + + base_dir = "." + simulated_output_dir = os.path.join(base_dir, "../simulated_output") + output_dir = os.path.join(base_dir, "../output") + + results = load_json_files(simulated_output_dir, output_dir, args.algo, args.arch) + + # Calculate and print performance statistics + df = calculate_performance_stats(results, args.algo, args.arch) + + # Create heatmap + output_filename = f"optimization_strategies_heatmap_{args.algo}_{args.arch}.png" + create_performance_heatmap(df, output_filename, args.algo, args.arch) + + +if __name__ == "__main__": + main() diff --git a/projects/rocrand/tuning/utils/utils.py b/projects/rocrand/tuning/utils/utils.py new file mode 100644 index 000000000000..6bb0e7d3a5bc --- /dev/null +++ b/projects/rocrand/tuning/utils/utils.py @@ -0,0 +1,92 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import argparse +from typing import Optional +import os + +BASE_DIR = f'{os.path.dirname(os.path.abspath(__file__))}/..' + +class Parser: + @staticmethod + def get_parser( + default_bytes: Optional[int] = None, + default_max_feval: int = 100, + default_output_dir: str = "../output", + default_strategy: str = "dual_annealing", + ): + parser = argparse.ArgumentParser() + + parser.add_argument( + "--size", + type=int, + default=default_bytes, + help=f"Size in bytes (default: {default_bytes})", + ) + parser.add_argument( + "--max-fevals", + type=int, + default=default_max_feval, + help=f"Maximum number of unique valid function evaluations (default: {default_max_feval})", + ) + parser.add_argument( + "--output-dir", + type=str, + default=default_output_dir, + help=f"Output directory for JSON files (default: {default_output_dir})", + ) + parser.add_argument( + "--exclude-default-config", + action="store_true", + default=False, + help="Include default configs of previous tuning in the new results (default: off)", + ) + parser.add_argument( + "--strategy", + type=str, + default=default_strategy, + help=f"Strategy Kernel Tuner will use (default: {default_strategy})", + ) + parser.add_argument( + "--simulation-mode", + action="store_true", + default=False, + help="Run strategy in simulation mode (you need the cache files of the full searchspace)", + ) + parser.add_argument( + "--arch-name", + required=False, + help="Specify arch, needed when running in simulation", + ) + parser.add_argument( + "--save-metadata", + action="store_true", + default=False, + help="Save Kernel Tuner metadata (default: off)", + ) + parser.add_argument( + "--seed", + type=int, + default=-1, + help=f"The initial seed to generate the input data with", + ) + return parser From 6a8e1128935d8f6faedbd296d67f1ff46fac96bf Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:05:20 -0400 Subject: [PATCH 02/38] feat(rocrand): ported over generate.py and updated it for rocrand --- projects/rocrand/tuning/confgen/generate.py | 162 ++++++++++++++++++++ 1 file changed, 162 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/generate.py diff --git a/projects/rocrand/tuning/confgen/generate.py b/projects/rocrand/tuning/confgen/generate.py new file mode 100644 index 000000000000..6b24c0a8aea9 --- /dev/null +++ b/projects/rocrand/tuning/confgen/generate.py @@ -0,0 +1,162 @@ +#!/usr/bin/env python3 + +import argparse +import datetime +import glob +import json +import logging +import os +import pathlib + +import typing_extensions as ty +import jinja2 + +log = logging.getLogger("confgen.generate") + +def make_config( + type_data: dict[str, ty.Any]): + import_keys = ["arch_name", "tune_params_keys", "tune_params"] + out = {k : type_data[k] for k in import_keys if k in type_data} + return out + +def main(): + cli = argparse.ArgumentParser() + cli.add_argument( + "--input", + "-i", + required=True, + help="Directory or file glob-pattern that points to kernel tuner results", + ) + cli.add_argument( + "--output", "-o", required=True, help="The output directory for config headerse" + ) + cli.add_argument( + "--existing", + "-e", + help="The directory with current config headers to merge new configs with.", + ) + cli.add_argument( + "--target-arch", + "-t", + action="append", + help="The target arch that these configs are for.", + ) + + args = cli.parse_args() + + gfx_target = [] + if args.target_arch: + gfx_target = args.target_arch + + # Read '--input' argument and suffix with '.json' if it's not '.json'. + input_glob = args.input + if not input_glob.endswith(".json"): + input_glob = f"{input_glob}/*.json" + + # Find all files. + file_paths = glob.glob(input_glob) + + algs: dict[str, dict[frozenset, dict[frozenset, dict[str, ty.Any]]]] = {} + for file_path in file_paths: + log.debug(f"Parsing: {file_path}") + with open(file_path, "r") as f: + try: + data = json.load(f) + except json.decoder.JSONDecodeError: + log.warning(f"Skipping due to JSONDecodeError: {file_path}") + continue + + # Filter/skip data entries that have no "time" column + data_filter = [c for c in data["cache"].values() if "time" in c] + filter_dif = len(data["cache"]) - len(data_filter) + if filter_dif > 0: + log.warning(f"Skipping failed compilations: {filter_dif}") + # Find best config + min_configs = min(data_filter, key=lambda c: c["time"]) + # Drop unrelated entries + config_ignored_names = [ + "time", + "times", + "compile_time", + "verification_time", + "framework_time", + "benchmark_time", + "strategy_time", + "timestamp", + ] + min_configs = { + k: min_configs[k] for k in min_configs if k not in config_ignored_names + } + + # Ensure algorithm entry exists + alg: str = data["algo_name"] + if alg not in algs: + algs[alg] = {} + + for target in gfx_target: + # Ensure target entry exists + if target not in algs[alg]: + algs[alg][target] = {} + + config = make_config(data) + config['tune_params'] = min_configs + # Add config entry + algs[alg][target] = config + + log.info(f"Parsed {len(algs)} different algorithm(s)!") + + script_dir = pathlib.Path(__file__).parent.absolute() + env = jinja2.Environment( + loader=jinja2.FileSystemLoader(script_dir / "templates"), + autoescape=jinja2.select_autoescape(), + trim_blocks=True, + lstrip_blocks=True + ) + + # Find existing configs + existing_dir = None + if args.existing: + existing_dir = pathlib.Path(args.existing) + if not existing_dir.exists(): + log.error(f"Could not find directory: {existing_dir}") + quit() + + # Merge newly generated configs to existing configs + if existing_dir: + import parse as config_parser + + for algo_name, targets in algs.items(): + config_file = existing_dir / f"{algo_name}_config.hpp" + log.debug(f"Parsing: {config_file}") + with open(config_file, "r") as file: + existing_data = config_parser.parse_lines(file.readlines()) + + merged_targets = {target_hash : make_config(target) for target_hash, target in existing_data.items()} + + for target_hash, config in targets.items(): + merged_targets[target_hash] = config + algs[algo_name] = merged_targets + + # Output generated configs + output_dir = pathlib.Path(args.output) + os.makedirs(output_dir, exist_ok=True) + for algo_name, configs in algs.items(): + + template_name = f"{algo_name}.h.jinja2" + template = env.get_template(template_name) + rendered = template.render( + { + "year": datetime.datetime.now().year, + "targets": configs, + "algo_name": algo_name, + } + ) + + with open(output_dir / f"{algo_name}_config.hpp", "w") as file: + file.write(rendered) + file.write(f'\n') + +if __name__ == "__main__": + logging.basicConfig(level=logging.DEBUG) + main() + \ No newline at end of file From e32717e692c70caa0605361898dc929c03c3a04b Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:05:42 -0400 Subject: [PATCH 03/38] feat(rocrand): ported over parse.py and updated it for rocrand --- projects/rocrand/tuning/confgen/parse.py | 36 ++++++++++++++++++++++++ 1 file changed, 36 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/parse.py diff --git a/projects/rocrand/tuning/confgen/parse.py b/projects/rocrand/tuning/confgen/parse.py new file mode 100644 index 000000000000..88b2fe02b86d --- /dev/null +++ b/projects/rocrand/tuning/confgen/parse.py @@ -0,0 +1,36 @@ +#!/usr/bin/env python3 + +def parse_lines(lines: list[str]): + + config = {} + + INSIDE_BLOCK = False + INSIDE_GRID = False + + for line in lines: + if 'ROCRAND_RNG_PSEUDO' in line and 'generator_config_selector' in line: + engine_name = line.split('<')[-1].split(',')[0].split('_')[-1].lower() + + if 'get_threads' in line: + INSIDE_BLOCK = True + if 'get_blocks' in line: + INSIDE_BLOCK = False + INSIDE_GRID = True + + if 'target_arch::' in line: + arch = (line.split()[1].split(':')[-2]) + if arch not in config: + config[arch] = {"arch_name" : arch, "engine" : engine_name} + if INSIDE_BLOCK: + block_size = int(line.split()[-1][:-1]) + + config[arch]['tune_param_keys'] = ['block_size_x'] + config[arch]['tune_params'] = {"block_size_x" : block_size} + + elif INSIDE_GRID: + grid_size = int(line.split()[-1][:-1]) + + config[arch]['tune_param_keys'].append('grid_size') + config[arch]['tune_params']["grid_size"] = grid_size + + return config From 7d641debd93129227e665403068ea6a1a4c86061 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:06:20 -0400 Subject: [PATCH 04/38] feat(rocrand): wrote the base template for rocrand config gen --- .../confgen/templates/common/base.h.jinja2 | 69 +++++++++++++++++++ 1 file changed, 69 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/common/base.h.jinja2 diff --git a/projects/rocrand/tuning/confgen/templates/common/base.h.jinja2 b/projects/rocrand/tuning/confgen/templates/common/base.h.jinja2 new file mode 100644 index 000000000000..c200edbb3838 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/common/base.h.jinja2 @@ -0,0 +1,69 @@ +// Copyright (c) {{year}} Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +{% block header_guard_pre -%} +#ifndef ROCRAND_RNG_CONFIG_{{algo_name|upper}}_CONFIG_HPP_ +#define ROCRAND_RNG_CONFIG_{{algo_name|upper}}_CONFIG_HPP_ +{%- endblock %} + +{% block header_includes -%} +#include "config_defaults.hpp" +#include +{%- endblock %} + +{% block generated_file_warning -%} +/* DO NOT EDIT THIS FILE + * This file is automatically generated by `/tuning/confgen/generate.py`. + */ +{%- endblock %} + +namespace rocrand_impl::host +{ +template +struct generator_config_selector +{ + __host__ __device__ + static constexpr unsigned int get_threads(const target_arch arch) + { + switch(arch) + { + {% for key, value in targets.items() %} + case target_arch::{{key}}: return {{value.tune_params.block_size_x}}; + {% endfor %} + default: return generator_config_defaults::threads; + } + } + + __host__ __device__ + static constexpr unsigned int get_blocks(const target_arch arch) + { + switch(arch) + { + {% for key, value in targets.items() %} + case target_arch::{{key}}: return {{value.tune_params.grid_size}}; + {% endfor %} + default: return generator_config_defaults::blocks; + } + } +}; + +} // end namespace rocrand_impl::host + +#endif // ROCRAND_RNG_CONFIG_{{algo_name|upper}}_CONFIG_HPP_ From 01d5dedec260d9c1d4f3c05fc8f4d15cde09a541 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:08:20 -0400 Subject: [PATCH 05/38] feat(rocrand): implement the base tuner for kernel tuning --- .../tuner/templates/common/base.cpp.jinja2 | 218 ++++++++++++++++++ 1 file changed, 218 insertions(+) create mode 100644 projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 diff --git a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 new file mode 100644 index 000000000000..100e9485d43c --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 @@ -0,0 +1,218 @@ +// Include rocRAND's benchmark +#include "benchmark_host_api.hpp" + +// Include primbench executor. +#include "primbench.hpp" + +{% block backend_include %} +// Include the right backend + #include "rng/{{algo_name}}.hpp" +{% endblock %} +#include +#include +#include + +using config_provider = rocrand_impl::host::static_config_provider<{{ config.block_size_x }}, {{ config.grid_size }}>; + +{% block generator_define %} +template +using generator_template + = rocrand_impl::host::{{ algo_name }}_generator_template; +{% endblock %} + + +// Disable ULL +template class GeneratorTemplate> +struct output_type_supported : std::true_type +{}; + +#define DISABLE_ULL(Generator) \ + template<> \ + struct output_type_supported : std::false_type \ + {} + +DISABLE_ULL(generator_template); +#undef DISABLE_ULL + +// QUEUING TEMPLATE +template class GeneratorTemplate, class ConfigProvider> +void queue_one(primbench::executor& executor, + size_t dimensions, + size_t offset, + bool benchmark_host, + double poisson_lambda) +{ + using generator_t = GeneratorTemplate; + constexpr auto rng_type = generator_t::type(); + + if constexpr(Dist == DISTRIBUTION_POISSON) + { + executor.queue>(rng_type, + RAND_ORDERING_PSEUDO_DYNAMIC, + dimensions, + offset, + benchmark_host, + poisson_lambda); + } + else + { + executor.queue>(rng_type, + RAND_ORDERING_PSEUDO_DEFAULT, + dimensions, + offset, + benchmark_host); + } +} + +template class GeneratorTemplate, class ConfigProvider> +void queue_distributions(primbench::executor& executor, + size_t dimensions, + size_t offset, + bool benchmark_host, + const std::vector& poisson_lambdas) +{ + if constexpr(!output_type_supported::value) + { + return; + } + else if constexpr(std::is_integral_v) + { + queue_one(executor, + dimensions, + offset, + benchmark_host, + 0.0); + + if constexpr(std::is_same_v) + { + for(double lambda : poisson_lambdas) + { + queue_one( + executor, + dimensions, + offset, + benchmark_host, + lambda); + } + } + } + else + { + queue_one(executor, + dimensions, + offset, + benchmark_host, + 0.0); + queue_one(executor, + dimensions, + offset, + benchmark_host, + 0.0); + queue_one(executor, + dimensions, + offset, + benchmark_host, + 0.0); + } +} + +template class GeneratorTemplate, class ConfigProvider> +void queue_all_types(primbench::executor& executor, + size_t dimensions, + size_t offset, + bool benchmark_host, + const std::vector& poisson_lambdas) +{ + queue_distributions(executor, + dimensions, + offset, + benchmark_host, + poisson_lambdas); + queue_distributions(executor, + dimensions, + offset, + benchmark_host, + poisson_lambdas); + queue_distributions(executor, + dimensions, + offset, + benchmark_host, + poisson_lambdas); + queue_distributions(executor, + dimensions, + offset, + benchmark_host, + poisson_lambdas); + queue_distributions(executor, + dimensions, + offset, + benchmark_host, + poisson_lambdas); + queue_distributions(executor, + dimensions, + offset, + benchmark_host, + poisson_lambdas); + queue_distributions(executor, + dimensions, + offset, + benchmark_host, + poisson_lambdas); +} + +{% block extern_c_wrapper %} +extern "C" float {{algo_name}}_wrapper( + {% block arguments %} + uint64_t bytes_size + {% endblock %} +) +{ + // Flush is required for fancy primbench printing + std::cout << std::flush; + + {% block primbench_settings %} + auto settings = primbench::settings(); + settings.json_out = "/dev/null"; + settings.max_warming_secs = 90; + settings.max_cooling_secs = 90; + settings.size = bytes_size; + settings.skip_header = true; + {% endblock %} + + {% block primbench_args %} + char* args[] = { "benchmark-executable" /* primbench requires an executable name for exposing --help */ }; + {% endblock %} + + {% block primbench_executor %} + auto executor = primbench::executor(1, args, settings, primbench::flags::sync); + {% endblock %} + + {% block set_up_settings %} + auto dimensions = 1; + auto offset = 0; + auto poisson_lambdas = {10.0}; + auto benchmark_host = false; + {% endblock %} + + queue_all_types(executor, + dimensions, + offset, + benchmark_host, + poisson_lambdas); + + + executor.run(); + + {% block calculate_time %} + std::vector bps = executor.get_all_bytes_per_second(); + + double total = 0.0; + for(double x : bps) + total += (1000.0 * primbench::GiB) / x; + + return total / static_cast(bps.size()); + {% endblock %} + +} +{% endblock %} \ No newline at end of file From 32161180c6c1ff39b2095cb461706091ef1630ff Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:09:10 -0400 Subject: [PATCH 06/38] feat(rocrand): implement kernel tuner for lfsr113 --- .../tuning/confgen/templates/lfsr113.h.jinja2 | 1 + .../tuning/tuner/templates/lfsr113.cpp.jinja2 | 1 + .../rocrand/tuning/tuner/tuning_lfsr113.py | 39 +++++++++++++++++++ 3 files changed, 41 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/lfsr113.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/lfsr113.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_lfsr113.py diff --git a/projects/rocrand/tuning/confgen/templates/lfsr113.h.jinja2 b/projects/rocrand/tuning/confgen/templates/lfsr113.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/lfsr113.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/lfsr113.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/lfsr113.cpp.jinja2 new file mode 100644 index 000000000000..01c55291f9b0 --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/lfsr113.cpp.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.cpp.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/tuning_lfsr113.py b/projects/rocrand/tuning/tuner/tuning_lfsr113.py new file mode 100644 index 000000000000..9c02d89ae197 --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_lfsr113.py @@ -0,0 +1,39 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='lfsr113') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + +if __name__ == "__main__": + Tuner.cli() From ce25602c26cea45fad22f20de73a1db3babf6321 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:10:06 -0400 Subject: [PATCH 07/38] feat(rocrand): implement kernel tuner for mrg31k3p --- .../confgen/templates/mrg31k3p.h.jinja2 | 1 + .../tuner/templates/mrg31k3p.cpp.jinja2 | 13 ++++++ .../rocrand/tuning/tuner/tuning_mrg31k3p.py | 40 +++++++++++++++++++ 3 files changed, 54 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/mrg31k3p.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/mrg31k3p.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_mrg31k3p.py diff --git a/projects/rocrand/tuning/confgen/templates/mrg31k3p.h.jinja2 b/projects/rocrand/tuning/confgen/templates/mrg31k3p.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/mrg31k3p.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/mrg31k3p.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/mrg31k3p.cpp.jinja2 new file mode 100644 index 000000000000..0343bfde67b4 --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/mrg31k3p.cpp.jinja2 @@ -0,0 +1,13 @@ +{% extends "common/base.cpp.jinja2" %} + +{% block backend_include %} + #include "rng/mrg.hpp" +{% endblock %} + +{% block generator_define %} + template + using generator_template + = rocrand_impl::host::mrg_generator_template; +{% endblock %} diff --git a/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py b/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py new file mode 100644 index 000000000000..836fb4acbaa0 --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py @@ -0,0 +1,40 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='mrg31k3p') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + +if __name__ == "__main__": + Tuner.cli() + \ No newline at end of file From 899b77ec8fdf654fd1770c71eb430ab2606712eb Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:10:39 -0400 Subject: [PATCH 08/38] feat(rocrand): implement kernel tuner for mrg32k3a --- .../confgen/templates/mrg32k3a.h.jinja2 | 1 + .../tuner/templates/mrg32k3a.cpp.jinja2 | 13 ++++++ .../rocrand/tuning/tuner/tuning_mrg32k3a.py | 40 +++++++++++++++++++ 3 files changed, 54 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/mrg32k3a.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/mrg32k3a.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_mrg32k3a.py diff --git a/projects/rocrand/tuning/confgen/templates/mrg32k3a.h.jinja2 b/projects/rocrand/tuning/confgen/templates/mrg32k3a.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/mrg32k3a.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/mrg32k3a.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/mrg32k3a.cpp.jinja2 new file mode 100644 index 000000000000..b727382436a0 --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/mrg32k3a.cpp.jinja2 @@ -0,0 +1,13 @@ +{% extends "common/base.cpp.jinja2" %} + +{% block backend_include %} + #include "rng/mrg.hpp" +{% endblock %} + +{% block generator_define %} + template + using generator_template + = rocrand_impl::host::mrg_generator_template; +{% endblock %} diff --git a/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py b/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py new file mode 100644 index 000000000000..de36e33d6923 --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py @@ -0,0 +1,40 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='mrg32k3a') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + +if __name__ == "__main__": + Tuner.cli() + \ No newline at end of file From 5d3b74af3a58209a854dc55d3b1e306dea8ddbfd Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:11:26 -0400 Subject: [PATCH 09/38] feat(rocrand): implement kernel tuner for mt19937 --- .../tuning/confgen/templates/mt19937.h.jinja2 | 1 + .../tuning/tuner/templates/mt19937.cpp.jinja2 | 6 ++ .../rocrand/tuning/tuner/tuning_mt19937.py | 55 +++++++++++++++++++ 3 files changed, 62 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/mt19937.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/mt19937.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_mt19937.py diff --git a/projects/rocrand/tuning/confgen/templates/mt19937.h.jinja2 b/projects/rocrand/tuning/confgen/templates/mt19937.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/mt19937.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/mt19937.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/mt19937.cpp.jinja2 new file mode 100644 index 000000000000..19e606358f23 --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/mt19937.cpp.jinja2 @@ -0,0 +1,6 @@ +{% extends "common/base.cpp.jinja2" %} + +{% block backend_include %} + #include "rng/mt19937.hpp" + #include "rocrand/rocrand_mt19937_precomputed.h" +{% endblock %} diff --git a/projects/rocrand/tuning/tuner/tuning_mt19937.py b/projects/rocrand/tuning/tuner/tuning_mt19937.py new file mode 100644 index 000000000000..353ce80b2b5e --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_mt19937.py @@ -0,0 +1,55 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='mt19937') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + + def _get_restrictions(self): + min_total_threads = 32768 + mt19937_jumps_radix = 256 + threads_per_generator = 8 + def validate(params): + threads = params['block_size_x'] + blocks = params['grid_size'] + + if threads * blocks < min_total_threads: + return False + + return (threads * blocks) // threads_per_generator <= (mt19937_jumps_radix * mt19937_jumps_radix) + + return validate + +if __name__ == "__main__": + Tuner.cli() + \ No newline at end of file From ffc23e7594054cc41df3ce059112a70c99fd712b Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:12:01 -0400 Subject: [PATCH 10/38] feat(rocrand): implement kernel tuner for mtgp32 --- .../tuning/confgen/templates/mtgp32.h.jinja2 | 1 + .../tuning/tuner/templates/mtgp32.cpp.jinja2 | 1 + .../rocrand/tuning/tuner/tuning_mtgp32.py | 54 +++++++++++++++++++ 3 files changed, 56 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/mtgp32.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/mtgp32.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_mtgp32.py diff --git a/projects/rocrand/tuning/confgen/templates/mtgp32.h.jinja2 b/projects/rocrand/tuning/confgen/templates/mtgp32.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/mtgp32.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/mtgp32.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/mtgp32.cpp.jinja2 new file mode 100644 index 000000000000..01c55291f9b0 --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/mtgp32.cpp.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.cpp.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/tuning_mtgp32.py b/projects/rocrand/tuning/tuner/tuning_mtgp32.py new file mode 100644 index 000000000000..259f8fb75926 --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_mtgp32.py @@ -0,0 +1,54 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='mtgp32') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + + def _get_restrictions(self): + min_total_threads = 32768 + def validate(params): + threads = params['block_size_x'] + blocks = params['grid_size'] + + if threads * blocks < min_total_threads: + return False + + + return blocks <= 512 and threads == 256 + + return validate + +if __name__ == "__main__": + Tuner.cli() + \ No newline at end of file From f0210455c6fb95c125d4fe2a10fcb18de16fc428 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:12:36 -0400 Subject: [PATCH 11/38] feat(rocrand): implement kernel tuner for philox4_32_10 --- .../confgen/templates/philox4_32_10.h.jinja2 | 1 + .../tuner/templates/philox4_32_10.cpp.jinja2 | 12 ++++++ .../tuning/tuner/tuning_philox4_32_10.py | 40 +++++++++++++++++++ 3 files changed, 53 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/philox4_32_10.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/philox4_32_10.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_philox4_32_10.py diff --git a/projects/rocrand/tuning/confgen/templates/philox4_32_10.h.jinja2 b/projects/rocrand/tuning/confgen/templates/philox4_32_10.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/philox4_32_10.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/philox4_32_10.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/philox4_32_10.cpp.jinja2 new file mode 100644 index 000000000000..b245027857a6 --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/philox4_32_10.cpp.jinja2 @@ -0,0 +1,12 @@ +{% extends "common/base.cpp.jinja2" %} + +{% block backend_include %} + #include "rng/philox4x32_10.hpp" +{% endblock %} + +{% block generator_define %} + template + using generator_template + = rocrand_impl::host::philox4x32_10_generator_template; +{% endblock %} diff --git a/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py b/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py new file mode 100644 index 000000000000..1465c47bc486 --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py @@ -0,0 +1,40 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='philox4_32_10') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + +if __name__ == "__main__": + Tuner.cli() + \ No newline at end of file From 9f97d4941bfd32d6320150943a095b8d5824e964 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:13:16 -0400 Subject: [PATCH 12/38] feat(rocrand): implement kernel tuner for threefry2_32_20 --- .../templates/threefry2_32_20.h.jinja2 | 1 + .../templates/threefry2_32_20.cpp.jinja2 | 14 +++++++ .../tuning/tuner/tuning_threefry2_32_20.py | 40 +++++++++++++++++++ 3 files changed, 55 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/threefry2_32_20.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/threefry2_32_20.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py diff --git a/projects/rocrand/tuning/confgen/templates/threefry2_32_20.h.jinja2 b/projects/rocrand/tuning/confgen/templates/threefry2_32_20.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/threefry2_32_20.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/threefry2_32_20.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/threefry2_32_20.cpp.jinja2 new file mode 100644 index 000000000000..9cc4cbfebe24 --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/threefry2_32_20.cpp.jinja2 @@ -0,0 +1,14 @@ +{% extends "common/base.cpp.jinja2" %} + +{% block backend_include %} + #include "rng/threefry.hpp" +{% endblock %} + +{% block generator_define %} + template + using generator_template + = rocrand_impl::host::threefry_generator_template< + rocrand_impl::system::device_system, + rocrand_impl::host::threefry_device_engine, + ConfigProvider>; +{% endblock %} diff --git a/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py b/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py new file mode 100644 index 000000000000..efd8b35b9693 --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py @@ -0,0 +1,40 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='threefry2_32_20') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + +if __name__ == "__main__": + Tuner.cli() + \ No newline at end of file From 33a32a8056e3871a64ce0dd0c8bccff172da4751 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:13:35 -0400 Subject: [PATCH 13/38] feat(rocrand): implement kernel tuner for threefry2_64_20 --- .../templates/threefry2_64_20.h.jinja2 | 1 + .../templates/threefry2_64_20.cpp.jinja2 | 14 +++++++ .../tuning/tuner/tuning_threefry2_64_20.py | 40 +++++++++++++++++++ 3 files changed, 55 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/threefry2_64_20.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py diff --git a/projects/rocrand/tuning/confgen/templates/threefry2_64_20.h.jinja2 b/projects/rocrand/tuning/confgen/templates/threefry2_64_20.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/threefry2_64_20.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 new file mode 100644 index 000000000000..fa561c571b1f --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 @@ -0,0 +1,14 @@ +{% extends "common/base.cpp.jinja2" %} + +{% block backend_include %} + #include "rng/threefry.hpp" +{% endblock %} + +{% block generator_define %} + template + using generator_template + = rocrand_impl::host::threefry_generator_template< + rocrand_impl::system::device_system, + rocrand_impl::host::threefry_device_engine, + ConfigProvider>; +{% endblock %} diff --git a/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py b/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py new file mode 100644 index 000000000000..d54a0a8b7a3b --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py @@ -0,0 +1,40 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='threefry2_64_20') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + +if __name__ == "__main__": + Tuner.cli() + \ No newline at end of file From 8482551479ceca64f42543369a3ef8c0b3b93b58 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:13:59 -0400 Subject: [PATCH 14/38] feat(rocrand): implement kernel tuner for threefry4_32_20 --- .../templates/threefry4_32_20.h.jinja2 | 1 + .../templates/threefry4_32_20.cpp.jinja2 | 14 +++++++ .../tuning/tuner/tuning_threefry4_32_20.py | 40 +++++++++++++++++++ 3 files changed, 55 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/threefry4_32_20.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/threefry4_32_20.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py diff --git a/projects/rocrand/tuning/confgen/templates/threefry4_32_20.h.jinja2 b/projects/rocrand/tuning/confgen/templates/threefry4_32_20.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/threefry4_32_20.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/threefry4_32_20.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/threefry4_32_20.cpp.jinja2 new file mode 100644 index 000000000000..e88b599ba1c0 --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/threefry4_32_20.cpp.jinja2 @@ -0,0 +1,14 @@ +{% extends "common/base.cpp.jinja2" %} + +{% block backend_include %} + #include "rng/threefry.hpp" +{% endblock %} + +{% block generator_define %} + template + using generator_template + = rocrand_impl::host::threefry_generator_template< + rocrand_impl::system::device_system, + rocrand_impl::host::threefry_device_engine, + ConfigProvider>; +{% endblock %} diff --git a/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py b/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py new file mode 100644 index 000000000000..c460a21aeae8 --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py @@ -0,0 +1,40 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='threefry4_32_20') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + +if __name__ == "__main__": + Tuner.cli() + \ No newline at end of file From 3b27607764c2d95023896a067628d10f347b2fc7 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:14:22 -0400 Subject: [PATCH 15/38] feat(rocrand): implement kernel tuner for threefry4_64_20 --- .../templates/threefry4_64_20.h.jinja2 | 1 + .../templates/threefry4_64_20.cpp.jinja2 | 14 +++++++ .../tuning/tuner/tuning_threefry4_64_20.py | 40 +++++++++++++++++++ 3 files changed, 55 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/threefry4_64_20.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py diff --git a/projects/rocrand/tuning/confgen/templates/threefry4_64_20.h.jinja2 b/projects/rocrand/tuning/confgen/templates/threefry4_64_20.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/threefry4_64_20.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 new file mode 100644 index 000000000000..010c548181ac --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 @@ -0,0 +1,14 @@ +{% extends "common/base.cpp.jinja2" %} + +{% block backend_include %} + #include "rng/threefry.hpp" +{% endblock %} + +{% block generator_define %} + template + using generator_template + = rocrand_impl::host::threefry_generator_template< + rocrand_impl::system::device_system, + rocrand_impl::host::threefry_device_engine, + ConfigProvider>; +{% endblock %} diff --git a/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py b/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py new file mode 100644 index 000000000000..07bf31f892bd --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py @@ -0,0 +1,40 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='threefry4_64_20') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + +if __name__ == "__main__": + Tuner.cli() + \ No newline at end of file From eacfce5a012b3d5295130974c7eabaa5f20e6e5a Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:14:45 -0400 Subject: [PATCH 16/38] feat(rocrand): implement kernel tuner for xorwow --- .../tuning/confgen/templates/xorwow.h.jinja2 | 1 + .../tuning/tuner/templates/xorwow.cpp.jinja2 | 1 + .../rocrand/tuning/tuner/tuning_xorwow.py | 39 +++++++++++++++++++ 3 files changed, 41 insertions(+) create mode 100644 projects/rocrand/tuning/confgen/templates/xorwow.h.jinja2 create mode 100644 projects/rocrand/tuning/tuner/templates/xorwow.cpp.jinja2 create mode 100644 projects/rocrand/tuning/tuner/tuning_xorwow.py diff --git a/projects/rocrand/tuning/confgen/templates/xorwow.h.jinja2 b/projects/rocrand/tuning/confgen/templates/xorwow.h.jinja2 new file mode 100644 index 000000000000..7003b499e2f3 --- /dev/null +++ b/projects/rocrand/tuning/confgen/templates/xorwow.h.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.h.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/templates/xorwow.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/xorwow.cpp.jinja2 new file mode 100644 index 000000000000..01c55291f9b0 --- /dev/null +++ b/projects/rocrand/tuning/tuner/templates/xorwow.cpp.jinja2 @@ -0,0 +1 @@ +{% extends "common/base.cpp.jinja2" %} diff --git a/projects/rocrand/tuning/tuner/tuning_xorwow.py b/projects/rocrand/tuning/tuner/tuning_xorwow.py new file mode 100644 index 000000000000..31d564a255f9 --- /dev/null +++ b/projects/rocrand/tuning/tuner/tuning_xorwow.py @@ -0,0 +1,39 @@ +#!/usr/bin/env python3 + +# Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +import sys +import os + +sys.path.append(f"{os.path.dirname(__file__)}/../") + +from tuner.base_tuner import BaseTuner, TunerArgs + +class Tuner(BaseTuner): + @classmethod + def _get_default_args(cls) -> TunerArgs: + return TunerArgs(algo_full_name='xorwow') + + def __init__(self, args: TunerArgs) -> None: + super().__init__(args) + +if __name__ == "__main__": + Tuner.cli() From baa1c50c806be6bd60124f8efbef1e36e8bba7b9 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:15:05 -0400 Subject: [PATCH 17/38] feat(rocrand): wrote README and added gitignore --- .../library/src/rng/config/lfsr113_config.hpp | 10 +- .../src/rng/config/mrg31k3p_config.hpp | 10 +- .../src/rng/config/mrg32k3a_config.hpp | 10 +- .../library/src/rng/config/mt19937_config.hpp | 12 +- .../library/src/rng/config/mtgp32_config.hpp | 10 +- .../src/rng/config/philox4_32_10_config.hpp | 10 +- .../src/rng/config/threefry2_32_20_config.hpp | 10 +- .../src/rng/config/threefry2_64_20_config.hpp | 10 +- .../src/rng/config/threefry4_32_20_config.hpp | 10 +- .../src/rng/config/threefry4_64_20_config.hpp | 10 +- .../library/src/rng/config/xorwow_config.hpp | 10 +- projects/rocrand/tuning/.gitignore | 5 + projects/rocrand/tuning/README.md | 133 ++++++++++++++++++ 13 files changed, 172 insertions(+), 78 deletions(-) create mode 100644 projects/rocrand/tuning/.gitignore create mode 100644 projects/rocrand/tuning/README.md diff --git a/projects/rocrand/library/src/rng/config/lfsr113_config.hpp b/projects/rocrand/library/src/rng/config/lfsr113_config.hpp index 5a48c9138434..d07914f4e787 100644 --- a/projects/rocrand/library/src/rng/config/lfsr113_config.hpp +++ b/projects/rocrand/library/src/rng/config/lfsr113_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_LFSR113_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_LFSR113_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -77,4 +73,4 @@ struct generator_config_selector } // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_LFSR113_HPP_ +#endif // ROCRAND_RNG_CONFIG_LFSR113_CONFIG_HPP_ diff --git a/projects/rocrand/library/src/rng/config/mrg31k3p_config.hpp b/projects/rocrand/library/src/rng/config/mrg31k3p_config.hpp index 422c3ec596ff..ae6fb2fabe2d 100644 --- a/projects/rocrand/library/src/rng/config/mrg31k3p_config.hpp +++ b/projects/rocrand/library/src/rng/config/mrg31k3p_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_MRG31K3P_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_MRG31K3P_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -77,4 +73,4 @@ struct generator_config_selector } // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_MRG31K3P_HPP_ +#endif // ROCRAND_RNG_CONFIG_MRG31K3P_CONFIG_HPP_ diff --git a/projects/rocrand/library/src/rng/config/mrg32k3a_config.hpp b/projects/rocrand/library/src/rng/config/mrg32k3a_config.hpp index d5cbb009c3b7..dd6fdd1fa918 100644 --- a/projects/rocrand/library/src/rng/config/mrg32k3a_config.hpp +++ b/projects/rocrand/library/src/rng/config/mrg32k3a_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_MRG32K3A_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_MRG32K3A_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -77,4 +73,4 @@ struct generator_config_selector } // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_MRG32K3A_HPP_ +#endif // ROCRAND_RNG_CONFIG_MRG32K3A_CONFIG_HPP_ diff --git a/projects/rocrand/library/src/rng/config/mt19937_config.hpp b/projects/rocrand/library/src/rng/config/mt19937_config.hpp index 6e8e35214a5d..7a1c069c6cd1 100644 --- a/projects/rocrand/library/src/rng/config/mt19937_config.hpp +++ b/projects/rocrand/library/src/rng/config/mt19937_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_MT19937_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_MT19937_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -71,6 +67,6 @@ struct generator_config_selector } }; -} // namespace rocrand_impl::host +} // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_MT19937_HPP_ +#endif // ROCRAND_RNG_CONFIG_MT19937_CONFIG_HPP_ diff --git a/projects/rocrand/library/src/rng/config/mtgp32_config.hpp b/projects/rocrand/library/src/rng/config/mtgp32_config.hpp index 1350a0f333e4..166c3eb9fc4b 100644 --- a/projects/rocrand/library/src/rng/config/mtgp32_config.hpp +++ b/projects/rocrand/library/src/rng/config/mtgp32_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_MTGP32_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_MTGP32_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -77,4 +73,4 @@ struct generator_config_selector } // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_MTGP32_HPP_ +#endif // ROCRAND_RNG_CONFIG_MTGP32_CONFIG_HPP_ diff --git a/projects/rocrand/library/src/rng/config/philox4_32_10_config.hpp b/projects/rocrand/library/src/rng/config/philox4_32_10_config.hpp index 83e0ddc18d67..044279c4b788 100644 --- a/projects/rocrand/library/src/rng/config/philox4_32_10_config.hpp +++ b/projects/rocrand/library/src/rng/config/philox4_32_10_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_PHILOX4_32_10_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_PHILOX4_32_10_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -77,4 +73,4 @@ struct generator_config_selector } // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_PHILOX4_32_10_HPP_ +#endif // ROCRAND_RNG_CONFIG_PHILOX4_32_10_CONFIG_HPP_ diff --git a/projects/rocrand/library/src/rng/config/threefry2_32_20_config.hpp b/projects/rocrand/library/src/rng/config/threefry2_32_20_config.hpp index d408a1e3dd35..42c9381a7e70 100644 --- a/projects/rocrand/library/src/rng/config/threefry2_32_20_config.hpp +++ b/projects/rocrand/library/src/rng/config/threefry2_32_20_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_THREEFRY2_32_20_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_THREEFRY2_32_20_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -79,4 +75,4 @@ struct generator_config_selector } // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_THREEFRY2_32_20_HPP_ +#endif // ROCRAND_RNG_CONFIG_THREEFRY2_32_20_CONFIG_HPP_ diff --git a/projects/rocrand/library/src/rng/config/threefry2_64_20_config.hpp b/projects/rocrand/library/src/rng/config/threefry2_64_20_config.hpp index be39bbd63e0c..07c4cebe4730 100644 --- a/projects/rocrand/library/src/rng/config/threefry2_64_20_config.hpp +++ b/projects/rocrand/library/src/rng/config/threefry2_64_20_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_THREEFRY2_64_20_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_THREEFRY2_64_20_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -79,4 +75,4 @@ struct generator_config_selector } // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_THREEFRY2_64_20_HPP_ +#endif // ROCRAND_RNG_CONFIG_THREEFRY2_64_20_CONFIG_HPP_ diff --git a/projects/rocrand/library/src/rng/config/threefry4_32_20_config.hpp b/projects/rocrand/library/src/rng/config/threefry4_32_20_config.hpp index e5e75904bee4..665f9b360f09 100644 --- a/projects/rocrand/library/src/rng/config/threefry4_32_20_config.hpp +++ b/projects/rocrand/library/src/rng/config/threefry4_32_20_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_THREEFRY4_32_20_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_THREEFRY4_32_20_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -79,4 +75,4 @@ struct generator_config_selector } // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_THREEFRY4_32_20_HPP_ +#endif // ROCRAND_RNG_CONFIG_THREEFRY4_32_20_CONFIG_HPP_ diff --git a/projects/rocrand/library/src/rng/config/threefry4_64_20_config.hpp b/projects/rocrand/library/src/rng/config/threefry4_64_20_config.hpp index f8eedce8c87c..53cec06e7fe8 100644 --- a/projects/rocrand/library/src/rng/config/threefry4_64_20_config.hpp +++ b/projects/rocrand/library/src/rng/config/threefry4_64_20_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_THREEFRY4_64_20_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_THREEFRY4_64_20_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -79,4 +75,4 @@ struct generator_config_selector } // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_THREEFRY4_64_20_HPP_ +#endif // ROCRAND_RNG_CONFIG_THREEFRY4_64_20_CONFIG_HPP_ diff --git a/projects/rocrand/library/src/rng/config/xorwow_config.hpp b/projects/rocrand/library/src/rng/config/xorwow_config.hpp index 73a90ecdfc67..e6448d24d48e 100644 --- a/projects/rocrand/library/src/rng/config/xorwow_config.hpp +++ b/projects/rocrand/library/src/rng/config/xorwow_config.hpp @@ -1,4 +1,4 @@ -// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. // // Permission is hereby granted, free of charge, to any person obtaining a copy // of this software and associated documentation files (the "Software"), to deal @@ -20,17 +20,13 @@ #ifndef ROCRAND_RNG_CONFIG_XORWOW_CONFIG_HPP_ #define ROCRAND_RNG_CONFIG_XORWOW_CONFIG_HPP_ - #include "config_defaults.hpp" #include - /* DO NOT EDIT THIS FILE - * This file is automatically generated by `/scripts/config-tuning/select_best_config.py`. + * This file is automatically generated by `/tuning/confgen/generate.py`. */ - namespace rocrand_impl::host { - template struct generator_config_selector { @@ -77,4 +73,4 @@ struct generator_config_selector } // end namespace rocrand_impl::host -#endif // ROCRAND_RNG_CONFIG_XORWOW_HPP_ +#endif // ROCRAND_RNG_CONFIG_XORWOW_CONFIG_HPP_ diff --git a/projects/rocrand/tuning/.gitignore b/projects/rocrand/tuning/.gitignore new file mode 100644 index 000000000000..966457cdffce --- /dev/null +++ b/projects/rocrand/tuning/.gitignore @@ -0,0 +1,5 @@ +output +__pycache__ +temp_*.cc +results.json +output \ No newline at end of file diff --git a/projects/rocrand/tuning/README.md b/projects/rocrand/tuning/README.md new file mode 100644 index 000000000000..3b9b1b116c2b --- /dev/null +++ b/projects/rocrand/tuning/README.md @@ -0,0 +1,133 @@ +# rocRAND Tuning + +This document is for rocRAND developers and contributors. + +The pseudo-random number generators (RNG engines) in rocRAND use tuned launch +configurations (block size and grid size) to improve performance on different GPU +architectures. These are located in `library/src/rng/config/_config.hpp`. +Manually finding optimal configurations is impractical, so the set of tools in this +folder is provided to automate the process. The tooling is built on top of +[Kernel Tuner](https://kerneltuner.github.io/kernel_tuner/stable/contents.html). + +## Supported engines + +The following engines have tuners (`tuner/tuning_.py`): + +* `lfsr113` +* `mrg31k3p` +* `mrg32k3a` +* `mt19937` +* `mtgp32` +* `philox4_32_10` +* `threefry2_32_20` +* `threefry2_64_20` +* `threefry4_32_20` +* `threefry4_64_20` +* `xorwow` + +Run `./run_tuning.py --list` to print the list that is actually discovered. + +## Provided tools + +The `tuning` folder contains mostly tools and templates. + +### Tuner + +Provides tuning scripts to tune individual engines. A tool to run multiple tuners is +provided as `run_tuning.py`; the individual tuners live in the `tuner` folder. + +Tuners are implemented by deriving from `tuner/base_tuner.py`, which contains reasonable +defaults for Kernel Tuner. The base class already tunes `block_size_x` and `grid_size` +over the ranges defined in `base_tuner.py`, so most engines only need a thin subclass +that sets the engine name (see `tuner/tuning_xorwow.py` for the simplest example). An +engine that needs extra constraints overrides `_get_restrictions(...)` to reject invalid +parameter combinations (see `tuner/tuning_mt19937.py` and `tuner/tuning_mtgp32.py`). + +Each engine also needs a benchmark wrapper template in `tuner/templates/.cpp.jinja2`, +which generates the C++ code that Kernel Tuner compiles and benchmarks. + +### Confgen + +Provides scripts and utilities to generate configuration headers. +`confgen/generate.py` reads the tuner output (JSON files, by default under `./output`) +and generates a config header per engine. It can optionally merge with the existing +config headers in `library/src/rng/config/` to keep previously tuned configurations for +other architectures. + +The configuration header is generated from templates. Most of the shared template +content is defined in `confgen/templates/common/base.h.jinja2`. This base template is +then inherited by `confgen/templates/.h.jinja2`, which provides the +engine-specific details. + +## Dependencies + +The tooling depends on the following Python libraries: + +* `numpy` +* `jinja2` +* `kernel-tuner` +* `hip-python` + +## Example: tuning xorwow + +As an example, the following commands can be used to tune the `xorwow` engine: + +```sh +# Change directory +cd projects/rocrand/tuning + +# Activate virtual environment +python -m venv .venv +source .venv/bin/activate + +# Install dependencies +pip3 install -i https://test.pypi.org/simple hip-python +pip3 install jinja2 kernel-tuner numpy + +# Run tuning (--algo-regex accepts a regex matched against engine names) +./run_tuning.py --help +./run_tuning.py --algo-regex xorwow + +# Generate configurations, merging with the existing headers. +# --target-arch/-t tells confgen which GPU architecture these results are for; +# this information is not present in the tuner output and must be supplied explicitly. +./confgen/generate.py --help +./confgen/generate.py \ + --input "./output/*.json" \ + --existing ../library/src/rng/config/ \ + --output ./tmp \ + --target-arch gfx90a +``` + +The generated headers are written to the `--output` directory (`./tmp` above). Once you +are satisfied with the results, copy them over the headers in +`library/src/rng/config/`. + +## Simulation mode and strategy comparison + +The tuner can run in simulation mode (`--simulation-mode`, requires `--arch-name`) to +replay a previously collected full search space cache under a different search strategy, +without touching the GPU. `utils/compare_simulated_strategies.py` compares the results of +multiple simulated strategies against the best-known times and renders a heatmap, which +is useful for choosing a tuning strategy: + +```sh +python3 utils/compare_simulated_strategies.py --arch gfx90a --algo xorwow +``` + +## Adding a new engine + +To implement tuning and configuration generation for a new engine, the following files +need to be created: + +* `tuner/tuning_.py` + * A subclass of `BaseTuner` that returns the engine name from `_get_default_args(...)`. + * Override `_get_restrictions(...)` if the engine has constraints on valid + `block_size_x` / `grid_size` combinations. +* `tuner/templates/.cpp.jinja2` + * The template used to generate the benchmark that Kernel Tuner runs. +* `confgen/templates/.h.jinja2` + * The template that specifies how to generate the config header. + +After creating these files, `run_tuning.py` will automatically discover the new engine +and `confgen/generate.py` can pick up and merge its configurations. From 668e4924427860d5775ecc6e585181e437f9c4ec Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 12:52:55 -0400 Subject: [PATCH 18/38] fix(rocrand): pre-checkin changes --- projects/rocrand/tuning/confgen/generate.py | 13 +++++---- projects/rocrand/tuning/confgen/parse.py | 28 +++++++++--------- projects/rocrand/tuning/run_tuning.py | 9 ++++-- projects/rocrand/tuning/tuner/base_tuner.py | 29 +++++++++---------- .../tuner/templates/common/base.cpp.jinja2 | 11 ++++--- .../tuner/templates/mrg31k3p.cpp.jinja2 | 2 +- .../tuner/templates/mrg32k3a.cpp.jinja2 | 2 +- .../tuning/tuner/templates/mt19937.cpp.jinja2 | 2 +- .../tuner/templates/philox4_32_10.cpp.jinja2 | 2 +- .../templates/threefry2_32_20.cpp.jinja2 | 2 +- .../templates/threefry2_64_20.cpp.jinja2 | 2 +- .../templates/threefry4_32_20.cpp.jinja2 | 2 +- .../templates/threefry4_64_20.cpp.jinja2 | 2 +- .../rocrand/tuning/tuner/tuning_lfsr113.py | 2 +- .../rocrand/tuning/tuner/tuning_mrg31k3p.py | 2 +- .../rocrand/tuning/tuner/tuning_mrg32k3a.py | 2 +- .../rocrand/tuning/tuner/tuning_mt19937.py | 10 ++++--- .../rocrand/tuning/tuner/tuning_mtgp32.py | 6 ++-- .../tuning/tuner/tuning_philox4_32_10.py | 2 +- .../tuning/tuner/tuning_threefry2_32_20.py | 2 +- .../tuning/tuner/tuning_threefry2_64_20.py | 2 +- .../tuning/tuner/tuning_threefry4_32_20.py | 2 +- .../tuning/tuner/tuning_threefry4_64_20.py | 2 +- .../rocrand/tuning/tuner/tuning_xorwow.py | 2 +- projects/rocrand/tuning/utils/utils.py | 4 +-- 25 files changed, 76 insertions(+), 68 deletions(-) diff --git a/projects/rocrand/tuning/confgen/generate.py b/projects/rocrand/tuning/confgen/generate.py index 6b24c0a8aea9..a8f70e8f2d9f 100644 --- a/projects/rocrand/tuning/confgen/generate.py +++ b/projects/rocrand/tuning/confgen/generate.py @@ -16,7 +16,7 @@ def make_config( type_data: dict[str, ty.Any]): import_keys = ["arch_name", "tune_params_keys", "tune_params"] - out = {k : type_data[k] for k in import_keys if k in type_data} + out = {k: type_data[k] for k in import_keys if k in type_data} return out def main(): @@ -99,7 +99,7 @@ def main(): algs[alg][target] = {} config = make_config(data) - config['tune_params'] = min_configs + config["tune_params"] = min_configs # Add config entry algs[alg][target] = config @@ -110,7 +110,7 @@ def main(): loader=jinja2.FileSystemLoader(script_dir / "templates"), autoescape=jinja2.select_autoescape(), trim_blocks=True, - lstrip_blocks=True + lstrip_blocks=True, ) # Find existing configs @@ -131,7 +131,10 @@ def main(): with open(config_file, "r") as file: existing_data = config_parser.parse_lines(file.readlines()) - merged_targets = {target_hash : make_config(target) for target_hash, target in existing_data.items()} + merged_targets = { + target_hash : make_config(target) + for target_hash, target in existing_data.items() + } for target_hash, config in targets.items(): merged_targets[target_hash] = config @@ -154,7 +157,7 @@ def main(): with open(output_dir / f"{algo_name}_config.hpp", "w") as file: file.write(rendered) - file.write(f'\n') + file.write(f"\n") if __name__ == "__main__": logging.basicConfig(level=logging.DEBUG) diff --git a/projects/rocrand/tuning/confgen/parse.py b/projects/rocrand/tuning/confgen/parse.py index 88b2fe02b86d..ac92166bc927 100644 --- a/projects/rocrand/tuning/confgen/parse.py +++ b/projects/rocrand/tuning/confgen/parse.py @@ -4,33 +4,33 @@ def parse_lines(lines: list[str]): config = {} - INSIDE_BLOCK = False + INSIDE_BLOCK = False INSIDE_GRID = False for line in lines: - if 'ROCRAND_RNG_PSEUDO' in line and 'generator_config_selector' in line: - engine_name = line.split('<')[-1].split(',')[0].split('_')[-1].lower() + if "ROCRAND_RNG_PSEUDO" in line and "generator_config_selector" in line: + engine_name = line.split("<")[-1].split(",")[0].split("_")[-1].lower() - if 'get_threads' in line: + if "get_threads" in line: INSIDE_BLOCK = True - if 'get_blocks' in line: + if "get_blocks" in line: INSIDE_BLOCK = False INSIDE_GRID = True - if 'target_arch::' in line: - arch = (line.split()[1].split(':')[-2]) + if "target_arch::" in line: + arch = (line.split()[1].split(":")[-2]) if arch not in config: - config[arch] = {"arch_name" : arch, "engine" : engine_name} + config[arch] = {"arch_name": arch, "engine": engine_name} if INSIDE_BLOCK: - block_size = int(line.split()[-1][:-1]) + block_size = int(line.split()[-1][:-1]) - config[arch]['tune_param_keys'] = ['block_size_x'] - config[arch]['tune_params'] = {"block_size_x" : block_size} + config[arch]["tune_param_keys"] = ["block_size_x"] + config[arch]["tune_params"] = {"block_size_x": block_size} elif INSIDE_GRID: - grid_size = int(line.split()[-1][:-1]) + grid_size = int(line.split()[-1][:-1]) - config[arch]['tune_param_keys'].append('grid_size') - config[arch]['tune_params']["grid_size"] = grid_size + config[arch]["tune_param_keys"].append("grid_size") + config[arch]["tune_params"]["grid_size"] = grid_size return config diff --git a/projects/rocrand/tuning/run_tuning.py b/projects/rocrand/tuning/run_tuning.py index b45a914287c7..d77b736fcb74 100644 --- a/projects/rocrand/tuning/run_tuning.py +++ b/projects/rocrand/tuning/run_tuning.py @@ -63,7 +63,12 @@ def import_module_from_file(file_path: str, module_name: str) -> Optional[Module def get_available_algorithms() -> List[str]: """Return list of supported tuning algorithms.""" - return sorted([re.sub(r'^tuning_|\.py', '', os.path.basename(f)) for f in glob(f'{BASE_DIR}/tuner/tuning_*')]) + return sorted( + [ + re.sub(r"^tuning_|\.py", "", os.path.basename(f)) + for f in glob(f"{BASE_DIR}/tuner/tuning_*") + ] + ) def filter_algorithms(available_algos: List[str], pattern: str) -> List[str]: """Filter algorithms based on regex pattern.""" @@ -196,4 +201,4 @@ def main(): if __name__ == "__main__": - main() \ No newline at end of file + main() diff --git a/projects/rocrand/tuning/tuner/base_tuner.py b/projects/rocrand/tuning/tuner/base_tuner.py index 00ceaff8882a..2fbc136efd4a 100644 --- a/projects/rocrand/tuning/tuner/base_tuner.py +++ b/projects/rocrand/tuning/tuner/base_tuner.py @@ -188,8 +188,8 @@ def cli(cls): def _get_tune_params(self) -> OrderedDict: params = OrderedDict() - params['block_size_x'] = BLOCK_SIZES - params['grid_size'] = GRID_SIZES + params["block_size_x"] = BLOCK_SIZES + params["grid_size"] = GRID_SIZES return params def _get_restrictions(self) -> Callable[[dict], bool] | List[str]: @@ -208,8 +208,8 @@ def _get_restrictions(self) -> Callable[[dict], bool] | List[str]: """ min_total_threads = 32768 def validate(params): - threads = params['block_size_x'] - blocks = params['grid_size'] + threads = params["block_size_x"] + blocks = params["grid_size"] return threads * blocks >= min_total_threads @@ -221,7 +221,7 @@ def _get_grid_div_x(self): def tune(self) -> None: """ - Run tuning + Run tuning """ print(f"\nTuning {self.algo_name}") @@ -248,8 +248,7 @@ def tune(self) -> None: print(f"Error: {str(e)}") raise - def _run_default_config( - self, tune_kernel_args: Dict) -> None: + def _run_default_config(self, tune_kernel_args: Dict) -> None: """Runs the default configuration if enabled.""" if self.exclude_default_config or self.simulation_mode: return @@ -263,7 +262,9 @@ def _run_default_config( return targets = [ - self.existing_config[k]['arch_name'] for k in self.existing_config if self.existing_config[k]['arch_name'] == target_arch + self.existing_config[k]["arch_name"] + for k in self.existing_config + if self.existing_config[k]["arch_name"] == target_arch ] target = None if len(targets) > 0: @@ -281,15 +282,13 @@ def _run_default_config( arch_config = self.existing_config[target] # Get first matching config - config = arch_config['tune_params'] + config = arch_config["tune_params"] if config is None: - warnings.warn( - f"No existing configuration found for {self.algo_name}'" - ) + warnings.warn(f"No existing configuration found for {self.algo_name}'") return - default_tune_params = { k: [v] for k, v in config.items()} + default_tune_params = {k: [v] for k, v in config.items()} # Get the base tuning archs and force set the range of the tune parameters # to the single-element lists 'default_tune_params'. We also change the @@ -370,7 +369,7 @@ def _get_base_tune_kernel_args( def _get_cache_file_name(self): """Return the name of the cache file based on algo name, arch name and key value types""" - cache_file_path = f'{self.algo_name}_{self.arch_name}_cache.json' + cache_file_path = f"{self.algo_name}_{self.arch_name}_cache.json" return cache_file_path @@ -424,4 +423,4 @@ def _get_compiler_options(self) -> List[str]: f"--offload-arch={self.arch_name}", "-lamd_smi", "-lrocrand", - ] \ No newline at end of file + ] diff --git a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 index 100e9485d43c..2f1451189f4d 100644 --- a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 @@ -4,8 +4,8 @@ // Include primbench executor. #include "primbench.hpp" -{% block backend_include %} -// Include the right backend +{% block backend_include %} +// Include the right backend #include "rng/{{algo_name}}.hpp" {% endblock %} #include @@ -35,7 +35,7 @@ struct output_type_supported : std::true_type DISABLE_ULL(generator_template); #undef DISABLE_ULL -// QUEUING TEMPLATE +// QUEUING TEMPLATE template class GeneratorTemplate, class ConfigProvider> void queue_one(primbench::executor& executor, size_t dimensions, @@ -209,10 +209,9 @@ extern "C" float {{algo_name}}_wrapper( double total = 0.0; for(double x : bps) - total += (1000.0 * primbench::GiB) / x; - + total += (1000.0 * primbench::GiB) / x; return total / static_cast(bps.size()); {% endblock %} } -{% endblock %} \ No newline at end of file +{% endblock %} diff --git a/projects/rocrand/tuning/tuner/templates/mrg31k3p.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/mrg31k3p.cpp.jinja2 index 0343bfde67b4..27cdd173aba8 100644 --- a/projects/rocrand/tuning/tuner/templates/mrg31k3p.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/mrg31k3p.cpp.jinja2 @@ -1,6 +1,6 @@ {% extends "common/base.cpp.jinja2" %} -{% block backend_include %} +{% block backend_include %} #include "rng/mrg.hpp" {% endblock %} diff --git a/projects/rocrand/tuning/tuner/templates/mrg32k3a.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/mrg32k3a.cpp.jinja2 index b727382436a0..2034b485bb6b 100644 --- a/projects/rocrand/tuning/tuner/templates/mrg32k3a.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/mrg32k3a.cpp.jinja2 @@ -1,6 +1,6 @@ {% extends "common/base.cpp.jinja2" %} -{% block backend_include %} +{% block backend_include %} #include "rng/mrg.hpp" {% endblock %} diff --git a/projects/rocrand/tuning/tuner/templates/mt19937.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/mt19937.cpp.jinja2 index 19e606358f23..863e149ff8e8 100644 --- a/projects/rocrand/tuning/tuner/templates/mt19937.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/mt19937.cpp.jinja2 @@ -1,6 +1,6 @@ {% extends "common/base.cpp.jinja2" %} -{% block backend_include %} +{% block backend_include %} #include "rng/mt19937.hpp" #include "rocrand/rocrand_mt19937_precomputed.h" {% endblock %} diff --git a/projects/rocrand/tuning/tuner/templates/philox4_32_10.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/philox4_32_10.cpp.jinja2 index b245027857a6..2bd3785ca143 100644 --- a/projects/rocrand/tuning/tuner/templates/philox4_32_10.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/philox4_32_10.cpp.jinja2 @@ -1,6 +1,6 @@ {% extends "common/base.cpp.jinja2" %} -{% block backend_include %} +{% block backend_include %} #include "rng/philox4x32_10.hpp" {% endblock %} diff --git a/projects/rocrand/tuning/tuner/templates/threefry2_32_20.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/threefry2_32_20.cpp.jinja2 index 9cc4cbfebe24..c29e2feb33db 100644 --- a/projects/rocrand/tuning/tuner/templates/threefry2_32_20.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/threefry2_32_20.cpp.jinja2 @@ -1,6 +1,6 @@ {% extends "common/base.cpp.jinja2" %} -{% block backend_include %} +{% block backend_include %} #include "rng/threefry.hpp" {% endblock %} diff --git a/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 index fa561c571b1f..544075783589 100644 --- a/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 @@ -1,6 +1,6 @@ {% extends "common/base.cpp.jinja2" %} -{% block backend_include %} +{% block backend_include %} #include "rng/threefry.hpp" {% endblock %} diff --git a/projects/rocrand/tuning/tuner/templates/threefry4_32_20.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/threefry4_32_20.cpp.jinja2 index e88b599ba1c0..df9768ea3868 100644 --- a/projects/rocrand/tuning/tuner/templates/threefry4_32_20.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/threefry4_32_20.cpp.jinja2 @@ -1,6 +1,6 @@ {% extends "common/base.cpp.jinja2" %} -{% block backend_include %} +{% block backend_include %} #include "rng/threefry.hpp" {% endblock %} diff --git a/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 index 010c548181ac..ff80fd25fa05 100644 --- a/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 @@ -1,6 +1,6 @@ {% extends "common/base.cpp.jinja2" %} -{% block backend_include %} +{% block backend_include %} #include "rng/threefry.hpp" {% endblock %} diff --git a/projects/rocrand/tuning/tuner/tuning_lfsr113.py b/projects/rocrand/tuning/tuner/tuning_lfsr113.py index 9c02d89ae197..1144eb8d13b4 100644 --- a/projects/rocrand/tuning/tuner/tuning_lfsr113.py +++ b/projects/rocrand/tuning/tuner/tuning_lfsr113.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='lfsr113') + return TunerArgs(algo_full_name="lfsr113") def __init__(self, args: TunerArgs) -> None: super().__init__(args) diff --git a/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py b/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py index 836fb4acbaa0..d6ba6df68431 100644 --- a/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py +++ b/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='mrg31k3p') + return TunerArgs(algo_full_name="mrg31k3p") def __init__(self, args: TunerArgs) -> None: super().__init__(args) diff --git a/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py b/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py index de36e33d6923..9859e0dd49d4 100644 --- a/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py +++ b/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='mrg32k3a') + return TunerArgs(algo_full_name="mrg32k3a") def __init__(self, args: TunerArgs) -> None: super().__init__(args) diff --git a/projects/rocrand/tuning/tuner/tuning_mt19937.py b/projects/rocrand/tuning/tuner/tuning_mt19937.py index 353ce80b2b5e..cff351e0962c 100644 --- a/projects/rocrand/tuning/tuner/tuning_mt19937.py +++ b/projects/rocrand/tuning/tuner/tuning_mt19937.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='mt19937') + return TunerArgs(algo_full_name="mt19937") def __init__(self, args: TunerArgs) -> None: super().__init__(args) @@ -40,13 +40,15 @@ def _get_restrictions(self): mt19937_jumps_radix = 256 threads_per_generator = 8 def validate(params): - threads = params['block_size_x'] - blocks = params['grid_size'] + threads = params["block_size_x"] + blocks = params["grid_size"] if threads * blocks < min_total_threads: return False - return (threads * blocks) // threads_per_generator <= (mt19937_jumps_radix * mt19937_jumps_radix) + return (threads * blocks) // threads_per_generator <= ( + mt19937_jumps_radix * mt19937_jumps_radix + ) return validate diff --git a/projects/rocrand/tuning/tuner/tuning_mtgp32.py b/projects/rocrand/tuning/tuner/tuning_mtgp32.py index 259f8fb75926..06b5adeda7c0 100644 --- a/projects/rocrand/tuning/tuner/tuning_mtgp32.py +++ b/projects/rocrand/tuning/tuner/tuning_mtgp32.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='mtgp32') + return TunerArgs(algo_full_name="mtgp32") def __init__(self, args: TunerArgs) -> None: super().__init__(args) @@ -38,8 +38,8 @@ def __init__(self, args: TunerArgs) -> None: def _get_restrictions(self): min_total_threads = 32768 def validate(params): - threads = params['block_size_x'] - blocks = params['grid_size'] + threads = params["block_size_x"] + blocks = params["grid_size"] if threads * blocks < min_total_threads: return False diff --git a/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py b/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py index 1465c47bc486..2b0c9f43b9d4 100644 --- a/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py +++ b/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='philox4_32_10') + return TunerArgs(algo_full_name="philox4_32_10") def __init__(self, args: TunerArgs) -> None: super().__init__(args) diff --git a/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py b/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py index efd8b35b9693..18e3e683ec56 100644 --- a/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py +++ b/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='threefry2_32_20') + return TunerArgs(algo_full_name="threefry2_32_20") def __init__(self, args: TunerArgs) -> None: super().__init__(args) diff --git a/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py b/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py index d54a0a8b7a3b..11a9f577260f 100644 --- a/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py +++ b/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='threefry2_64_20') + return TunerArgs(algo_full_name="threefry2_64_20") def __init__(self, args: TunerArgs) -> None: super().__init__(args) diff --git a/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py b/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py index c460a21aeae8..e2b6999e66f7 100644 --- a/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py +++ b/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='threefry4_32_20') + return TunerArgs(algo_full_name="threefry4_32_20") def __init__(self, args: TunerArgs) -> None: super().__init__(args) diff --git a/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py b/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py index 07bf31f892bd..94d668a2bac1 100644 --- a/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py +++ b/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='threefry4_64_20') + return TunerArgs(algo_full_name="threefry4_64_20") def __init__(self, args: TunerArgs) -> None: super().__init__(args) diff --git a/projects/rocrand/tuning/tuner/tuning_xorwow.py b/projects/rocrand/tuning/tuner/tuning_xorwow.py index 31d564a255f9..f01067d6103c 100644 --- a/projects/rocrand/tuning/tuner/tuning_xorwow.py +++ b/projects/rocrand/tuning/tuner/tuning_xorwow.py @@ -30,7 +30,7 @@ class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: - return TunerArgs(algo_full_name='xorwow') + return TunerArgs(algo_full_name="xorwow") def __init__(self, args: TunerArgs) -> None: super().__init__(args) diff --git a/projects/rocrand/tuning/utils/utils.py b/projects/rocrand/tuning/utils/utils.py index 6bb0e7d3a5bc..15579d912dbd 100644 --- a/projects/rocrand/tuning/utils/utils.py +++ b/projects/rocrand/tuning/utils/utils.py @@ -22,9 +22,9 @@ import argparse from typing import Optional -import os +import os -BASE_DIR = f'{os.path.dirname(os.path.abspath(__file__))}/..' +BASE_DIR = f"{os.path.dirname(os.path.abspath(__file__))}/.." class Parser: @staticmethod From a12eab548b7b1ee2ef97386169f67cd0247b3d0e Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 13:05:09 -0400 Subject: [PATCH 19/38] fix(rocrand): more pre-checkin changes --- projects/rocrand/tuning/.gitignore | 2 +- projects/rocrand/tuning/confgen/generate.py | 5 ++--- projects/rocrand/tuning/confgen/parse.py | 4 ++-- projects/rocrand/tuning/tuner/base_tuner.py | 2 +- .../rocrand/tuning/tuner/templates/common/base.cpp.jinja2 | 2 +- 5 files changed, 7 insertions(+), 8 deletions(-) diff --git a/projects/rocrand/tuning/.gitignore b/projects/rocrand/tuning/.gitignore index 966457cdffce..af94260a83a4 100644 --- a/projects/rocrand/tuning/.gitignore +++ b/projects/rocrand/tuning/.gitignore @@ -2,4 +2,4 @@ output __pycache__ temp_*.cc results.json -output \ No newline at end of file +output diff --git a/projects/rocrand/tuning/confgen/generate.py b/projects/rocrand/tuning/confgen/generate.py index a8f70e8f2d9f..73b5422175c5 100644 --- a/projects/rocrand/tuning/confgen/generate.py +++ b/projects/rocrand/tuning/confgen/generate.py @@ -13,8 +13,7 @@ log = logging.getLogger("confgen.generate") -def make_config( - type_data: dict[str, ty.Any]): +def make_config(type_data: dict[str, ty.Any]): import_keys = ["arch_name", "tune_params_keys", "tune_params"] out = {k: type_data[k] for k in import_keys if k in type_data} return out @@ -133,7 +132,7 @@ def main(): merged_targets = { target_hash : make_config(target) - for target_hash, target in existing_data.items() + for target_hash, target in existing_data.items() } for target_hash, config in targets.items(): diff --git a/projects/rocrand/tuning/confgen/parse.py b/projects/rocrand/tuning/confgen/parse.py index ac92166bc927..a1a150ceb437 100644 --- a/projects/rocrand/tuning/confgen/parse.py +++ b/projects/rocrand/tuning/confgen/parse.py @@ -15,10 +15,10 @@ def parse_lines(lines: list[str]): INSIDE_BLOCK = True if "get_blocks" in line: INSIDE_BLOCK = False - INSIDE_GRID = True + INSIDE_GRID = True if "target_arch::" in line: - arch = (line.split()[1].split(":")[-2]) + arch = line.split()[1].split(":")[-2] if arch not in config: config[arch] = {"arch_name": arch, "engine": engine_name} if INSIDE_BLOCK: diff --git a/projects/rocrand/tuning/tuner/base_tuner.py b/projects/rocrand/tuning/tuner/base_tuner.py index 2fbc136efd4a..6b392313aa10 100644 --- a/projects/rocrand/tuning/tuner/base_tuner.py +++ b/projects/rocrand/tuning/tuner/base_tuner.py @@ -263,7 +263,7 @@ def _run_default_config(self, tune_kernel_args: Dict) -> None: targets = [ self.existing_config[k]["arch_name"] - for k in self.existing_config + for k in self.existing_config if self.existing_config[k]["arch_name"] == target_arch ] target = None diff --git a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 index 2f1451189f4d..867ffd42725a 100644 --- a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 @@ -209,7 +209,7 @@ extern "C" float {{algo_name}}_wrapper( double total = 0.0; for(double x : bps) - total += (1000.0 * primbench::GiB) / x; + total += (1000.0 * primbench::GiB) / x; return total / static_cast(bps.size()); {% endblock %} From 9d7bf59ed1f5b498434150b891ebcaee1a2ec2f6 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 13:08:24 -0400 Subject: [PATCH 20/38] fix(rocrand): even more pre-checkin changes --- projects/rocrand/tuning/confgen/generate.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rocrand/tuning/confgen/generate.py b/projects/rocrand/tuning/confgen/generate.py index 73b5422175c5..5315298c628d 100644 --- a/projects/rocrand/tuning/confgen/generate.py +++ b/projects/rocrand/tuning/confgen/generate.py @@ -131,7 +131,7 @@ def main(): existing_data = config_parser.parse_lines(file.readlines()) merged_targets = { - target_hash : make_config(target) + target_hash: make_config(target) for target_hash, target in existing_data.items() } From e5fffa620c8a8a10928245a8c744578321d7ceb9 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 16:45:11 -0400 Subject: [PATCH 21/38] feat(rocrand): updated grid and block size defaults --- projects/rocrand/tuning/tuner/base_tuner.py | 31 ++++++++++++++++++--- 1 file changed, 27 insertions(+), 4 deletions(-) diff --git a/projects/rocrand/tuning/tuner/base_tuner.py b/projects/rocrand/tuning/tuner/base_tuner.py index 6b392313aa10..4b7c952932f5 100644 --- a/projects/rocrand/tuning/tuner/base_tuner.py +++ b/projects/rocrand/tuning/tuner/base_tuner.py @@ -34,6 +34,8 @@ from dataclasses import dataclass import confgen.parse import pathlib +import re +import subprocess """ The following base class is used when implementing the tuning for new algorithms @@ -87,8 +89,7 @@ """ Inclusive range for params tuning, edit these to adjust tuning grid range. """ -BLOCK_SIZES = [32, 64, 128, 256, 512, 1024, 2048] -GRID_SIZES = [32, 64, 128, 256, 512, 1024] +BLOCK_SIZES = [32, 64, 128, 256, 512, 1024] @dataclass class TunerArgs: @@ -185,11 +186,33 @@ def cli(cls): ) defaults.update_with_kwargs(**args) cls(defaults).tune() + + def _get_grid_sizes(self) -> str: + rocminfo_out = subprocess.check_output('rocminfo', encoding='utf-8') + match = re.search(r'^\s*Name:\s*gfx\d+.*?^\s*Compute Unit:\s*(\d+)', + rocminfo_out, flags=re.MULTILINE | re.DOTALL) + if not match: + raise Exception('Could not find Compute Unit info in rocminfo output') + num_compute_units = int(match.group(1)) + compute_unit_multipliers = [4, 5, 8, 10, 16, 32] + min_grid_size = 128 + max_grid_size = 4096 + + grid_sizes = [128, 256, 512, 1024, 2048] + for cu_mul in compute_unit_multipliers: + new_grid_size = cu_mul * num_compute_units + if new_grid_size >= min_grid_size and new_grid_size <= max_grid_size: + grid_sizes.append(new_grid_size) + grid_sizes = list(set(grid_sizes)) # Unique + grid_sizes.sort() + return grid_sizes def _get_tune_params(self) -> OrderedDict: params = OrderedDict() - params["block_size_x"] = BLOCK_SIZES - params["grid_size"] = GRID_SIZES + + params["block_size_x"] = BLOCK_SIZES + params["grid_size"] = self._get_grid_sizes() + return params def _get_restrictions(self) -> Callable[[dict], bool] | List[str]: From 992311188e3c12b465b8e2de17e5b4840b1131cf Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 16:51:07 -0400 Subject: [PATCH 22/38] fix(rocrand): allow threefryx_64_20 to tune ull --- projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 | 3 ++- .../rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 | 3 +++ .../rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 | 3 +++ 3 files changed, 8 insertions(+), 1 deletion(-) diff --git a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 index 867ffd42725a..dd6676d2b985 100644 --- a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 @@ -21,7 +21,7 @@ using generator_template ConfigProvider>; {% endblock %} - +{% block disable_ull %} // Disable ULL template class GeneratorTemplate> struct output_type_supported : std::true_type @@ -34,6 +34,7 @@ struct output_type_supported : std::true_type DISABLE_ULL(generator_template); #undef DISABLE_ULL +{% endblock %} // QUEUING TEMPLATE template class GeneratorTemplate, class ConfigProvider> diff --git a/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 index 544075783589..d0330bd8da14 100644 --- a/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/threefry2_64_20.cpp.jinja2 @@ -12,3 +12,6 @@ rocrand_impl::host::threefry_device_engine, ConfigProvider>; {% endblock %} + +{% block disable_ull %} +{% endblock %} diff --git a/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 index ff80fd25fa05..9d7d6ff4511b 100644 --- a/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/threefry4_64_20.cpp.jinja2 @@ -12,3 +12,6 @@ rocrand_impl::host::threefry_device_engine, ConfigProvider>; {% endblock %} + +{% block disable_ull %} +{% endblock %} From 155f76fb403b6b1fa7c6478bffe9a35a24582ea4 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 1 Oct 2026 17:07:07 -0400 Subject: [PATCH 23/38] fix(rocrand): prevent arch result being overwritten --- projects/rocrand/tuning/confgen/generate.py | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/projects/rocrand/tuning/confgen/generate.py b/projects/rocrand/tuning/confgen/generate.py index 5315298c628d..38f9a8156432 100644 --- a/projects/rocrand/tuning/confgen/generate.py +++ b/projects/rocrand/tuning/confgen/generate.py @@ -38,6 +38,7 @@ def main(): "--target-arch", "-t", action="append", + required=True, help="The target arch that these configs are for.", ) @@ -91,8 +92,10 @@ def main(): alg: str = data["algo_name"] if alg not in algs: algs[alg] = {} - + arch = data["arch_name"] for target in gfx_target: + if target != arch: + continue # Ensure target entry exists if target not in algs[alg]: algs[alg][target] = {} From f086ec9890d999f44d3b445d2a56e10cb7981bdd Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Fri, 2 Oct 2026 15:22:44 -0400 Subject: [PATCH 24/38] fix(rocrand): fix help text typo --- projects/rocrand/tuning/utils/utils.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rocrand/tuning/utils/utils.py b/projects/rocrand/tuning/utils/utils.py index 15579d912dbd..2f8b63be3d50 100644 --- a/projects/rocrand/tuning/utils/utils.py +++ b/projects/rocrand/tuning/utils/utils.py @@ -58,7 +58,7 @@ def get_parser( "--exclude-default-config", action="store_true", default=False, - help="Include default configs of previous tuning in the new results (default: off)", + help="Exclude default configs of previous tuning in the new results (default: off)", ) parser.add_argument( "--strategy", From c4b04351b9b534aa9838ef84a27a0468e45ecc1e Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Fri, 2 Oct 2026 15:25:40 -0400 Subject: [PATCH 25/38] fix(rocrand): made error message 1 instead of 1 --- projects/rocrand/tuning/confgen/generate.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rocrand/tuning/confgen/generate.py b/projects/rocrand/tuning/confgen/generate.py index 38f9a8156432..e663ae260d89 100644 --- a/projects/rocrand/tuning/confgen/generate.py +++ b/projects/rocrand/tuning/confgen/generate.py @@ -121,7 +121,7 @@ def main(): existing_dir = pathlib.Path(args.existing) if not existing_dir.exists(): log.error(f"Could not find directory: {existing_dir}") - quit() + quit(1) # Merge newly generated configs to existing configs if existing_dir: From ed54be506134ce5245a9b03a1980d5d3cdfdf907 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Fri, 2 Oct 2026 15:27:43 -0400 Subject: [PATCH 26/38] fix(rocrand): removed dead cli flag --- projects/rocrand/tuning/utils/utils.py | 6 ------ 1 file changed, 6 deletions(-) diff --git a/projects/rocrand/tuning/utils/utils.py b/projects/rocrand/tuning/utils/utils.py index 2f8b63be3d50..e9e8796bdd61 100644 --- a/projects/rocrand/tuning/utils/utils.py +++ b/projects/rocrand/tuning/utils/utils.py @@ -77,12 +77,6 @@ def get_parser( required=False, help="Specify arch, needed when running in simulation", ) - parser.add_argument( - "--save-metadata", - action="store_true", - default=False, - help="Save Kernel Tuner metadata (default: off)", - ) parser.add_argument( "--seed", type=int, From 25964fadfd0eb3d6f73398c2c6ee7cf97117b064 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Mon, 5 Oct 2026 17:52:47 -0400 Subject: [PATCH 27/38] fix(rocrand): ran pre-checkin script --- projects/rocrand/tuning/confgen/generate.py | 4 +++- projects/rocrand/tuning/confgen/parse.py | 1 + projects/rocrand/tuning/run_tuning.py | 2 ++ projects/rocrand/tuning/tuner/base_tuner.py | 17 +++++++++++------ .../tuner/templates/common/base.cpp.jinja2 | 6 +++--- projects/rocrand/tuning/tuner/tuning_lfsr113.py | 2 ++ .../rocrand/tuning/tuner/tuning_mrg31k3p.py | 3 ++- .../rocrand/tuning/tuner/tuning_mrg32k3a.py | 3 ++- projects/rocrand/tuning/tuner/tuning_mt19937.py | 4 +++- projects/rocrand/tuning/tuner/tuning_mtgp32.py | 5 +++-- .../tuning/tuner/tuning_philox4_32_10.py | 3 ++- .../tuning/tuner/tuning_threefry2_32_20.py | 3 ++- .../tuning/tuner/tuning_threefry2_64_20.py | 3 ++- .../tuning/tuner/tuning_threefry4_32_20.py | 3 ++- .../tuning/tuner/tuning_threefry4_64_20.py | 3 ++- projects/rocrand/tuning/tuner/tuning_xorwow.py | 2 ++ projects/rocrand/tuning/utils/utils.py | 1 + 17 files changed, 45 insertions(+), 20 deletions(-) diff --git a/projects/rocrand/tuning/confgen/generate.py b/projects/rocrand/tuning/confgen/generate.py index e663ae260d89..cff9aa380bd9 100644 --- a/projects/rocrand/tuning/confgen/generate.py +++ b/projects/rocrand/tuning/confgen/generate.py @@ -13,11 +13,13 @@ log = logging.getLogger("confgen.generate") + def make_config(type_data: dict[str, ty.Any]): import_keys = ["arch_name", "tune_params_keys", "tune_params"] out = {k: type_data[k] for k in import_keys if k in type_data} return out + def main(): cli = argparse.ArgumentParser() cli.add_argument( @@ -161,7 +163,7 @@ def main(): file.write(rendered) file.write(f"\n") + if __name__ == "__main__": logging.basicConfig(level=logging.DEBUG) main() - \ No newline at end of file diff --git a/projects/rocrand/tuning/confgen/parse.py b/projects/rocrand/tuning/confgen/parse.py index a1a150ceb437..6837b0325611 100644 --- a/projects/rocrand/tuning/confgen/parse.py +++ b/projects/rocrand/tuning/confgen/parse.py @@ -1,5 +1,6 @@ #!/usr/bin/env python3 + def parse_lines(lines: list[str]): config = {} diff --git a/projects/rocrand/tuning/run_tuning.py b/projects/rocrand/tuning/run_tuning.py index d77b736fcb74..2200184084a1 100644 --- a/projects/rocrand/tuning/run_tuning.py +++ b/projects/rocrand/tuning/run_tuning.py @@ -34,6 +34,7 @@ from tuner.base_tuner import TunerArgs + @contextmanager def working_directory(path: Path): """Context manager for changing the current working directory.""" @@ -70,6 +71,7 @@ def get_available_algorithms() -> List[str]: ] ) + def filter_algorithms(available_algos: List[str], pattern: str) -> List[str]: """Filter algorithms based on regex pattern.""" try: diff --git a/projects/rocrand/tuning/tuner/base_tuner.py b/projects/rocrand/tuning/tuner/base_tuner.py index 4b7c952932f5..8d86ff300088 100644 --- a/projects/rocrand/tuning/tuner/base_tuner.py +++ b/projects/rocrand/tuning/tuner/base_tuner.py @@ -91,6 +91,7 @@ """ BLOCK_SIZES = [32, 64, 128, 256, 512, 1024] + @dataclass class TunerArgs: algo_full_name: Optional[str] = None @@ -188,11 +189,14 @@ def cli(cls): cls(defaults).tune() def _get_grid_sizes(self) -> str: - rocminfo_out = subprocess.check_output('rocminfo', encoding='utf-8') - match = re.search(r'^\s*Name:\s*gfx\d+.*?^\s*Compute Unit:\s*(\d+)', - rocminfo_out, flags=re.MULTILINE | re.DOTALL) + rocminfo_out = subprocess.check_output("rocminfo", encoding="utf-8") + match = re.search( + r"^\s*Name:\s*gfx\d+.*?^\s*Compute Unit:\s*(\d+)", + rocminfo_out, + flags=re.MULTILINE | re.DOTALL, + ) if not match: - raise Exception('Could not find Compute Unit info in rocminfo output') + raise Exception("Could not find Compute Unit info in rocminfo output") num_compute_units = int(match.group(1)) compute_unit_multipliers = [4, 5, 8, 10, 16, 32] min_grid_size = 128 @@ -206,11 +210,11 @@ def _get_grid_sizes(self) -> str: grid_sizes = list(set(grid_sizes)) # Unique grid_sizes.sort() return grid_sizes - + def _get_tune_params(self) -> OrderedDict: params = OrderedDict() - params["block_size_x"] = BLOCK_SIZES + params["block_size_x"] = BLOCK_SIZES params["grid_size"] = self._get_grid_sizes() return params @@ -230,6 +234,7 @@ def _get_restrictions(self) -> Callable[[dict], bool] | List[str]: See Kernel Tuner's documentation for more details """ min_total_threads = 32768 + def validate(params): threads = params["block_size_x"] blocks = params["grid_size"] diff --git a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 index dd6676d2b985..1a50c4cc405e 100644 --- a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 @@ -171,7 +171,7 @@ extern "C" float {{algo_name}}_wrapper( { // Flush is required for fancy primbench printing std::cout << std::flush; - + {% block primbench_settings %} auto settings = primbench::settings(); settings.json_out = "/dev/null"; @@ -184,7 +184,7 @@ extern "C" float {{algo_name}}_wrapper( {% block primbench_args %} char* args[] = { "benchmark-executable" /* primbench requires an executable name for exposing --help */ }; {% endblock %} - + {% block primbench_executor %} auto executor = primbench::executor(1, args, settings, primbench::flags::sync); {% endblock %} @@ -207,7 +207,7 @@ extern "C" float {{algo_name}}_wrapper( {% block calculate_time %} std::vector bps = executor.get_all_bytes_per_second(); - + double total = 0.0; for(double x : bps) total += (1000.0 * primbench::GiB) / x; diff --git a/projects/rocrand/tuning/tuner/tuning_lfsr113.py b/projects/rocrand/tuning/tuner/tuning_lfsr113.py index 1144eb8d13b4..c1a6aca59748 100644 --- a/projects/rocrand/tuning/tuner/tuning_lfsr113.py +++ b/projects/rocrand/tuning/tuner/tuning_lfsr113.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -35,5 +36,6 @@ def _get_default_args(cls) -> TunerArgs: def __init__(self, args: TunerArgs) -> None: super().__init__(args) + if __name__ == "__main__": Tuner.cli() diff --git a/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py b/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py index d6ba6df68431..c8d22fac4542 100644 --- a/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py +++ b/projects/rocrand/tuning/tuner/tuning_mrg31k3p.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -35,6 +36,6 @@ def _get_default_args(cls) -> TunerArgs: def __init__(self, args: TunerArgs) -> None: super().__init__(args) + if __name__ == "__main__": Tuner.cli() - \ No newline at end of file diff --git a/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py b/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py index 9859e0dd49d4..9a252979f4e2 100644 --- a/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py +++ b/projects/rocrand/tuning/tuner/tuning_mrg32k3a.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -35,6 +36,6 @@ def _get_default_args(cls) -> TunerArgs: def __init__(self, args: TunerArgs) -> None: super().__init__(args) + if __name__ == "__main__": Tuner.cli() - \ No newline at end of file diff --git a/projects/rocrand/tuning/tuner/tuning_mt19937.py b/projects/rocrand/tuning/tuner/tuning_mt19937.py index cff351e0962c..9b62cfa63263 100644 --- a/projects/rocrand/tuning/tuner/tuning_mt19937.py +++ b/projects/rocrand/tuning/tuner/tuning_mt19937.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -39,6 +40,7 @@ def _get_restrictions(self): min_total_threads = 32768 mt19937_jumps_radix = 256 threads_per_generator = 8 + def validate(params): threads = params["block_size_x"] blocks = params["grid_size"] @@ -52,6 +54,6 @@ def validate(params): return validate + if __name__ == "__main__": Tuner.cli() - \ No newline at end of file diff --git a/projects/rocrand/tuning/tuner/tuning_mtgp32.py b/projects/rocrand/tuning/tuner/tuning_mtgp32.py index 06b5adeda7c0..602c4f233dc1 100644 --- a/projects/rocrand/tuning/tuner/tuning_mtgp32.py +++ b/projects/rocrand/tuning/tuner/tuning_mtgp32.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -37,18 +38,18 @@ def __init__(self, args: TunerArgs) -> None: def _get_restrictions(self): min_total_threads = 32768 + def validate(params): threads = params["block_size_x"] blocks = params["grid_size"] if threads * blocks < min_total_threads: return False - return blocks <= 512 and threads == 256 return validate + if __name__ == "__main__": Tuner.cli() - \ No newline at end of file diff --git a/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py b/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py index 2b0c9f43b9d4..08bac0b25b0b 100644 --- a/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py +++ b/projects/rocrand/tuning/tuner/tuning_philox4_32_10.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -35,6 +36,6 @@ def _get_default_args(cls) -> TunerArgs: def __init__(self, args: TunerArgs) -> None: super().__init__(args) + if __name__ == "__main__": Tuner.cli() - \ No newline at end of file diff --git a/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py b/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py index 18e3e683ec56..22b6ab7f0b25 100644 --- a/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py +++ b/projects/rocrand/tuning/tuner/tuning_threefry2_32_20.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -35,6 +36,6 @@ def _get_default_args(cls) -> TunerArgs: def __init__(self, args: TunerArgs) -> None: super().__init__(args) + if __name__ == "__main__": Tuner.cli() - \ No newline at end of file diff --git a/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py b/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py index 11a9f577260f..58e3ba0e7c6e 100644 --- a/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py +++ b/projects/rocrand/tuning/tuner/tuning_threefry2_64_20.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -35,6 +36,6 @@ def _get_default_args(cls) -> TunerArgs: def __init__(self, args: TunerArgs) -> None: super().__init__(args) + if __name__ == "__main__": Tuner.cli() - \ No newline at end of file diff --git a/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py b/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py index e2b6999e66f7..83673c5d0f8f 100644 --- a/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py +++ b/projects/rocrand/tuning/tuner/tuning_threefry4_32_20.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -35,6 +36,6 @@ def _get_default_args(cls) -> TunerArgs: def __init__(self, args: TunerArgs) -> None: super().__init__(args) + if __name__ == "__main__": Tuner.cli() - \ No newline at end of file diff --git a/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py b/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py index 94d668a2bac1..31e7355d195d 100644 --- a/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py +++ b/projects/rocrand/tuning/tuner/tuning_threefry4_64_20.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -35,6 +36,6 @@ def _get_default_args(cls) -> TunerArgs: def __init__(self, args: TunerArgs) -> None: super().__init__(args) + if __name__ == "__main__": Tuner.cli() - \ No newline at end of file diff --git a/projects/rocrand/tuning/tuner/tuning_xorwow.py b/projects/rocrand/tuning/tuner/tuning_xorwow.py index f01067d6103c..6219e0aaf1c2 100644 --- a/projects/rocrand/tuning/tuner/tuning_xorwow.py +++ b/projects/rocrand/tuning/tuner/tuning_xorwow.py @@ -27,6 +27,7 @@ from tuner.base_tuner import BaseTuner, TunerArgs + class Tuner(BaseTuner): @classmethod def _get_default_args(cls) -> TunerArgs: @@ -35,5 +36,6 @@ def _get_default_args(cls) -> TunerArgs: def __init__(self, args: TunerArgs) -> None: super().__init__(args) + if __name__ == "__main__": Tuner.cli() diff --git a/projects/rocrand/tuning/utils/utils.py b/projects/rocrand/tuning/utils/utils.py index e9e8796bdd61..fef98080771c 100644 --- a/projects/rocrand/tuning/utils/utils.py +++ b/projects/rocrand/tuning/utils/utils.py @@ -26,6 +26,7 @@ BASE_DIR = f"{os.path.dirname(os.path.abspath(__file__))}/.." + class Parser: @staticmethod def get_parser( From 2a288d5321fa31f28685bbe26a34b24751891677 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Mon, 5 Oct 2026 18:40:40 -0400 Subject: [PATCH 28/38] fix(rocrand): moved default output_type_supported def outside of code block --- .../rocrand/tuning/tuner/templates/common/base.cpp.jinja2 | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 index 1a50c4cc405e..79b1194a911b 100644 --- a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 @@ -21,12 +21,13 @@ using generator_template ConfigProvider>; {% endblock %} -{% block disable_ull %} -// Disable ULL template class GeneratorTemplate> struct output_type_supported : std::true_type {}; +{% block disable_ull %} +// Disable ULL + #define DISABLE_ULL(Generator) \ template<> \ struct output_type_supported : std::false_type \ From 4d84e5df4595d1cca8fe4636bc7fe3df811720e3 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Wed, 7 Oct 2026 12:03:17 -0400 Subject: [PATCH 29/38] fix(rocrand): pass self.seed to kernel tuner --- projects/rocrand/tuning/tuner/base_tuner.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/rocrand/tuning/tuner/base_tuner.py b/projects/rocrand/tuning/tuner/base_tuner.py index 8d86ff300088..9f2c414483b7 100644 --- a/projects/rocrand/tuning/tuner/base_tuner.py +++ b/projects/rocrand/tuning/tuner/base_tuner.py @@ -264,7 +264,7 @@ def tune(self) -> None: try: tune_kernel_args = self._get_base_tune_kernel_args() # Run main tuning - results, _ = kernel_tuner.tune_kernel(**tune_kernel_args) + results, _ = kernel_tuner.tune_kernel(**tune_kernel_args, seed=self.seed) # Run default config if enabled self._run_default_config(tune_kernel_args) @@ -331,7 +331,7 @@ def _run_default_config(self, tune_kernel_args: Dict) -> None: ) print(f"Found existing configuration: {default_tune_params}") - kernel_tuner.tune_kernel(**tune_kernel_args) + kernel_tuner.tune_kernel(**tune_kernel_args, seed=self.seed) def generate_wrapper( self, From a1874cd44a7288bd2d4a6dbec0ee49b310728976 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Wed, 7 Oct 2026 12:21:25 -0400 Subject: [PATCH 30/38] fix(rocrand) use built in functions to get c.u instead of regex --- projects/rocrand/tuning/tuner/base_tuner.py | 12 +++--------- 1 file changed, 3 insertions(+), 9 deletions(-) diff --git a/projects/rocrand/tuning/tuner/base_tuner.py b/projects/rocrand/tuning/tuner/base_tuner.py index 9f2c414483b7..f2cc815d4594 100644 --- a/projects/rocrand/tuning/tuner/base_tuner.py +++ b/projects/rocrand/tuning/tuner/base_tuner.py @@ -189,15 +189,8 @@ def cli(cls): cls(defaults).tune() def _get_grid_sizes(self) -> str: - rocminfo_out = subprocess.check_output("rocminfo", encoding="utf-8") - match = re.search( - r"^\s*Name:\s*gfx\d+.*?^\s*Compute Unit:\s*(\d+)", - rocminfo_out, - flags=re.MULTILINE | re.DOTALL, - ) - if not match: - raise Exception("Could not find Compute Unit info in rocminfo output") - num_compute_units = int(match.group(1)) + num_compute_units = self.device_properties.multiProcessorCount + compute_unit_multipliers = [4, 5, 8, 10, 16, 32] min_grid_size = 128 max_grid_size = 4096 @@ -209,6 +202,7 @@ def _get_grid_sizes(self) -> str: grid_sizes.append(new_grid_size) grid_sizes = list(set(grid_sizes)) # Unique grid_sizes.sort() + return grid_sizes def _get_tune_params(self) -> OrderedDict: From bfba25b68e82bbd0251d4ec3cd41da123cdda6da Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Wed, 7 Oct 2026 12:26:26 -0400 Subject: [PATCH 31/38] fix(rocrand): Made running compare simulated strategies anywhere possible --- projects/rocrand/tuning/utils/compare_simulated_strategies.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/rocrand/tuning/utils/compare_simulated_strategies.py b/projects/rocrand/tuning/utils/compare_simulated_strategies.py index bcb321042a39..10e4f94f9d2a 100644 --- a/projects/rocrand/tuning/utils/compare_simulated_strategies.py +++ b/projects/rocrand/tuning/utils/compare_simulated_strategies.py @@ -28,6 +28,7 @@ import pandas as pd import seaborn as sns import matplotlib.pyplot as plt +from utils import BASE_DIR """ This script performs comparative analysis of different optimization strategies for kernel tuning simulations. @@ -239,7 +240,7 @@ def create_performance_heatmap(df, output_filename, algo, arch): def main(): args = parse_arguments() - base_dir = "." + base_dir = BASE_DIR simulated_output_dir = os.path.join(base_dir, "../simulated_output") output_dir = os.path.join(base_dir, "../output") From 8202c38a76fef3fb3c6bddf282470076a2c79b35 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Wed, 7 Oct 2026 14:06:47 -0400 Subject: [PATCH 32/38] fix(rocrand): Make sure existing archs are not overwritten --- projects/rocrand/tuning/confgen/generate.py | 36 +++++++++++++-------- 1 file changed, 22 insertions(+), 14 deletions(-) diff --git a/projects/rocrand/tuning/confgen/generate.py b/projects/rocrand/tuning/confgen/generate.py index cff9aa380bd9..d098227e8167 100644 --- a/projects/rocrand/tuning/confgen/generate.py +++ b/projects/rocrand/tuning/confgen/generate.py @@ -73,6 +73,15 @@ def main(): filter_dif = len(data["cache"]) - len(data_filter) if filter_dif > 0: log.warning(f"Skipping failed compilations: {filter_dif}") + + # Filter/skip data entries that has an acrchitecture thats not targeted with -t + arch = data["arch_name"] + if arch not in gfx_target: + log.warning( + f"Skipping {file_path}: arch {arch} not in --target-arch {gfx_target}" + ) + continue + # Find best config min_configs = min(data_filter, key=lambda c: c["time"]) # Drop unrelated entries @@ -92,20 +101,19 @@ def main(): # Ensure algorithm entry exists alg: str = data["algo_name"] - if alg not in algs: - algs[alg] = {} - arch = data["arch_name"] - for target in gfx_target: - if target != arch: - continue - # Ensure target entry exists - if target not in algs[alg]: - algs[alg][target] = {} - - config = make_config(data) - config["tune_params"] = min_configs - # Add config entry - algs[alg][target] = config + + config = make_config(data) + config["tune_params"] = min_configs + algs.setdefault(alg, {})[arch] = config + + if not algs: + log.error(f"No tuning reslults found for targets {gfx_target} in {input_glob}") + quit(1) + + found_archs = {arch for targets in algs.values() for arch in targets} + for target in gfx_target: + if target not in found_archs: + log.warning(f"No tuning results found for --target-arch {target}") log.info(f"Parsed {len(algs)} different algorithm(s)!") From d347f8d5118e6fc66052ab5f968f17ca43a72576 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Wed, 7 Oct 2026 14:11:14 -0400 Subject: [PATCH 33/38] fix(rocrand): use geometric mean instead of arithmetic mean --- .../rocrand/tuning/tuner/templates/common/base.cpp.jinja2 | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 index 79b1194a911b..95a18aaa062b 100644 --- a/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 +++ b/projects/rocrand/tuning/tuner/templates/common/base.cpp.jinja2 @@ -209,10 +209,10 @@ extern "C" float {{algo_name}}_wrapper( {% block calculate_time %} std::vector bps = executor.get_all_bytes_per_second(); - double total = 0.0; + double prod = 1.0; for(double x : bps) - total += (1000.0 * primbench::GiB) / x; - return total / static_cast(bps.size()); + prod *= (1000.0 * primbench::GiB) / x; + return static_cast(std::pow(prod, 1.0 / static_cast(bps.size()))); {% endblock %} } From 4cca9e00e05de93446aba8c70cb37edc6ba42c3f Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Wed, 7 Oct 2026 14:17:25 -0400 Subject: [PATCH 34/38] fix(rocrand): ensure lfsr113 only get pow2 grid size --- projects/rocrand/tuning/tuner/tuning_lfsr113.py | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/projects/rocrand/tuning/tuner/tuning_lfsr113.py b/projects/rocrand/tuning/tuner/tuning_lfsr113.py index c1a6aca59748..95848986787d 100644 --- a/projects/rocrand/tuning/tuner/tuning_lfsr113.py +++ b/projects/rocrand/tuning/tuner/tuning_lfsr113.py @@ -36,6 +36,13 @@ def _get_default_args(cls) -> TunerArgs: def __init__(self, args: TunerArgs) -> None: super().__init__(args) + def _get_restrictions(self): + def validate(params): + g_size = params["grid_size"] + return g_size > 0 and g_size & (g_size - 1) == 0 + + return validate + if __name__ == "__main__": Tuner.cli() From bad2c243813da68e806904cf04eac07bf19676a6 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 8 Oct 2026 11:14:07 -0400 Subject: [PATCH 35/38] fix(rocrand): fix lfsr min thread check --- projects/rocrand/tuning/tuner/tuning_lfsr113.py | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/projects/rocrand/tuning/tuner/tuning_lfsr113.py b/projects/rocrand/tuning/tuner/tuning_lfsr113.py index 95848986787d..5248f89f2b17 100644 --- a/projects/rocrand/tuning/tuner/tuning_lfsr113.py +++ b/projects/rocrand/tuning/tuner/tuning_lfsr113.py @@ -37,8 +37,15 @@ def __init__(self, args: TunerArgs) -> None: super().__init__(args) def _get_restrictions(self): + min_total_threads = 32768 + def validate(params): g_size = params["grid_size"] + b_size = params["block_size_x"] + + if b_size * g_size < min_total_threads: + return False + return g_size > 0 and g_size & (g_size - 1) == 0 return validate From 8bda42ad2083721bc08c4ec9035209cfe666958f Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 8 Oct 2026 12:17:15 -0400 Subject: [PATCH 36/38] fix(rocrand): resolve output path inconsistency --- projects/rocrand/tuning/run_tuning.py | 21 ++++++++++++++++----- projects/rocrand/tuning/utils/utils.py | 2 +- 2 files changed, 17 insertions(+), 6 deletions(-) diff --git a/projects/rocrand/tuning/run_tuning.py b/projects/rocrand/tuning/run_tuning.py index 2200184084a1..6a668d5bd431 100644 --- a/projects/rocrand/tuning/run_tuning.py +++ b/projects/rocrand/tuning/run_tuning.py @@ -85,7 +85,7 @@ def filter_algorithms(available_algos: List[str], pattern: str) -> List[str]: def run_tuning( algorithms: List[str], size: int | None = None, - output_dir: str | None = None, + output_dir: Path | None = None, exclude_default_config: bool | None = None, max_fevals: int | None = None, strategy: str | None = None, @@ -93,14 +93,17 @@ def run_tuning( arch_name: str | None = None, seed: int | None = None, ) -> None: + + if not output_dir: + output_dir = Path(f"{BASE_DIR}/output") + """Run tuning for specified algorithms.""" - current_dir = Path(__file__).parent - tuning_dir = current_dir / "tuner" + tuning_dir = f"{BASE_DIR}/tuner" with working_directory(tuning_dir): for algo in algorithms: print(f"\nRunning tuning for {algo}...") - tuning_file = tuning_dir / f"tuning_{algo}.py" + tuning_file = Path(f"{tuning_dir}/tuning_{algo}.py") if not tuning_file.exists(): print(f"Error: Tuning file not found for {algo}") @@ -114,6 +117,7 @@ def run_tuning( continue print(f"Working directory: {Path.cwd()}") + print(f"Output directory: {output_dir.absolute()}") try: if hasattr(tuning_module, "Tuner"): args: TunerArgs = tuning_module.Tuner._get_default_args() @@ -187,12 +191,19 @@ def main(): print(f" - {algo}") output_dir = Path(args.output_dir) + + if not output_dir.is_absolute(): + output_dir = output_dir.absolute() + else: + # gets rid of any ../ + output_dir = output_dir.resolve() + output_dir.mkdir(parents=True, exist_ok=True) run_tuning( matched_algos, args.size, - args.output_dir, + output_dir, args.exclude_default_config, args.max_fevals, args.strategy, diff --git a/projects/rocrand/tuning/utils/utils.py b/projects/rocrand/tuning/utils/utils.py index fef98080771c..f68df0d7be9a 100644 --- a/projects/rocrand/tuning/utils/utils.py +++ b/projects/rocrand/tuning/utils/utils.py @@ -32,7 +32,7 @@ class Parser: def get_parser( default_bytes: Optional[int] = None, default_max_feval: int = 100, - default_output_dir: str = "../output", + default_output_dir: str = f"{BASE_DIR}/output", default_strategy: str = "dual_annealing", ): parser = argparse.ArgumentParser() From 3e91afcc7b0f100b5cb5aeab26ab39cb6de066f4 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 8 Oct 2026 12:29:23 -0400 Subject: [PATCH 37/38] fix(rocrand): prevent old timing ussage when --size is used --- projects/rocrand/tuning/tuner/base_tuner.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rocrand/tuning/tuner/base_tuner.py b/projects/rocrand/tuning/tuner/base_tuner.py index f2cc815d4594..cb0f2f23370e 100644 --- a/projects/rocrand/tuning/tuner/base_tuner.py +++ b/projects/rocrand/tuning/tuner/base_tuner.py @@ -370,7 +370,7 @@ def _get_base_tune_kernel_args( "strategy": self.strategy, "grid_div_x": self._get_grid_div_x(), "cache": str(self._get_cache_file_path()), - "problem_size": 10, + "problem_size": self.bytes_size, "lang": "C", "compiler": "hipcc", "compiler_options": self._get_compiler_options(), From 417e46a2c9bea8feba17749861f2be60e0d1d15d Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Thu, 8 Oct 2026 13:05:46 -0400 Subject: [PATCH 38/38] fix(rocrand): fix grid size cache while in simulation mode --- projects/rocrand/tuning/run_tuning.py | 2 +- projects/rocrand/tuning/tuner/base_tuner.py | 8 ++++++++ 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/projects/rocrand/tuning/run_tuning.py b/projects/rocrand/tuning/run_tuning.py index 6a668d5bd431..ba00e1f25619 100644 --- a/projects/rocrand/tuning/run_tuning.py +++ b/projects/rocrand/tuning/run_tuning.py @@ -117,7 +117,7 @@ def run_tuning( continue print(f"Working directory: {Path.cwd()}") - print(f"Output directory: {output_dir.absolute()}") + print(f"Output directory: {output_dir.absolute()}") try: if hasattr(tuning_module, "Tuner"): args: TunerArgs = tuning_module.Tuner._get_default_args() diff --git a/projects/rocrand/tuning/tuner/base_tuner.py b/projects/rocrand/tuning/tuner/base_tuner.py index cb0f2f23370e..76ffc31ca3d5 100644 --- a/projects/rocrand/tuning/tuner/base_tuner.py +++ b/projects/rocrand/tuning/tuner/base_tuner.py @@ -147,11 +147,15 @@ def __init__(self, args: TunerArgs): self.device_properties = hip.hipDeviceProp_t() hip.hipGetDeviceProperties(self.device_properties, self.device_id) self.arch_name = self.device_properties.gcnArchName.decode().split(":")[0] + self.cached_grid_sizes = None else: assert ( args.arch_name ), "Simulation mode requires --arch-name to be specified" self.arch_name = args.arch_name + with open(self._get_cache_file_path()) as f: + self.cached_grid_sizes = json.load(f)["tune_params"]["grid_size"] + self.exclude_default_config = args.exclude_default_config self.max_fevals = args.max_fevals self.save_metadata = args.save_metadata @@ -189,6 +193,10 @@ def cli(cls): cls(defaults).tune() def _get_grid_sizes(self) -> str: + + if self.cached_grid_sizes: + return self.cached_grid_sizes + num_compute_units = self.device_properties.multiProcessorCount compute_unit_multipliers = [4, 5, 8, 10, 16, 32]