Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
111 commits
Select commit Hold shift + click to select a range
6334abe
Bump main to 2.9.0.
matyas-streamhpc Oct 9, 2025
1b29c61
Bump main to 3.0.0.
matyas-streamhpc Oct 9, 2025
6e278ad
Require C++17 for compiling hipCUB
matyas-streamhpc Oct 9, 2025
8f719d8
Drop `BinaryFlip` operator
matyas-streamhpc Oct 9, 2025
04b66a2
Deprecate `hipcub::Swap`
matyas-streamhpc Oct 13, 2025
bc64f0b
Drop hipCUB APIs with a debug_synchronous parameter
matyas-streamhpc Oct 15, 2025
9e10195
Update deprecated warning for hipcub::DivideAndRoundUp
matyas-streamhpc Oct 27, 2025
ff14e19
Add `hip::std` support
matyas-streamhpc Nov 24, 2025
a1e67e3
Adds support for large number of items and large number of segments t…
matyas-streamhpc Oct 27, 2025
bdbfba7
Deprecate hipcub::min and hipcub:max and replace internal uses of std…
matyas-streamhpc Dec 5, 2025
b48d6c9
Remove deprecated hipcub::min and hipcub::max
matyas-streamhpc Dec 5, 2025
c86f6a0
Replace rocprim counting iterator to thrust for nvcc compatibility
matyas-streamhpc Oct 30, 2025
712a64b
Deprecate and replace `HIPCUB_IS_INT128_ENABLED`
matyas-streamhpc Nov 3, 2025
8d5cf87
Adds support for large num items to ``DeviceMerge``
matyas-streamhpc Nov 3, 2025
ec5d13f
Remove ``DeviceSpmv``
matyas-streamhpc Nov 3, 2025
f84d66f
Adds support for large number of segments to DeviceSegmentedReduce
matyas-streamhpc Dec 4, 2025
dfaaa89
Remove ``LEGACY_PTX_ARCH``
matyas-streamhpc Nov 3, 2025
ff9736e
Fix hipCUB trait deprecations
matyas-streamhpc Nov 14, 2025
7e435cf
Drop deprecated entities from util_type
matyas-streamhpc Nov 13, 2025
86074bb
Drop ``GridBarrier``
matyas-streamhpc Nov 13, 2025
fe0b4f0
Replace pre-c++17 traits with modern ones
matyas-streamhpc Nov 13, 2025
0a9561b
Replace ``HIPCUB_IF_CONSTEXPR``
matyas-streamhpc Nov 13, 2025
221328e
Drop small deprecated entites
matyas-streamhpc Nov 13, 2025
a66efa1
Remove deprecated iterators
matyas-streamhpc Nov 13, 2025
6355a79
Add `__int128`, `__half` and `hip_bfloat16` overloads to `to_bits` in…
matyas-streamhpc Nov 17, 2025
9a91f43
Fix add implicit conversion for read access for `conditional_discard_…
matyas-streamhpc Nov 24, 2025
ce90eaf
Drop deprecated features from util_ptx.cuh
matyas-streamhpc Dec 1, 2025
2e40dd8
Fix typo in TEST_UTILS_INCLUDE_GUARD
matyas-streamhpc Dec 2, 2025
0e6c33e
Use hip::std limits
matyas-streamhpc Dec 1, 2025
330befa
Move iterator traits to hip::std::iterator_traits
matyas-streamhpc Dec 9, 2025
c6c72bf
Avoid int overflow during multiplication
matyas-streamhpc Dec 3, 2025
5533866
Remove `HIPCUB_MIN` and `HIPCUB_MAX`
matyas-streamhpc Dec 3, 2025
b061442
Minimize usage of `hipcub::Traits`
matyas-streamhpc Dec 8, 2025
d4afe09
Drop deprecated CUB macros
matyas-streamhpc Dec 3, 2025
e7d2dc1
Update mdspan support
matyas-streamhpc Dec 3, 2025
20e875f
Replace util_arch.cuh macros with inline constexpr variables
matyas-streamhpc Dec 3, 2025
2d76d93
Use `_CCCL_PRAGMA_UNROLL_FULL()` and `_CCCL_PRAGMA_NOUNROLL()`
matyas-streamhpc Dec 3, 2025
7b185a8
Allow rapids to avoid unrolling some loops in sort
matyas-streamhpc Dec 3, 2025
a679a27
Adds support for large num items to DeviceMerge
matyas-streamhpc Dec 3, 2025
53749bd
Adds support for large number of buffers to `DeviceCpy::Batched` and …
matyas-streamhpc Dec 3, 2025
269fcbf
Fix hip std namespace
matyas-streamhpc Dec 3, 2025
e084472
Increase test coverage for hipCUB batched device copy and add fix for…
matyas-streamhpc Dec 10, 2025
5fe566b
fix(hipcub): fix broken bitwise representation derivation for extende…
Naraenda Dec 16, 2025
ae865bd
Resolve "Fix is_floating_point value in hipCUB/rocPRIM"
cenxuantian Dec 18, 2025
b5c514a
Update copyright year
matyas-streamhpc Jan 5, 2026
c575d71
Update CHANGELOG
matyas-streamhpc Jan 26, 2026
3929408
Fix clang format
matyas-streamhpc Jan 26, 2026
58ae8ca
Update counting, discard and transform iterator CUB backend for tests
matyas-streamhpc Oct 30, 2025
0e8e87f
Fix assert_near
matyas-streamhpc Jan 14, 2026
67aa414
Fix is_floating_point_v error for half/bfloat16 with CUDA
matyas-streamhpc Jan 7, 2026
045f692
Replace util_arch.cuh macros with inline constexpr variables
matyas-streamhpc Jan 7, 2026
67fda6a
Add cuda headers
matyas-streamhpc Jan 7, 2026
80f7801
Fix device histogram
matyas-streamhpc Jan 7, 2026
87a1424
Fix to_bits
matyas-streamhpc Jan 8, 2026
1f0ccd5
Fix ForEachInExtentsAPI test
matyas-streamhpc Jan 8, 2026
b617dfa
Fix device_test_enabled_for_warp_size_v for CUDA
matyas-streamhpc Jan 8, 2026
5304a56
Fix device_reduce CUB backend
matyas-streamhpc Jan 9, 2026
ee2f24e
Fix cub util_type include
matyas-streamhpc Jan 14, 2026
7380c48
Update WarpExchange rocPRIM backend for CUB compatibility
matyas-streamhpc Jan 15, 2026
3fcd617
Fix for ArgMax/Min
matyas-streamhpc Jan 9, 2026
27b56d2
Fix type in test
matyas-streamhpc Jan 20, 2026
c65b9ff
Fix device select UniqueByKey
matyas-streamhpc Jan 20, 2026
c51d76f
Fix clang format
matyas-streamhpc Jan 20, 2026
ae5c051
Fix type in benchmark segmented sort
matyas-streamhpc Jan 20, 2026
3172f47
Fix types in segmented sort
matyas-streamhpc Jan 20, 2026
491a62d
Device segmented sort - support large indices
matyas-streamhpc Jan 20, 2026
edffcb5
Patch thread operators - add missing headers
matyas-streamhpc Jan 7, 2026
f5a456c
Patch thread operators - add cub backend
matyas-streamhpc Jan 7, 2026
114d47a
Patch thread operators - fix
matyas-streamhpc Jan 22, 2026
9b1e23b
Fix types and formatting in generate resource spec
matyas-streamhpc Jan 26, 2026
6e4b449
Remove internal `hipcub::Division` usage
matyas-streamhpc Jan 28, 2026
18170f8
Internal removal of `hipcub::Inequality`
matyas-streamhpc Jan 28, 2026
27b09d2
Fix sort keys over 4GB
matyas-streamhpc Jan 28, 2026
e185eb8
Fix device radix sort large sizes test
matyas-streamhpc Jan 28, 2026
42688d4
Fix device reduce ArgMin/Max CUB backend
matyas-streamhpc Feb 11, 2026
696d9fb
Fix block merge sort default value management for CUB backend compati…
matyas-streamhpc Feb 17, 2026
dd0380f
Fix block radix rank test for CUB backend
matyas-streamhpc Feb 19, 2026
085a90b
Fix static linking for Windows
matyas-streamhpc Feb 19, 2026
21e5537
Fix Windows build by forcing static MSVC runtime to match vcpkg stati…
matyas-streamhpc Feb 20, 2026
41dc251
Fix Windows HIP cannot support allocations over 4 GiB
matyas-streamhpc Feb 23, 2026
7666541
Remove internal ``hipcub::Difference`` usage
matyas-streamhpc Feb 26, 2026
f67e0a5
Remove internal ``hipcub::InequalityWrapper`` usage
matyas-streamhpc Feb 26, 2026
e38fb41
Remove internal ``hipcub::Equality`` usage
matyas-streamhpc Feb 26, 2026
4a9dd08
Remove internal ``hipcub::Max`` usage
matyas-streamhpc Feb 26, 2026
6ca32ed
Remove internal usage of ``hipcub::Sum``
matyas-streamhpc Mar 2, 2026
41b261a
Remove headers that include internally used thread operators
matyas-streamhpc Mar 5, 2026
433c2e3
Deprecate thread operators
matyas-streamhpc Mar 6, 2026
ca164b7
Remove tests of deprecated thread operators
matyas-streamhpc Mar 9, 2026
f529d58
Deprecate and replace `hipcub::BFE`
matyas-streamhpc Mar 9, 2026
7a62537
Fix libhipcxx version check
matyas-streamhpc Mar 9, 2026
8c8dd1f
Add sort_last for custom test type
matyas-streamhpc Apr 7, 2026
4ce2f2d
Update device reduce test cases
matyas-streamhpc Apr 8, 2026
f548a7f
Add numeric_limits support for custom test types
matyas-streamhpc Apr 8, 2026
19939ba
Skip unsupported bfloat16/half histogram tests
matyas-streamhpc Apr 16, 2026
a1236e6
feat(hipcub): add infrastructure to expose the compatible cccl version
Naraenda May 20, 2026
92a0d72
fix(hipcub): Fix address sanitizer test header for CUDA
matyas-streamhpc Jul 9, 2026
c86f56c
feat(hipcub): use primbench in benchmark_device_batch_copy.cpp
MyNameIsTrez Jul 9, 2026
e5f3a4c
feat(hipcub): Google Benchmark is no longer a dependency
MyNameIsTrez Jul 9, 2026
c7fddcb
feat(hipcub): use primbench in 6 block algorithms
MyNameIsTrez Jul 9, 2026
4f2f28f
feat(hipcub): use primbench in 5 block algorithms
MyNameIsTrez Jul 9, 2026
05aa139
feat(hipcub): use primbench in 5 device algorithms
MyNameIsTrez Jul 9, 2026
59644ed
feat(hipcub): use primbench in 5 device algorithms
MyNameIsTrez Jul 9, 2026
d4f0dcb
feat(hipcub): use primbench in 7 device algorithms
MyNameIsTrez Jul 9, 2026
ee2ae71
feat(hipcub): use primbench in 6 warp algorithms
MyNameIsTrez Jul 9, 2026
2088805
chore(hipcub): remove misleading comments
MyNameIsTrez Jul 13, 2026
95c300b
fix(hipcub): hipMalloc allocated the wrong size
MyNameIsTrez Jul 13, 2026
bc41ee2
Merge branch 'develop' into users/Nemi26/use-primbench-in-hipcub
Naraenda Jul 14, 2026
e965109
fix(hipcub): use int64_t instead of long long
MyNameIsTrez Jul 16, 2026
0427ff4
feat(hipcub): add run_benchmarks.py
MyNameIsTrez Jul 16, 2026
08b0136
fix(hipcub): replace usage of uint64_t with unsigned long long
MyNameIsTrez Jul 22, 2026
d07f10f
fix(hipcub): remove includes that break portability
MyNameIsTrez Jul 22, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions projects/hipcub/.clang-format
Original file line number Diff line number Diff line change
Expand Up @@ -167,4 +167,10 @@ Macros:
- HIPCUB_DETAIL_DEPRECATED_DEBUG_SYNCHRONOUS=[[DETAIL_DEPRECATED_DEBUG_SYNCHRONOUS___]]
BreakAfterAttributes: Always

WhitespaceSensitiveMacros: [
'HIPCUB_HAS_INCLUDE',
'_HIPCUB_LIBCXX_INCLUDE',
'_HIPCUB_STD_INCLUDE'
]

---
369 changes: 369 additions & 0 deletions projects/hipcub/.gitlab/run_benchmarks.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,369 @@
#!/usr/bin/env python3

# Copyright (c) 2022-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
import json
import os
import re
import stat
import subprocess
import sys


def run_benchmarks(
benchmark_executables_dir,
gpu_architecture,
json_out_dir,
csv_out_dir,
benchmark_executable_filter,
skip_gathered,
size,
hot,
seed,
filter,
dry,
min_gpu_ms_per_batch,
min_secs,
noise_timeout_secs,
batch_window_size,
noise_tolerance_percent,
min_gpu_temp,
max_gpu_temp,
max_warming_secs,
max_cooling_secs,
output_hip_device_properties_context,
output_amdsmi_context,
output_batches,
spaces_per_indent,
stream_blocking_timeout_secs,
):
def is_benchmark_executable(filename):
if not re.search(benchmark_executable_filter, filename):
return False

path = os.path.join(benchmark_executables_dir, filename)
st_mode = os.stat(path).st_mode

# Return True when any execution flag is set, AND it is a regular file.
# Doesn't check permissions.
return (st_mode & (stat.S_IXUSR | stat.S_IXGRP | stat.S_IXOTH)) and (
st_mode & stat.S_IFREG
)

def should_skip(json_out_path):
if not skip_gathered:
return False

try:
with open(json_out_path) as f:
json.load(f)
except (FileNotFoundError, json.JSONDecodeError):
return False

return True

# Create output directories if they don't exist
os.makedirs(json_out_dir, exist_ok=True)
if csv_out_dir:
os.makedirs(csv_out_dir, exist_ok=True)

benchmark_names = [
name
for name in os.listdir(benchmark_executables_dir)
if is_benchmark_executable(name)
]
print(
"The following benchmarks will be run:\n{}".format("\n".join(benchmark_names)),
file=sys.stderr,
flush=True,
)

success = True

for benchmark_name in benchmark_names:
results_json_name = f"{benchmark_name}_{gpu_architecture}.json"
json_out_path = os.path.join(json_out_dir, results_json_name)

if should_skip(json_out_path):
print(
f"Skipping {benchmark_name}, because its results have already been gathered at {json_out_path}",
file=sys.stderr,
flush=True,
)
continue

benchmark_path = os.path.join(benchmark_executables_dir, benchmark_name)

args = [benchmark_path]
args += ["--json-out", json_out_path]

if size:
args += ["--size", size]
if hot:
args += ["--hot"]
if seed:
args += ["--seed", seed]
if csv_out_dir:
results_csv_name = f"{benchmark_name}_{gpu_architecture}.csv"
csv_out_path = os.path.join(csv_out_dir, results_csv_name)
args += ["--csv-out", csv_out_path]
if filter:
args += ["--filter", filter]
if dry:
args += ["--dry"]
if min_gpu_ms_per_batch:
args += ["--min-gpu-ms-per-batch", min_gpu_ms_per_batch]
if min_secs:
args += ["--min-secs", min_secs]
if noise_timeout_secs:
args += ["--noise-timeout-secs", noise_timeout_secs]
if batch_window_size:
args += ["--batch-window-size", batch_window_size]
if noise_tolerance_percent:
args += ["--noise-tolerance-percent", noise_tolerance_percent]
if min_gpu_temp:
args += ["--min-gpu-temp", min_gpu_temp]
if max_gpu_temp:
args += ["--max-gpu-temp", max_gpu_temp]
if max_warming_secs:
args += ["--max-warming-secs", max_warming_secs]
if max_cooling_secs:
args += ["--max-cooling-secs", max_cooling_secs]
if output_hip_device_properties_context:
args += ["--output-hip-device-properties-context"]
if output_amdsmi_context:
args += ["--output-amdsmi-context"]
if output_batches:
args += ["--output-batches"]
if spaces_per_indent:
args += ["--spaces-per-indent", spaces_per_indent]
if stream_blocking_timeout_secs:
args += ["--stream-blocking-timeout-secs", stream_blocking_timeout_secs]

try:
subprocess.check_call(args)
except subprocess.CalledProcessError as error:
print(
f'Could not run benchmark at {benchmark_path}. Error: "{error}"',
file=sys.stderr,
flush=True,
)
success = False

return success


def main():
parser = argparse.ArgumentParser()

# Python arguments
parser.add_argument(
"--benchmark-executables-dir",
help="Benchmark executables directory",
required=True,
)
parser.add_argument(
"--gpu-architecture",
help="Architecture of the currently enabled GPU",
required=True,
)
parser.add_argument(
"--json-out-dir",
help="Directory to write JSON result files to, creating parent directories if necessary",
required=True,
)
parser.add_argument(
"--csv-out-dir",
help="Directory to write CSV result files to, creating parent directories if necessary",
required=False,
)
parser.add_argument(
"--benchmark-executable-filter",
help="Regular expression that controls the list of benchmark executables to run",
default=r"^benchmark",
required=False,
)
parser.add_argument(
"--skip-gathered",
help="Skip running benchmarks whose JSON file has already been gathered",
default=False,
action="store_true",
required=False,
)

# primbench arguments
parser.add_argument(
"--size",
help="Input size. Benchmarks decide what this represents, but it is commonly the number of bytes or items",
default="",
required=False,
)
parser.add_argument(
"--hot",
help="Skip clearing the GPU cache between batch iterations",
default=False,
action="store_true",
required=False,
)
parser.add_argument(
"--seed", help="Seed used for input generation", default="", required=False
)
parser.add_argument(
"--filter",
help="Regex filter of specialization names to benchmark",
default="",
required=False,
)
parser.add_argument(
"--dry",
help="Perform a dry run. The benchmark setup is still run, and JSON and CSV files are still output, but state.run() immediately returns",
default=False,
action="store_true",
required=False,
)
parser.add_argument(
"--min-gpu-ms-per-batch",
help="Minimum duration of a batch in milliseconds (GPU time)",
default="",
required=False,
)
parser.add_argument(
"--min-secs",
help="Minimum total benchmark duration in seconds (wall time)",
default="",
required=False,
)
parser.add_argument(
"--noise-timeout-secs",
help="Maximum total benchmark duration in seconds before timing out a noisy run (wall time)",
default="",
required=False,
)
parser.add_argument(
"--batch-window-size",
help="Number of batch times used in the noise (coefficient of variation) window to decide early stopping",
default="",
required=False,
)
parser.add_argument(
"--noise-tolerance-percent",
help="Noise tolerance of batch times in percent, used to determine early benchmark stopping",
default="",
required=False,
)
parser.add_argument(
"--min-gpu-temp",
help="Minimum GPU temperature in °C. Too low slows benchmarks; too high increases noise",
default="",
required=False,
)
parser.add_argument(
"--max-gpu-temp",
help="Maximum GPU temperature in °C. Too low slows benchmarks; too high increases noise",
default="",
required=False,
)
parser.add_argument(
"--max-warming-secs",
help="Maximum seconds allowed for GPU warming before an error is thrown",
default="",
required=False,
)
parser.add_argument(
"--max-cooling-secs",
help="Maximum seconds allowed for GPU cooling before an error is thrown",
default="",
required=False,
)
parser.add_argument(
"--output-hip-device-properties-context",
help="Output a `hip_device_properties` object in the context, containing GPU details",
default=False,
action="store_true",
required=False,
)
parser.add_argument(
"--output-amdsmi-context",
help="Output an `amdsmi` object in the context, containing GPU details",
default=False,
action="store_true",
required=False,
)
parser.add_argument(
"--output-batches",
help="Output a `batches` array for each specialization, containing per-batch details",
default=False,
action="store_true",
required=False,
)
parser.add_argument(
"--spaces-per-indent",
help="Number of spaces per indentation level in JSON output. Set to 0 for no indentation",
default="",
required=False,
)
parser.add_argument(
"--stream-blocking-timeout-secs",
help="Maximum stream blocking duration in seconds before timing out. Stream is blocked while queueing kernel calls. Use `primbench::flags::sync` if kernel is synchronous",
default="",
required=False,
)

args = parser.parse_args()

benchmark_run_successful = run_benchmarks(
args.benchmark_executables_dir,
args.gpu_architecture,
args.json_out_dir,
args.csv_out_dir,
args.benchmark_executable_filter,
args.skip_gathered,
args.size,
args.hot,
args.seed,
args.filter,
args.dry,
args.min_gpu_ms_per_batch,
args.min_secs,
args.noise_timeout_secs,
args.batch_window_size,
args.noise_tolerance_percent,
args.min_gpu_temp,
args.max_gpu_temp,
args.max_warming_secs,
args.max_cooling_secs,
args.output_hip_device_properties_context,
args.output_amdsmi_context,
args.output_batches,
args.spaces_per_indent,
args.stream_blocking_timeout_secs,
)

return benchmark_run_successful


if __name__ == "__main__":
success = main()
if success:
exit(0)
else:
exit(1)
Loading
Loading