Thrust

One address space means a device_vector IS a host vector; the algorithms are the <algorithm>/<numeric> ones. Execution-policy tags (thrust::host/device) are accepted and ignored.

Defined in code/simulator.h.

This page documents 47 symbols: 47 functions.

Functions

Signature	Description	Availability	Location	Example
inline Out copy(In f, In l, Out o)		sim only	simulator.h:2706	View
inline Out copy(simExecPolicy, In f, In l, Out o)		sim only	simulator.h:2742	View
inline std::ptrdiff_t count(It f, It l, const T& v)		sim only	simulator.h:2734	View
inline std::ptrdiff_t count_if(It f, It l, P p)		sim only	simulator.h:2735	View
inline device_ptr<T> device_pointer_cast(T* p)		sim only	simulator.h:2695
explicit device_ptr(T* q) : p(q)		sim only	simulator.h:2683
inline Out exclusive_scan(In f, In l, Out o)		sim only	simulator.h:2724
inline Out exclusive_scan(simExecPolicy, In f, In l, Out o)		sim only	simulator.h:2745
inline void fill(It f, It l, const T& v)		sim only	simulator.h:2707	View
std::vector<std::ptrdiff_t> idx(static_cast<size_t>(n))		sim only	simulator.h:2749	View
inline Out inclusive_scan(In f, In l, Out o)		sim only	simulator.h:2723	View
inline Out inclusive_scan(simExecPolicy, In f, In l, Out o)		sim only	simulator.h:2744	View
inline bool is_sorted(It f, It l)		sim only	simulator.h:2717	View
inline bool is_sorted(It f, It l, C c)		sim only	simulator.h:2718	View
std::vector<typename std::iterator_traits<KIt>::value_type> k2(static_cast<size_t>(n))		sim only	simulator.h:2753
inline It max_element(It f, It l)		sim only	simulator.h:2737
inline It min_element(It f, It l)		sim only	simulator.h:2736
bool operator!=(const device_ptr& o) const		sim only	simulator.h:2693
T& operator*() const		sim only	simulator.h:2684
device_ptr operator+(std::ptrdiff_t i) const		sim only	simulator.h:2688
device_ptr& operator++()		sim only	simulator.h:2691
device_ptr operator-(std::ptrdiff_t i) const		sim only	simulator.h:2689
std::ptrdiff_t operator-(const device_ptr& o) const		sim only	simulator.h:2690
T* operator->() const		sim only	simulator.h:2685
bool operator==(const device_ptr& o) const		sim only	simulator.h:2692
T& operator[](std::ptrdiff_t i) const		sim only	simulator.h:2686
inline T* raw_pointer_cast(const device_ptr<T>& p)		sim only	simulator.h:2696
inline T* raw_pointer_cast(T* p)		sim only	simulator.h:2697
inline T* raw_pointer_cast(std::vector<T>& v)		sim only	simulator.h:2698
inline T reduce(It f, It l, T init)		sim only	simulator.h:2719	View
inline typename std::iterator_traits<It>::value_type reduce(It f, It l)		sim only	simulator.h:2720	View
inline T reduce(simExecPolicy, It f, It l, T init)		sim only	simulator.h:2743	View
inline It remove(It f, It l, const T& v)		sim only	simulator.h:2700	View
inline It remove(simExecPolicy, It f, It l, const T& v)	Execution-policy overloads: the tag is accepted and dropped.	sim only	simulator.h:2739	View
inline It remove_if(It f, It l, P p)		sim only	simulator.h:2701
inline It remove_if(simExecPolicy, It f, It l, P p)		sim only	simulator.h:2740
inline void sequence(It f, It l, T init)		sim only	simulator.h:2708	View
inline void sequence(It f, It l)	2-argument form: the value type is deduced from the iterator (a defaulted template parameter cannot be, which is why both spellings are needed).	sim only	simulator.h:2713	View
inline void sort(It f, It l)		sim only	simulator.h:2702	View
inline void sort(It f, It l, C c)		sim only	simulator.h:2703	View
inline void sort(simExecPolicy, It f, It l)		sim only	simulator.h:2741	View
inline void sort_by_key(KIt kf, KIt kl, VIt vf)	sort_by_key: co-sort a values range by its keys.	sim only	simulator.h:2747	View
inline void stable_sort(It f, It l)		sim only	simulator.h:2704
inline void stable_sort(It f, It l, C c)		sim only	simulator.h:2705
inline Out transform(In f, In l, Out o, Op op)		sim only	simulator.h:2729
inline Out transform(In1 f, In1 l, In2 g, Out o, Op op)		sim only	simulator.h:2732
std::vector<typename std::iterator_traits<VIt>::value_type> v2(static_cast<size_t>(n))		sim only	simulator.h:2754

Examples

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

copy (test_web_cs_radixsort_thrust.cu:119)

        float time = 0;
        check(cudaEventElapsedTime(&time, start_event, stop_event), "cudaEventElapsedTime");
        totalTime += time;
    }
    totalTime /= (1.0e3f * numIterations);
    std::printf("radixSortThrust, Throughput = %.4f MElements/s, Time = %.5f s, Size = %u elements\n",
                1.0e-6f * numElements / totalTime, totalTime, numElements);
    check(cudaGetLastError(), "after radixsort");

    // Get results back to host for correctness checking
    thrust::copy(d_keys.begin(), d_keys.end(), h_keysSorted.begin());
    if (!keysOnly) thrust::copy(d_values.begin(), d_values.end(), h_values.begin());
    check(cudaGetLastError(), "copying results to host memory");

    check(cudaEventDestroy(start_event), "cudaEventDestroy(start)");
    check(cudaEventDestroy(stop_event), "cudaEventDestroy(stop)");

    // Check 1 (the sample's): the keys come back in order.
    bool bTestResult = std::is_sorted(h_keysSorted.begin(), h_keysSorted.end());
    if (!bTestResult) std::printf("  output is not sorted\n");

    // Check 2 [SIM]: sorted keys are a permutation of the input.

copy (test_web_cs_radixsort_thrust.cu:119)

        float time = 0;
        check(cudaEventElapsedTime(&time, start_event, stop_event), "cudaEventElapsedTime");
        totalTime += time;
    }
    totalTime /= (1.0e3f * numIterations);
    std::printf("radixSortThrust, Throughput = %.4f MElements/s, Time = %.5f s, Size = %u elements\n",
                1.0e-6f * numElements / totalTime, totalTime, numElements);
    check(cudaGetLastError(), "after radixsort");

    // Get results back to host for correctness checking
    thrust::copy(d_keys.begin(), d_keys.end(), h_keysSorted.begin());
    if (!keysOnly) thrust::copy(d_values.begin(), d_values.end(), h_values.begin());
    check(cudaGetLastError(), "copying results to host memory");

    check(cudaEventDestroy(start_event), "cudaEventDestroy(start)");
    check(cudaEventDestroy(stop_event), "cudaEventDestroy(stop)");

    // Check 1 (the sample's): the keys come back in order.
    bool bTestResult = std::is_sorted(h_keysSorted.begin(), h_keysSorted.end());
    if (!bTestResult) std::printf("  output is not sorted\n");

    // Check 2 [SIM]: sorted keys are a permutation of the input.

count (test_web_ndb_cpp11.cu:111)

  cudaMemcpy(d_text, corpus, len, cudaMemcpyHostToDevice);

  int count = 0;
  int *d_count;
  cudaMalloc(&d_count, sizeof(int));
  cudaMemset(d_count, 0, sizeof(int));

  LAUNCH(xyzw_frequency, 8, 256, d_count, d_text, len);   // [SIM]
  cudaMemcpy(&count, d_count, sizeof(int), cudaMemcpyDeviceToHost);

  std::cout << "counted " << count << " instances of 'x', 'y', 'z', or 'w' "
            << "(expected " << expected << ")" << std::endl;
  const bool ok = (count == expected);
  std::printf("C++11 grid-stride count_if: %s\n", ok ? "PASSED" : "FAILED");

  cudaFree(d_count);
  cudaFree(d_text);
  return ok ? 0 : 1;
}

count_if (test_web_ndb_cpp11.cu:52)

    __host__ __device__ bool operator!=(const iter& o) const { return cur < o.cur; }
  };
  __host__ __device__ iter begin() const { return { begin_, step_ }; }
  __host__ __device__ iter end()   const { return { end_,   step_ }; }
};
/////////////////////////////////////////////////////////////////

// verbatim device function template from the source repo
template <typename T, typename Predicate>
__device__
void count_if(int *count, T *data, int n, Predicate p)
{
  for (auto i : grid_stride_range(0, n)) {
    if (p(data[i])) atomicAdd(count, 1);
  }
}

// verbatim kernel: lambda that searches for x, y, z or w (range-based for +
// initializer_list inside the functor; auto so the array/functor types are implicit)
__global__
void xyzw_frequency(int *count, char *text, int n)
{

fill (test_web_cs_transpose.cu:289)

        case 1: kernel = &copySharedMem;            kernelName = "shared memory copy"; break;
        case 2: kernel = &transposeNaive;           kernelName = "naive             "; break;
        case 3: kernel = &transposeCoalesced;       kernelName = "coalesced         "; break;
        case 4: kernel = &transposeNoBankConflicts; kernelName = "optimized         "; break;
        case 5: kernel = &transposeCoarseGrained;   kernelName = "coarse-grained    "; break;
        case 6: kernel = &transposeFineGrained;     kernelName = "fine-grained      "; break;
        case 7: kernel = &transposeDiagonal;        kernelName = "diagonal          "; break;
        }

        // reset output between kernels (as the sample does)
        std::fill(h_odata.begin(), h_odata.end(), 0.0f);
        cudaMemcpy(d_odata, h_odata.data(), mem, cudaMemcpyHostToDevice);

        LAUNCH(kernel, grid, threads, d_odata, d_idata, size_x, size_y);   // [SIM]

        cudaMemcpy(h_odata.data(), d_odata, mem, cudaMemcpyDeviceToHost);

        // pick the reference: copies -> identity; coarse/fine -> partial (skip);
        // everything else -> full transpose.
        const bool isCopy    = (kernel == &copy || kernel == &copySharedMem);
        const bool isPartial = (kernel == &transposeCoarseGrained || kernel == &transposeFineGrained);
        const float *gold    = isCopy ? h_idata.data() : gold_T.data();

idx (test_llama.cu:190)

        float e = expf(dot * scale - m);
        denom += e;
        num   += e * V[(static_cast<size_t>(j) * N_KV + kvhead) * HEAD_DIM + d];
    }
    out[idx] = num / denom;
}

// Element-wise residual add: out = a + b. One thread per element.
__global__ void addKernel(const float* a, const float* b, float* out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) out[idx] = a[idx] + b[idx];
}

// SwiGLU gate: out = SiLU(gate) ⊙ up, with SiLU(x) = x * sigmoid(x). One
// thread per element of the (T, 1408) hidden tensor.
__global__ void swigluKernel(const float* gate, const float* up, float* out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float g = gate[idx];
        float silu = g / (1.0f + expf(-g));
        out[idx] = silu * up[idx];
    }

inclusive_scan (test_web_bf_thrust_scan.cu:72)

            g_ok = false;
        }
}

// --- thrust_scan_vector.cu: scan a thrust::device_vector in place ------------
static void scan_vector() {
    thrust::device_vector<int> x(N, 0);
    thrust::device_vector<int> y(N, 0);
    for (int i = 0; i < N; ++i) x[i] = i + 1;

    thrust::inclusive_scan(x.begin(), x.end(), y.begin());

    std::vector<int> got(N);
    std::printf("device_vector: ");
    for (int i = 0; i < N; ++i) { got[i] = int(y[i]); std::printf("%d ", got[i]); }
    std::printf("\n");
    verify(got, "device_vector");
}

// --- thrust_scan_pointer.cu: scan raw device pointers, thrust::device policy -
static void scan_pointer() {
    int *x = nullptr, *y = nullptr;

inclusive_scan (test_web_bf_thrust_scan.cu:72)

            g_ok = false;
        }
}

// --- thrust_scan_vector.cu: scan a thrust::device_vector in place ------------
static void scan_vector() {
    thrust::device_vector<int> x(N, 0);
    thrust::device_vector<int> y(N, 0);
    for (int i = 0; i < N; ++i) x[i] = i + 1;

    thrust::inclusive_scan(x.begin(), x.end(), y.begin());

    std::vector<int> got(N);
    std::printf("device_vector: ");
    for (int i = 0; i < N; ++i) { got[i] = int(y[i]); std::printf("%d ", got[i]); }
    std::printf("\n");
    verify(got, "device_vector");
}

// --- thrust_scan_pointer.cu: scan raw device pointers, thrust::device policy -
static void scan_pointer() {
    int *x = nullptr, *y = nullptr;

is_sorted (test_web_cs_radixsort_thrust.cu:127)

    // Get results back to host for correctness checking
    thrust::copy(d_keys.begin(), d_keys.end(), h_keysSorted.begin());
    if (!keysOnly) thrust::copy(d_values.begin(), d_values.end(), h_values.begin());
    check(cudaGetLastError(), "copying results to host memory");

    check(cudaEventDestroy(start_event), "cudaEventDestroy(start)");
    check(cudaEventDestroy(stop_event), "cudaEventDestroy(stop)");

    // Check 1 (the sample's): the keys come back in order.
    bool bTestResult = std::is_sorted(h_keysSorted.begin(), h_keysSorted.end());
    if (!bTestResult) std::printf("  output is not sorted\n");

    // Check 2 [SIM]: sorted keys are a permutation of the input.
    {
        std::vector<T> a(keysOriginal), b(h_keysSorted.begin(), h_keysSorted.end());
        std::sort(a.begin(), a.end());
        if (a != b) { std::printf("  output is not a permutation of the input\n"); bTestResult = false; }
    }

    // Check 3 [SIM]: each payload still points at its own key.
    if (!keysOnly) {

is_sorted (test_web_cs_radixsort_thrust.cu:127)

    // Get results back to host for correctness checking
    thrust::copy(d_keys.begin(), d_keys.end(), h_keysSorted.begin());
    if (!keysOnly) thrust::copy(d_values.begin(), d_values.end(), h_values.begin());
    check(cudaGetLastError(), "copying results to host memory");

    check(cudaEventDestroy(start_event), "cudaEventDestroy(start)");
    check(cudaEventDestroy(stop_event), "cudaEventDestroy(stop)");

    // Check 1 (the sample's): the keys come back in order.
    bool bTestResult = std::is_sorted(h_keysSorted.begin(), h_keysSorted.end());
    if (!bTestResult) std::printf("  output is not sorted\n");

    // Check 2 [SIM]: sorted keys are a permutation of the input.
    {
        std::vector<T> a(keysOriginal), b(h_keysSorted.begin(), h_keysSorted.end());
        std::sort(a.begin(), a.end());
        if (a != b) { std::printf("  output is not a permutation of the input\n"); bTestResult = false; }
    }

    // Check 3 [SIM]: each payload still points at its own key.
    if (!keysOnly) {

reduce (test_coop_groups.cu:73)

    auto tile  = cg::tiled_partition<32>(block);
    int v = in[gid];
    for (unsigned d = tile.size() / 2; d > 0; d >>= 1) v += tile.shfl_down(v, d);
    if (tile.thread_rank() == 0) outWarp[gid / 32u] = v;
}

// cg::reduce broadcasts the warp sum to every lane.
__global__ void kCgReduce(const int* in, int* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    out[gid] = cg::reduce(tile, in[gid], cg::plus<int>());
}

// tile.thread_rank() and meta_group_rank() identify a lane and its warp.
__global__ void kTileInfo(int* rank, int* warp) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    rank[gid] = static_cast<int>(tile.thread_rank());
    warp[gid] = static_cast<int>(tile.meta_group_rank());
}

static int* upload(const std::vector<int>& h) {

reduce (test_coop_groups.cu:73)

    auto tile  = cg::tiled_partition<32>(block);
    int v = in[gid];
    for (unsigned d = tile.size() / 2; d > 0; d >>= 1) v += tile.shfl_down(v, d);
    if (tile.thread_rank() == 0) outWarp[gid / 32u] = v;
}

// cg::reduce broadcasts the warp sum to every lane.
__global__ void kCgReduce(const int* in, int* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    out[gid] = cg::reduce(tile, in[gid], cg::plus<int>());
}

// tile.thread_rank() and meta_group_rank() identify a lane and its warp.
__global__ void kTileInfo(int* rank, int* warp) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    rank[gid] = static_cast<int>(tile.thread_rank());
    warp[gid] = static_cast<int>(tile.meta_group_rank());
}

static int* upload(const std::vector<int>& h) {

reduce (test_coop_groups.cu:73)

    auto tile  = cg::tiled_partition<32>(block);
    int v = in[gid];
    for (unsigned d = tile.size() / 2; d > 0; d >>= 1) v += tile.shfl_down(v, d);
    if (tile.thread_rank() == 0) outWarp[gid / 32u] = v;
}

// cg::reduce broadcasts the warp sum to every lane.
__global__ void kCgReduce(const int* in, int* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    out[gid] = cg::reduce(tile, in[gid], cg::plus<int>());
}

// tile.thread_rank() and meta_group_rank() identify a lane and its warp.
__global__ void kTileInfo(int* rank, int* warp) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    auto tile = cg::tiled_partition<32>(cg::this_thread_block());
    rank[gid] = static_cast<int>(tile.thread_rank());
    warp[gid] = static_cast<int>(tile.meta_group_rank());
}

static int* upload(const std::vector<int>& h) {

remove (test_web_olcf_hw9_task2.cu:168)

    LAUNCH_COOP(my_remove_if<mytype>, grid, block, c_idata, remove_val, d_odata, d_idxs, ds);
    cudaDeviceSynchronize();
    cudaMemcpy(h_data.data(), d_odata, tsize, cudaMemcpyDeviceToHost);
    for (int i = 0; i < val; i++)
        if (h_data[i] != i) {
            std::printf("mismatch 2 at %d, was: %d, should be: %d\n", i, h_data[i], i);
            ok = false; break;
        }

    // Cross-check against thrust::remove.
    thrust::remove(t_data.begin(), t_data.end(), -1);
    thrust::host_vector<mytype> th_data = t_data;
    for (int i = 0; i < val; i++)
        if (h_data[i] != th_data[i]) {
            std::printf("mismatch 3 at %d, was: %d, should be: %d\n", i, th_data[i], h_data[i]);
            ok = false; break;
        }

    cudaFree(d_idata);
    cudaFree(d_odata);
    cudaFree(d_idxs);

remove (test_web_olcf_hw9_task2.cu:168)

    LAUNCH_COOP(my_remove_if<mytype>, grid, block, c_idata, remove_val, d_odata, d_idxs, ds);
    cudaDeviceSynchronize();
    cudaMemcpy(h_data.data(), d_odata, tsize, cudaMemcpyDeviceToHost);
    for (int i = 0; i < val; i++)
        if (h_data[i] != i) {
            std::printf("mismatch 2 at %d, was: %d, should be: %d\n", i, h_data[i], i);
            ok = false; break;
        }

    // Cross-check against thrust::remove.
    thrust::remove(t_data.begin(), t_data.end(), -1);
    thrust::host_vector<mytype> th_data = t_data;
    for (int i = 0; i < val; i++)
        if (h_data[i] != th_data[i]) {
            std::printf("mismatch 3 at %d, was: %d, should be: %d\n", i, th_data[i], h_data[i]);
            ok = false; break;
        }

    cudaFree(d_idata);
    cudaFree(d_odata);
    cudaFree(d_idxs);

sequence (test_web_cs_radixsort_thrust.cu:84)

    std::mt19937 rng(20220607u);
    if (floatKeys) {
        std::uniform_real_distribution<float> u01(0.0f, 1.0f);
        for (unsigned int i = 0; i < numElements; i++) h_keys[i] = T(u01(rng));
    } else {
        std::uniform_int_distribution<unsigned int> u(0, UINT_MAX);
        for (unsigned int i = 0; i < numElements; i++) h_keys[i] = T(u(rng));
    }
    // [SIM] the sample writes thrust::sequence(first, last); the sim only has the
    // three-argument form, so the start value 0u is spelled out (legal in both).
    if (!keysOnly) thrust::sequence(h_values.begin(), h_values.end(), 0u);

    // Keep a copy of the unsorted keys so the permutation can be checked.
    const std::vector<T> keysOriginal(h_keys.begin(), h_keys.end());

    thrust::device_vector<T>            d_keys;
    thrust::device_vector<unsigned int> d_values;

    cudaEvent_t start_event, stop_event;
    check(cudaEventCreate(&start_event), "cudaEventCreate(start)");
    check(cudaEventCreate(&stop_event), "cudaEventCreate(stop)");

sequence (test_web_cs_radixsort_thrust.cu:84)

    std::mt19937 rng(20220607u);
    if (floatKeys) {
        std::uniform_real_distribution<float> u01(0.0f, 1.0f);
        for (unsigned int i = 0; i < numElements; i++) h_keys[i] = T(u01(rng));
    } else {
        std::uniform_int_distribution<unsigned int> u(0, UINT_MAX);
        for (unsigned int i = 0; i < numElements; i++) h_keys[i] = T(u(rng));
    }
    // [SIM] the sample writes thrust::sequence(first, last); the sim only has the
    // three-argument form, so the start value 0u is spelled out (legal in both).
    if (!keysOnly) thrust::sequence(h_values.begin(), h_values.end(), 0u);

    // Keep a copy of the unsorted keys so the permutation can be checked.
    const std::vector<T> keysOriginal(h_keys.begin(), h_keys.end());

    thrust::device_vector<T>            d_keys;
    thrust::device_vector<unsigned int> d_values;

    cudaEvent_t start_event, stop_event;
    check(cudaEventCreate(&start_event), "cudaEventCreate(start)");
    check(cudaEventCreate(&stop_event), "cudaEventCreate(stop)");

sort (test_web_cs_bitonic.cu:140)

    cudaMemcpy(d_SrcVal, h_val.data(), bytes, cudaMemcpyHostToDevice);

    // Single block, SHARED_SIZE_LIMIT/2 threads (one thread per comparator)
    LAUNCH(bitonicSortShared, 1, SHARED_SIZE_LIMIT / 2,
           d_DstKey, d_DstVal, d_SrcKey, d_SrcVal, arrayLength, dir);   // [SIM]

    cudaMemcpy(h_res.data(), d_DstKey, bytes, cudaMemcpyDeviceToHost);

    // CPU reference: std::sort of the same keys
    std::vector<uint> ref(h_key);
    std::sort(ref.begin(), ref.end());

    bool ok = true;
    for (uint i = 0; i < arrayLength; i++)
        if (h_res[i] != ref[i]) { ok = false; break; }

    std::printf("bitonic sort of %u keys (dir=%u) vs std::sort: %s\n",
                arrayLength, dir, ok ? "PASSED" : "FAILED");

    cudaFree(d_SrcKey); cudaFree(d_SrcVal); cudaFree(d_DstKey); cudaFree(d_DstVal);
    return ok ? 0 : 1;
}

sort (test_web_cs_bitonic.cu:140)

    cudaMemcpy(d_SrcVal, h_val.data(), bytes, cudaMemcpyHostToDevice);

    // Single block, SHARED_SIZE_LIMIT/2 threads (one thread per comparator)
    LAUNCH(bitonicSortShared, 1, SHARED_SIZE_LIMIT / 2,
           d_DstKey, d_DstVal, d_SrcKey, d_SrcVal, arrayLength, dir);   // [SIM]

    cudaMemcpy(h_res.data(), d_DstKey, bytes, cudaMemcpyDeviceToHost);

    // CPU reference: std::sort of the same keys
    std::vector<uint> ref(h_key);
    std::sort(ref.begin(), ref.end());

    bool ok = true;
    for (uint i = 0; i < arrayLength; i++)
        if (h_res[i] != ref[i]) { ok = false; break; }

    std::printf("bitonic sort of %u keys (dir=%u) vs std::sort: %s\n",
                arrayLength, dir, ok ? "PASSED" : "FAILED");

    cudaFree(d_SrcKey); cudaFree(d_SrcVal); cudaFree(d_DstKey); cudaFree(d_DstVal);
    return ok ? 0 : 1;
}

sort (test_web_cs_bitonic.cu:140)

    cudaMemcpy(d_SrcVal, h_val.data(), bytes, cudaMemcpyHostToDevice);

    // Single block, SHARED_SIZE_LIMIT/2 threads (one thread per comparator)
    LAUNCH(bitonicSortShared, 1, SHARED_SIZE_LIMIT / 2,
           d_DstKey, d_DstVal, d_SrcKey, d_SrcVal, arrayLength, dir);   // [SIM]

    cudaMemcpy(h_res.data(), d_DstKey, bytes, cudaMemcpyDeviceToHost);

    // CPU reference: std::sort of the same keys
    std::vector<uint> ref(h_key);
    std::sort(ref.begin(), ref.end());

    bool ok = true;
    for (uint i = 0; i < arrayLength; i++)
        if (h_res[i] != ref[i]) { ok = false; break; }

    std::printf("bitonic sort of %u keys (dir=%u) vs std::sort: %s\n",
                arrayLength, dir, ok ? "PASSED" : "FAILED");

    cudaFree(d_SrcKey); cudaFree(d_SrcVal); cudaFree(d_DstKey); cudaFree(d_DstVal);
    return ok ? 0 : 1;
}

sort_by_key (test_web_cs_radixsort_thrust.cu:104)

    check(cudaEventCreate(&stop_event), "cudaEventCreate(stop)");

    float totalTime = 0;
    for (unsigned int i = 0; i < numIterations; i++) {
        d_keys = h_keys;                           // reset data before sort
        if (!keysOnly) d_values = h_values;

        check(cudaEventRecord(start_event, 0), "cudaEventRecord(start)");

        if (keysOnly) thrust::sort(d_keys.begin(), d_keys.end());
        else          thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());

        check(cudaEventRecord(stop_event, 0), "cudaEventRecord(stop)");
        check(cudaEventSynchronize(stop_event), "cudaEventSynchronize");

        float time = 0;
        check(cudaEventElapsedTime(&time, start_event, stop_event), "cudaEventElapsedTime");
        totalTime += time;
    }
    totalTime /= (1.0e3f * numIterations);
    std::printf("radixSortThrust, Throughput = %.4f MElements/s, Time = %.5f s, Size = %u elements\n",
                1.0e-6f * numElements / totalTime, totalTime, numElements);