cuda thrust library

**CUDA Thrust Library** is **the high-level C++ template library providing STL-like parallel algorithms and data structures for CUDA** — offering 40+ algorithms including sort (100-300 GB/s), reduce (500-1000 GB/s), scan (400-800 GB/s), transform, and unique that automatically handle memory management, kernel launches, and optimization, achieving 60-90% of hand-optimized CUDA performance while reducing development time by 5-10× through expressive syntax like thrust::reduce(data.begin(), data.end()) that replaces 50-100 lines of custom CUDA code, making Thrust essential for rapid prototyping and production deployment where developer productivity matters and the 10-40% performance gap versus hand-tuned kernels is acceptable trade-off for 90% reduction in code complexity and maintenance burden. **Thrust Architecture:** - **Execution Policies**: thrust::device (GPU), thrust::host (CPU), thrust::omp (OpenMP); explicit control over execution location - **Containers**: thrust::device_vector, thrust::host_vector; automatic memory management; RAII semantics; seamless CPU-GPU transfers - **Iterators**: random access, constant, counting, transform, zip iterators; composable; enable complex operations without temporary storage - **Algorithms**: 40+ parallel algorithms; sort, reduce, scan, transform, copy, unique, partition; optimized implementations; 60-90% of hand-tuned performance **Core Algorithms:** - **Sort**: thrust::sort(), thrust::stable_sort(); radix sort for integers, merge sort for general; 100-300 GB/s; 60-80% of CUB performance - **Reduce**: thrust::reduce(), thrust::reduce_by_key(); sum, max, min, custom operators; 500-1000 GB/s; 70-90% of hand-tuned - **Scan**: thrust::inclusive_scan(), thrust::exclusive_scan(); prefix sum; 400-800 GB/s; 60-80% of hand-tuned - **Transform**: thrust::transform(); element-wise operations; 1-2 TB/s; near-optimal for memory-bound operations **Device Vectors:** - **Allocation**: thrust::device_vector d_vec(N); automatic cudaMalloc; RAII cleanup; exception-safe - **Access**: d_vec[i] for individual elements (slow); d_vec.data() for raw pointer; thrust::copy for bulk transfer - **Resize**: d_vec.resize(new_size); automatic reallocation; preserves existing data; amortized O(1) for growth - **Performance**: same as manual cudaMalloc; no overhead; automatic memory management eliminates leaks **Iterators:** - **Counting Iterator**: thrust::counting_iterator(0); generates sequence 0, 1, 2, ...; no storage; useful for indices - **Transform Iterator**: thrust::make_transform_iterator(iter, func); applies function on-the-fly; no temporary storage; 2-5× memory savings - **Zip Iterator**: thrust::make_zip_iterator(thrust::make_tuple(iter1, iter2)); combines multiple sequences; enables multi-array operations - **Constant Iterator**: thrust::constant_iterator(value); infinite sequence of same value; no storage; useful for fills **Functional Programming:** - **Functors**: thrust::plus(), thrust::multiplies(), thrust::maximum(); predefined operators; custom functors supported - **Lambda Expressions**: C++11 lambdas work with Thrust; [=] __device__ (int x) { return x * x; }; concise custom operations - **Composition**: combine iterators and functors; complex operations without temporaries; 2-10× memory savings - **Type Safety**: compile-time type checking; catches errors early; safer than raw CUDA **Performance Characteristics:** - **Sort**: 100-300 GB/s; 60-80% of CUB; 80-95% of hand-tuned; acceptable for most applications - **Reduce**: 500-1000 GB/s; 70-90% of hand-tuned; near-optimal for large arrays (>1M elements) - **Scan**: 400-800 GB/s; 60-80% of hand-tuned; good for large arrays; overhead for small arrays (<10K elements) - **Transform**: 1-2 TB/s; 90-100% of hand-tuned; memory-bound operations achieve peak bandwidth **Memory Management:** - **Automatic**: device_vector handles allocation, deallocation; no manual cudaMalloc/cudaFree; eliminates memory leaks - **RAII**: resource acquisition is initialization; exception-safe; automatic cleanup on scope exit - **Transfers**: thrust::copy(h_vec.begin(), h_vec.end(), d_vec.begin()); automatic cudaMemcpy; type-safe - **Pinned Memory**: thrust::system::cuda::experimental::pinned_allocator; 2-10× faster transfers; limited resource **Common Patterns:** - **Reduction**: float sum = thrust::reduce(d_vec.begin(), d_vec.end(), 0.0f, thrust::plus()); 500-1000 GB/s; 5-10 lines vs 50-100 for custom - **Transform**: thrust::transform(d_in.begin(), d_in.end(), d_out.begin(), thrust::negate()); 1-2 TB/s; 1 line vs 20-30 for custom - **Sort**: thrust::sort(d_vec.begin(), d_vec.end()); 100-300 GB/s; 1 line vs 100-200 for custom radix sort - **Scan**: thrust::exclusive_scan(d_in.begin(), d_in.end(), d_out.begin()); 400-800 GB/s; 1 line vs 50-100 for custom **Advanced Algorithms:** - **Reduce By Key**: thrust::reduce_by_key(keys.begin(), keys.end(), values.begin(), ...); groups by key; 300-600 GB/s; useful for histograms - **Unique**: thrust::unique(d_vec.begin(), d_vec.end()); removes duplicates; 200-400 GB/s; requires sorted input - **Partition**: thrust::partition(d_vec.begin(), d_vec.end(), predicate); splits by condition; 300-600 GB/s; stable variant available - **Merge**: thrust::merge(d_vec1.begin(), d_vec1.end(), d_vec2.begin(), d_vec2.end(), d_out.begin()); 200-400 GB/s **Custom Operations:** - **Custom Functors**: struct my_op { __device__ float operator()(float x) { return x * x; } }; use with transform, reduce - **Binary Operations**: struct my_binary_op { __device__ float operator()(float a, float b) { return a + b * b; } }; use with reduce, transform - **Predicates**: struct is_positive { __device__ bool operator()(float x) { return x > 0; } }; use with partition, remove_if - **Performance**: custom functors inline; no overhead vs hand-written; full optimization **Integration with CUDA:** - **Raw Pointers**: thrust::device_ptr d_ptr = thrust::device_pointer_cast(raw_ptr); wraps raw CUDA pointers - **Custom Kernels**: mix Thrust algorithms with custom kernels; use d_vec.data() for raw pointer; seamless integration - **Streams**: thrust::cuda::par.on(stream); execute Thrust algorithms on specific stream; enables concurrency - **Memory**: Thrust and CUDA share same memory; no copies; interoperable **Comparison with Alternatives:** - **vs CUB**: CUB 10-40% faster; lower-level; more complex API; Thrust easier to use; acceptable performance gap - **vs cuBLAS/cuDNN**: specialized libraries faster for their domains; Thrust more general; use specialized libraries when available - **vs Hand-Tuned**: hand-tuned 10-40% faster; 5-10× more code; Thrust faster development; use hand-tuned for critical kernels only - **vs STL**: Thrust 10-100× faster than CPU STL; same API; easy migration; parallel by default **Development Productivity:** - **Code Reduction**: 5-10× less code than custom CUDA; 1-10 lines vs 50-200 for complex algorithms - **Maintainability**: high-level abstractions; easier to understand; fewer bugs; 50-80% reduction in debugging time - **Portability**: same code runs on CPU (thrust::host) or GPU (thrust::device); easy testing; gradual migration - **Learning Curve**: familiar STL-like API; easier than raw CUDA; 2-4 weeks to proficiency vs 2-3 months for CUDA **Performance Optimization:** - **Execution Policy**: thrust::device for GPU, thrust::host for CPU; explicit control; choose based on data size - **Fused Operations**: combine operations to reduce kernel launches; thrust::transform_reduce(); 20-40% faster than separate - **Iterator Composition**: use transform iterators to avoid temporaries; 2-10× memory savings; 20-50% faster - **Custom Allocators**: thrust::device_malloc_allocator, custom allocators; control memory management; 10-30% improvement **Profiling Thrust:** - **Nsight Systems**: shows Thrust kernel launches; identifies bottlenecks; visualizes execution timeline - **Nsight Compute**: profiles individual Thrust kernels; memory, compute metrics; guides optimization - **Overhead**: Thrust overhead <5% for large arrays (>1M elements); 10-30% for small arrays (<10K elements) - **Optimization**: profile to identify slow operations; replace with custom kernels if needed; 90% of code uses Thrust, 10% custom **Best Practices:** - **Use Thrust First**: start with Thrust; optimize only if profiling shows bottleneck; 90% of code doesn't need hand-tuning - **Fuse Operations**: combine operations to reduce launches; transform_reduce, transform_scan; 20-40% faster - **Iterator Composition**: avoid temporaries with transform iterators; 2-10× memory savings; 20-50% faster - **Profile**: measure actual performance; compare with requirements; hand-tune only critical 10% - **Mix with CUDA**: use Thrust for general algorithms, custom kernels for specialized operations; best of both worlds **Performance Targets:** - **Sort**: 100-300 GB/s; 60-80% of hand-tuned; acceptable for most applications; 1 line of code - **Reduce**: 500-1000 GB/s; 70-90% of hand-tuned; near-optimal for large arrays; 1 line of code - **Scan**: 400-800 GB/s; 60-80% of hand-tuned; good for large arrays; 1 line of code - **Transform**: 1-2 TB/s; 90-100% of hand-tuned; memory-bound operations achieve peak; 1 line of code **Real-World Usage:** - **Data Processing**: sort, filter, transform pipelines; 5-10× faster development; 60-90% of hand-tuned performance - **Scientific Computing**: reductions, scans, transformations; 80-95% of hand-tuned; 90% less code - **Machine Learning**: data preprocessing, feature extraction; 70-90% of hand-tuned; rapid prototyping - **Graph Algorithms**: sorting edges, reducing vertices; 60-80% of hand-tuned; 5-10× less code CUDA Thrust Library represents **the productivity revolution in GPU programming** — by providing STL-like parallel algorithms that achieve 60-90% of hand-optimized performance while reducing code by 5-10× and development time by similar factors, Thrust makes GPU programming accessible to developers without deep CUDA expertise and enables rapid prototyping and production deployment where the 10-40% performance gap versus hand-tuned kernels is acceptable trade-off for 90% reduction in code complexity, making Thrust the default choice for GPU algorithms where developer time is more valuable than the last 10-40% of performance.