CUDA Argsort Optimization Guide

// Slow: copy-based approach
void argsort_naive(float* vals, int* idx, int n) {
    for (int i = 0; i < n; i++) idx[i] = i;
    thrust::sort(idx, idx + n, [vals](int a, int b) {
        return vals[a] > vals[b];  // Indirect access - slow!
    });
}

#include <cub/cub.cuh>

void argsort_cub(float* vals, int* indices, int n) {
    thrust::sequence(thrust::device, indices, indices + n);

    // Allocate temporary storage
    size_t temp_bytes = 0;
    cub::DeviceRadixSort::SortPairsDescending(
        nullptr, temp_bytes, vals, vals_out, indices, indices_out, n);

    void* d_temp;
    cudaMalloc(&d_temp, temp_bytes);

    cub::DeviceRadixSort::SortPairsDescending(
        d_temp, temp_bytes, vals, vals_out, indices, indices_out, n);
}