Thrust & CUB

Fancy iterators: counting_iterator, transform_iterator, zip_iterator, permutation_iterator, constant_iterator — compose operations without temporary arrays.

31.1 Thrust — STL-like parallel algorithms

Fancy iterators: counting_iterator, transform_iterator, zip_iterator, permutation_iterator, constant_iterator — compose operations without temporary arrays.

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/sort.h>
#include <thrust/reduce.h>
#include <thrust/transform.h>
#include <thrust/sequence.h>
#include <thrust/count.h>
#include <thrust/functional.h>

thrust::host_vector<float> h(1 << 20);
for (auto& v : h) v = rand() / (float)RAND_MAX;
thrust::device_vector<float> d = h;                          // copies to the GPU

thrust::sort(d.begin(), d.end());                            // GPU radix sort
float sum = thrust::reduce(d.begin(), d.end(), 0.0f, thrust::plus<float>());
thrust::transform(d.begin(), d.end(), d.begin(), [] __device__ (float x) { return x * x; });
int big = thrust::count_if(d.begin(), d.end(), [] __device__ (float x) { return x > 0.5f; });
thrust::inclusive_scan(d.begin(), d.end(), d.begin());
thrust::sequence(d.begin(), d.end());                        // 0, 1, 2, ...
thrust::sort_by_key(keys.begin(), keys.end(), values.begin());
thrust::reduce_by_key(...); thrust::copy_if(...); thrust::unique(...); thrust::gather(...);
float* raw = thrust::raw_pointer_cast(d.data());             // pass to custom kernels
thrust::sort(thrust::cuda::par.on(stream), d.begin(), d.end());   // execution
…

31.2 CUB — high-performance building blocks

Levels: device-wide (DeviceReduce, DeviceScan, DeviceRadixSort, DeviceSelect, DeviceHistogram, DeviceRunLengthEncode, DeviceSegmentedReduce), block-level (BlockReduce, BlockScan, BlockLoad/Store, BlockRadixSort, BlockExchange), warp-level (WarpReduce, WarpScan).

#include <cub/cub.cuh>

// Device-wide reduction (two-phase: query temp storage size, then run)
void* d_temp = nullptr; size_t tempBytes = 0;
cub::DeviceReduce::Sum(d_temp, tempBytes, d_in, d_out, n);
cudaMalloc(&d_temp, tempBytes);
cub::DeviceReduce::Sum(d_temp, tempBytes, d_in, d_out, n);

// Block-level primitive inside your own kernel
__global__ void blockSumKernel(const int* in, int* out) {
    using BlockReduce = cub::BlockReduce<int, 256>;
    __shared__ typename BlockReduce::TempStorage temp;
    int v = in[blockIdx.x * 256 + threadIdx.x];
    int total = BlockReduce(temp).Sum(v);
    if (threadIdx.x == 0) out[blockIdx.x] = total;
}