amd-jmahovsky · Oct 19, 2026 · Oct 19, 2026 · Oct 19, 2026 · Oct 19, 2026 · Oct 19, 2026
@@ -101,6 +101,7 @@ ROCPRIM_KERNEL ROCPRIM_LAUNCH_BOUNDS(ROCPRIM_DEFAULT_MAX_BLOCK_SIZE) void
 template<lookback_scan_determinism Determinism,
          typename Config,
          typename Selector,
+         typename AccumulatorType,
          typename KeysInputIterator,
          typename ValuesInputIterator,
          typename UniqueOutputIterator,
@@ -121,8 +122,6 @@ hipError_t reduce_by_key_impl_wrapped_config(void*                     temporary
                                              const hipStream_t         stream,
                                              const bool                debug_synchronous)
 {
-    using accumulator_type = reduce_by_key::accumulator_type_t<ValuesInputIterator, BinaryFunction>;
-
     bool use_atomic_block_id;
     ROCPRIM_RETURN_ON_ERROR(check_if_using_atomic_block_id(stream, use_atomic_block_id));
     const auto use_atomic_block_id_variant
@@ -142,7 +141,7 @@ hipError_t reduce_by_key_impl_wrapped_config(void*                     temporary
             const auto params = get_config<Selector>(Config{}, current_target);
 
             using scan_state_type
-                = reduce_by_key::lookback_scan_state_t<accumulator_type, use_sleepy_scan>;
+                = reduce_by_key::lookback_scan_state_t<AccumulatorType, use_sleepy_scan>;
             const unsigned int block_size      = params.kernel_config.block_size;
             const unsigned int items_per_block = block_size * params.kernel_config.items_per_thread;
 
@@ -162,7 +161,7 @@ hipError_t reduce_by_key_impl_wrapped_config(void*                     temporary
             // The number of segment heads in previous launches.
             std::size_t* d_global_head_count = nullptr;
             // The running accumulation across the launch boundary.
-            accumulator_type* d_previous_accumulated = nullptr;
+            AccumulatorType* d_previous_accumulated = nullptr;
 
             detail::temp_storage::layout layout{};
             ROCPRIM_RETURN_ON_ERROR(
@@ -291,6 +290,7 @@ hipError_t reduce_by_key_impl_wrapped_config(void*                     temporary
 
 template<lookback_scan_determinism Determinism,
          typename Config,
+         typename AccumulatorType,
          typename KeysInputIterator,
          typename ValuesInputIterator,
          typename UniqueOutputIterator,
@@ -311,23 +311,23 @@ hipError_t reduce_by_key_impl(void*                     temporary_storage,
                               const hipStream_t         stream,
                               const bool                debug_synchronous)
 {
-    using key_type         = ::rocprim::detail::value_type_t<KeysInputIterator>;
-    using accumulator_type = reduce_by_key::accumulator_type_t<ValuesInputIterator, BinaryFunction>;
-    using Selector = reduce_by_key_config_selector<key_type, accumulator_type, BinaryFunction>;
-
-    return detail::reduce_by_key_impl_wrapped_config<Determinism, Config, Selector>(
-        temporary_storage,
-        storage_size,
-        keys_input,
-        values_input,
-        size,
-        unique_output,
-        aggregates_output,
-        unique_count_output,
-        reduce_op,
-        key_compare_op,
-        stream,
-        debug_synchronous);
+    using key_type = ::rocprim::detail::value_type_t<KeysInputIterator>;
+    using Selector = reduce_by_key_config_selector<key_type, AccumulatorType, BinaryFunction>;
+
+    return detail::
+        reduce_by_key_impl_wrapped_config<Determinism, Config, Selector, AccumulatorType>(
+            temporary_storage,
+            storage_size,
+            keys_input,
+            values_input,
+            size,
+            unique_output,
+            aggregates_output,
+            unique_count_output,
+            reduce_op,
+            key_compare_op,
+            stream,
+            debug_synchronous);
 }
 
 } // namespace detail
@@ -466,19 +466,22 @@ inline hipError_t reduce_by_key(void*                     temporary_storage,
                                 hipStream_t               stream            = 0,
                                 bool                      debug_synchronous = false)
 {
-    return detail::reduce_by_key_impl<detail::lookback_scan_determinism::default_determinism,
-                                      Config>(temporary_storage,
-                                              storage_size,
-                                              keys_input,
-                                              values_input,
-                                              size,
-                                              unique_output,
-                                              aggregates_output,
-                                              unique_count_output,
-                                              reduce_op,
-                                              key_compare_op,
-                                              stream,
-                                              debug_synchronous);
+    return detail::reduce_by_key_impl<
+        detail::lookback_scan_determinism::default_determinism,
+        Config,
+        detail::reduce_by_key::accumulator_type_t<ValuesInputIterator, BinaryFunction>>(
+        temporary_storage,
+        storage_size,
+        keys_input,
+        values_input,
+        size,
+        unique_output,
+        aggregates_output,
+        unique_count_output,
+        reduce_op,
+        key_compare_op,
+        stream,
+        debug_synchronous);
 }
 
 /// \brief Bitwise-reproducible parallel reduce-by-key primitive for device level.
@@ -512,7 +515,10 @@ inline hipError_t deterministic_reduce_by_key(void*                     temporar
                                               hipStream_t stream            = 0,
                                               bool        debug_synchronous = false)
 {
-    return detail::reduce_by_key_impl<detail::lookback_scan_determinism::deterministic, Config>(
+    return detail::reduce_by_key_impl<
+        detail::lookback_scan_determinism::deterministic,
+        Config,
+        detail::reduce_by_key::accumulator_type_t<ValuesInputIterator, BinaryFunction>>(
         temporary_storage,
         storage_size,
         keys_input,

@@ -83,18 +83,19 @@ hipError_t run_length_encode_impl(void*                     temporary_storage,
     return detail::reduce_by_key_impl_wrapped_config<
         detail::lookback_scan_determinism::nondeterministic,
         typename select_reduce_by_key_config<Config>::type,
-        Selector>(temporary_storage,
-                  storage_size,
-                  keys_input,
-                  values_input,
-                  size,
-                  unique_output,
-                  aggregates_output,
-                  unique_count_output,
-                  reduce_op,
-                  key_compare_op,
-                  stream,
-                  debug_synchronous);
+        Selector,
+        accumulator_type>(temporary_storage,
+                          storage_size,
+                          keys_input,
+                          values_input,
+                          size,
+                          unique_output,
+                          aggregates_output,
+                          unique_count_output,
+                          reduce_op,
+                          key_compare_op,
+                          stream,
+                          debug_synchronous);
 }
 
 template<typename Config,

@@ -119,6 +119,7 @@ add_thrust_benchmark(bench/copy/if.cu)
 add_thrust_benchmark(bench/equal/basic.cu)
 add_thrust_benchmark(bench/fill/basic.cu)
 add_thrust_benchmark(bench/for_each/basic.cu)
+add_thrust_benchmark(bench/generate/basic.cu)
 add_thrust_benchmark(bench/inner_product/basic.cu)
 add_thrust_benchmark(bench/merge/basic.cu)
 add_thrust_benchmark(bench/partition/basic.cu)
@@ -145,7 +146,11 @@ add_thrust_benchmark(bench/sort/pairs.cu)
 add_thrust_benchmark(bench/sort/pairs_custom.cu)
 add_thrust_benchmark(bench/tabulate/basic.cu)
 add_thrust_benchmark(bench/transform/basic.cu)
+add_thrust_benchmark(bench/transform/if.cu)
+add_thrust_benchmark(bench/transform/if_stencil.cu)
+add_thrust_benchmark(bench/transform/zip_transform.cu)
 add_thrust_benchmark(bench/transform_reduce/sum.cu)
+add_thrust_benchmark(bench/uninitialized_copy/basic.cu)
 add_thrust_benchmark(bench/unique/basic.cu)
 add_thrust_benchmark(bench/unique/by_key.cu)
 add_thrust_benchmark(bench/vectorized_search/basic.cu)

@@ -1,30 +1,6 @@
-/******************************************************************************
- * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc.  All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
- *     * Redistributions of source code must retain the above copyright
- *       notice, this list of conditions and the following disclaimer.
- *     * Redistributions in binary form must reproduce the above copyright
- *       notice, this list of conditions and the following disclaimer in the
- *       documentation and/or other materials provided with the distribution.
- *     * Neither the name of the NVIDIA CORPORATION nor the
- *       names of its contributors may be used to endorse or promote products
- *       derived from this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
- * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
- * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- *
- ******************************************************************************/
+// SPDX-FileCopyrightText: Copyright (c) 2011-2023, NVIDIA CORPORATION. All rights reserved.
+// SPDX-FileCopyrightText: Modifications Copyright (c) 2024-2026 Advanced Micro Devices, Inc. All rights reserved.
+// SPDX-License-Identifier: BSD-3
 
 #include <thrust/adjacent_difference.h>
 #include <thrust/device_vector.h>

@@ -1,30 +1,6 @@
-/******************************************************************************
- * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc.  All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
- *     * Redistributions of source code must retain the above copyright
- *       notice, this list of conditions and the following disclaimer.
- *     * Redistributions in binary form must reproduce the above copyright
- *       notice, this list of conditions and the following disclaimer in the
- *       documentation and/or other materials provided with the distribution.
- *     * Neither the name of the NVIDIA CORPORATION nor the
- *       names of its contributors may be used to endorse or promote products
- *       derived from this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
- * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
- * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- *
- ******************************************************************************/
+// SPDX-FileCopyrightText: Copyright (c) 2011-2023, NVIDIA CORPORATION. All rights reserved.
+// SPDX-FileCopyrightText: Modifications Copyright (c) 2024-2026 Advanced Micro Devices, Inc. All rights reserved.
+// SPDX-License-Identifier: BSD-3
 
 #include <thrust/adjacent_difference.h>
 #include <thrust/device_vector.h>

@@ -1,30 +1,6 @@
-/******************************************************************************
- * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc.  All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
- *     * Redistributions of source code must retain the above copyright
- *       notice, this list of conditions and the following disclaimer.
- *     * Redistributions in binary form must reproduce the above copyright
- *       notice, this list of conditions and the following disclaimer in the
- *       documentation and/or other materials provided with the distribution.
- *     * Neither the name of the NVIDIA CORPORATION nor the
- *       names of its contributors may be used to endorse or promote products
- *       derived from this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
- * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
- * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- *
- ******************************************************************************/
+// SPDX-FileCopyrightText: Copyright (c) 2011-2023, NVIDIA CORPORATION. All rights reserved.
+// SPDX-FileCopyrightText: Modifications Copyright (c) 2024-2026 Advanced Micro Devices, Inc. All rights reserved.
+// SPDX-License-Identifier: BSD-3
 
 #include <thrust/adjacent_difference.h>
 #include <thrust/device_vector.h>

@@ -1,30 +1,6 @@
-/******************************************************************************
- * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc.  All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
- *     * Redistributions of source code must retain the above copyright
- *       notice, this list of conditions and the following disclaimer.
- *     * Redistributions in binary form must reproduce the above copyright
- *       notice, this list of conditions and the following disclaimer in the
- *       documentation and/or other materials provided with the distribution.
- *     * Neither the name of the NVIDIA CORPORATION nor the
- *       names of its contributors may be used to endorse or promote products
- *       derived from this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
- * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
- * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- *
- ******************************************************************************/
+// SPDX-FileCopyrightText: Copyright (c) 2011-2023, NVIDIA CORPORATION. All rights reserved.
+// SPDX-FileCopyrightText: Modifications Copyright (c) 2024-2026 Advanced Micro Devices, Inc. All rights reserved.
+// SPDX-License-Identifier: BSD-3
 
 // Benchmark utils
 #include "bench_utils.hpp"

@@ -1,30 +1,6 @@
-/******************************************************************************
- * Copyright (c) 2011-2023, NVIDIA CORPORATION.  All rights reserved.
- * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc.  All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
- *     * Redistributions of source code must retain the above copyright
- *       notice, this list of conditions and the following disclaimer.
- *     * Redistributions in binary form must reproduce the above copyright
- *       notice, this list of conditions and the following disclaimer in the
- *       documentation and/or other materials provided with the distribution.
- *     * Neither the name of the NVIDIA CORPORATION nor the
- *       names of its contributors may be used to endorse or promote products
- *       derived from this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
- * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
- * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- *
- ******************************************************************************/
+// SPDX-FileCopyrightText: Copyright (c) 2011-2023, NVIDIA CORPORATION. All rights reserved.
+// SPDX-FileCopyrightText: Modifications Copyright (c) 2024-2026 Advanced Micro Devices, Inc. All rights reserved.
+// SPDX-License-Identifier: BSD-3
 
 #include <thrust/copy.h>
 #include <thrust/count.h>
+12 −7		test/libcudacxx/cuda/atomics/atomic.local.pass.cpp
+68 −25		test/utils/amd/hiprtc/hiprtcc.cpp