[RFE] `DeviceSegmentedSort`: user requests, workloads & feature prioritization

### Help Us Collect and Prioritize Workloads for `DeviceSegmentedSort`

We are reworking `cub::DeviceSegmentedSort` (see [#10577](https://github.com/NVIDIA/cccl/issues/10577)) and we'd like your input on **which workloads and which features to prioritize**.

Please **copy & paste the form below into a new comment** and fill it out. Use the checkboxes (`[x]`) to select
what applies (multiple selections are fine), feel free to modify the form as you see fit, and add clarifying notes inline. Partial answers are also welcome, e.g., even just the segment-size distribution and the data type is useful to us.

🗒️ Segmented Sort Workload & Feature Specification Form

````markdown

# 🧮 Segmented Sort Workload & Feature Specification Form

**Use case:**
```
Briefly describe how or where you sort segments (library/application, what the segments represent,
where they come from, what happens to the result).
```

**What do you use today?**
- [ ] `cub::DeviceSegmentedSort`
- [ ] `cub::DeviceSegmentedRadixSort`
- [ ] `cub::DeviceRadixSort` / `cub::DeviceMergeSort` with the segment id packed into the key
- [ ] `thrust::sort` / `thrust::stable_sort` with a custom comparator
- [ ] A hand-written kernel
- [ ] Nothing yet — this is a new use case
- [ ] Other: ☐ (please specify)

```
If you moved *away* from cub::DeviceSegmentedSort, or evaluated it and chose something else,
please tell us why — this is the single most useful thing you can share with us.
```

---

## 📦 Data Layout

**How are the segments laid out in memory?**
- [ ] **Contiguous** — all segments live in one buffer, described by offsets (i.e., what
`DeviceSegmentedSort` supports today)
- [ ] **Disaggregate / batched** — each segment is its own allocation, addressed through an
array of pointers or an iterator-of-iterators (what `DeviceBatchedTopK` supports)
- [ ] Both, depending on the call site

▶️ If contiguous, how are the offsets given?

- [ ] A single consecutive offsets array (`d_offsets`, `d_offsets + 1` — segment `i` is `[o[i], o[i+1])`)
- [ ] Two independent `begin` / `end` arrays (segments may have gaps between them, or overlap)
- [ ] Computed on the fly (transform / counting iterator, e.g. fixed stride)

**Can total items exceed 2^31?**
- [ ] Yes
- [ ] No

**Are empty (zero-length) segments common in your input?**
- [ ] Yes, frequently
- [ ] Rarely
- [ ] Never

---

## 🔑 What Is Being Sorted

- [ ] **Keys only**
- [ ] **Key–value pairs** — value type / size: ☐ (please specify)
- [ ] **Index sort (argsort)** — I only want the *permutation* (gather map); the sorted keys themselves
are discarded

**Number of key columns:**
- [ ] One key per element
- [ ] Multiple keys per element, compared lexicographically (i.e. a multi-column / tuple sort)

---

## 💾 Key Data Types

**Common key types** (select all that apply):
- [ ] **Built-in / fundamental types**: standard integer and floating-point types (`uint8_t`, `int32_t`,
`int64_t`, `float`, `double`, …) — please list which: ☐
- [ ] **Extended floating-point types**: `__half`, `__nv_bfloat16`, fp8, …
- [ ] **128-bit types**: `__int128`, `__uint128`, fixed-point decimals
- [ ] **Custom data types** (structs, packed keys, …)
- [ ] **Specific subset only**: I primarily care about a few specific types: ☐ (please list them)

▶️ If custom types, how should they be ordered?

- [ ] Via a **decomposer** (the type can be reduced to a tuple of fundamental fields, so a radix sort works)
- [ ] Via a **custom comparator**

**Do you need control over the ordering of special values (NaN, ±0, nulls/sentinels)?**
- [ ] Yes — please describe: ☐
- [ ] No

---

## ⚙️ Basic Algorithmic Properties

**Sort type:**
- [ ] **Radix-based**
- [ ] **Comparison-based** (I need to supply a custom comparator)
- [ ] Either — I only care about performance

▶️ If radix-based, please specify:

**Bit range:** would you like to restrict the sort to a sub-range of the key's bits (`begin_bit` / `end_bit`,
as `DeviceSegmentedRadixSort` offers but `DeviceSegmentedSort` currently does not)?
- [ ] Yes — my keys have known constant high bits, so this would save passes
- [ ] No

**Stability** (equal keys retain their relative input order):
- [ ] Required
- [ ] Not required — I would take a faster unstable sort
- [ ] Depends on the call site
- [ ] Varies **per segment** (some segments ascending, some descending, within one call)

**Sort order:**
- [ ] Varies **per segment** (some segments ascending, some descending, within one call)

**Determinism** (bit-identical results run to run and across GPUs):
- [ ] Required
- [ ] Not required — I would take a faster non-deterministic variant

---

## 🏷️ Parameter Variability (argument annotations)

CCCL's newer batched APIs let you annotate *what you know* about a parameter, and specialize on it. The more
precisely you can describe a parameter, the more the algorithm can specialize. The four forms are:

| Form | Meaning |
| --- | --- |
| `cuda::args::constant{}` | Fixed at **compile time** |
| `cuda::args::immediate{v}` | A single value known on the **host** when the call is made |
| `cuda::args::deferred{it}` | A single value **produced on the device** by a preceding launch (same for all segments) |
| `cuda::args::deferred_sequence{it}` | A **distinct value per segment**, also read in stream order |

Please mark, for each parameter the form you can provide. (If a parameter is compile-time constant
for you, say so — that is the most valuable answer.)

| Parameter | compile-time | host value | device value, uniform | per-segment |
| --------- | :----------: | :--------: | :-------------------: | :---------: |
| Number of segments | ☐ | ☐ | ☐ | n/a |
| Segment size | ☐ | ☐ | ☐ | ☐ |
| Total number of items | ☐ | ☐ | ☐ | n/a |
| Sort order (asc/desc) | ☐ | ☐ | ☐ | ☐ |
| Bit range (if radix) | ☐ | ☐ | ☐ | ☐ |

**Do you know a tight compile-time upper bound on the maximum segment size?**
- [ ] Yes — the bound is: ☐
- [ ] Only a runtime bound (known on the host before the call): ☐
- [ ] Only a runtime bound (known on the device): ☐
- [ ] No bound at all

**Are all segments the same size?**
- [ ] Yes, and the size is a compile-time constant: ☐
- [ ] Yes, but the size is only known at runtime
- [ ] No, sizes vary

---

## 📏 Typical Problem Sizes

**Total number of items across all segments (`N`):**
- [ ] 1 – 999
- [ ] 1 000 – 99 999
- [ ] 100 000 – 999 999
- [ ] 1 000 000 – 9 999 999
- [ ] 10 000 000 – 99 999 999
- [ ] 100 000 000 – 999 999 999
- [ ] ≥ 1 000 000 000
- [ ] Variable / depends on workload
- [ ] Custom: ☐

**Number of segments:**
- [ ] 1 – 9
- [ ] 10 – 99
- [ ] 100 – 999
- [ ] 1 000 – 9 999
- [ ] 10 000 – 99 999
- [ ] 100 000 – 999 999
- [ ] 1 000 000 – 9 999 999
- [ ] ≥ 10 000 000
- [ ] Variable / depends on workload
- [ ] Custom: ☐

**Largest segment size (upper bound across all segments):**
- [ ] 1 – 9
- [ ] 10 – 99
- [ ] 100 – 999
- [ ] 1 000 – 9 999
- [ ] 10 000 – 99 999
- [ ] 100 000 – 999 999
- [ ] 1 000 000 – 9 999 999
- [ ] ≥ 10 000 000
- [ ] Variable / depends on workload
- [ ] Custom: ☐

**Typical (median or mean) segment size:** ☐

---

## 📊 Segment Size Distribution

This is the part that most directly drives our design, so please be as specific as you can.

**Shape of the distribution:**
- [ ] All segments (roughly) the same size
- [ ] Uniform within a range
- [ ] Power-law / Zipf (many small, few large)
- [ ] Heavy tail — mostly tiny segments plus a handful of very large ones
- [ ] Bimodal / clustered around a few sizes
- [ ] Unknown / data dependent
- [ ] Other: ☐

---

## 🚀 Execution & Integration

**Do you need CUDA graph capture support?**
(`DeviceSegmentedSort` currently synchronizes the stream mid-algorithm, so it **cannot** be stream-captured.)
- [ ] Yes — this is a blocker for me today
- [ ] Yes — nice to have, not blocking
- [ ] No

**Are you latency-bound or throughput-bound?**
- [ ] **Latency** — my problem often does not fill the GPU, and end-to-end call time is what matters
- [ ] **Throughput** — my problems are large enough to saturate the device
- [ ] Both, depending on the call site

**How often is the sort called?**
- [ ] Once / rarely
- [ ] Once per iteration of a hot loop — please state a rough call rate: ☐
- [ ] Many independent sorts in flight concurrently (multiple streams)

**Which GPU architectures do you care about?**
- [ ] Volta (SM 70)
- [ ] Turing (SM 75)
- [ ] Ampere (SM 80/86)
- [ ] Ada (SM 89)
- [ ] Hopper (SM 90)
- [ ] Blackwell (SM 100)
- [ ] Blackwell (SM 120)
- [ ] Other: ☐

---

## 🎯 Prioritization

Feel free to re-arrange based on your priorities:
- [ ] Less overprovisioning for small segments (finer size classes, thread/warp-level specializations)
- [ ] CUDA graph capture support / removing the host round-trip
- [ ] A latency-optimized path for small problems
- [ ] Support for the non-contiguous / batched layout
- [ ] A dedicated argsort (index-only) entry point
- [ ] Custom comparator support
- [ ] Custom / extended key types
- [ ] Faster large segments
- [ ] Other: ☐

```
Anything else we should know? Benchmarks, profiles, reproducers, or a link to your call site are all welcome.
```

````

---

_Related `DeviceSegmentedSort` issues:_

* 📄 [`[EPIC] Improve performance of DeviceSegmentedSort` #10577](https://github.com/NVIDIA/cccl/issues/10577)
*

_If you are interested in other segmented / batched algorithms, please see the corresponding issues:_

* 📄 [`DeviceSegmentedTopK`: User requests & feature prioritization](https://github.com/NVIDIA/cccl/issues/6391)

Start with the DeviceSegmentedSort rework described in issue #10577, then review the workload and feature specification form in this issue. Gather concrete use-case details, including layout, key types, segment sizes, execution needs, and priorities, and post them in a comment. Done means providing a completed or partially completed form that helps prioritize the redesign.

Written by the indexing model from the issue text.