56fd68e7dd
CCCL (CUDA C++ Core Libraries) provides: - CUB: device/block/warp-level GPU primitives (reduce, scan, sort, topk) - Thrust: high-level parallel algorithms (transform_reduce, sort, scan) - libcudacxx: CUDA C++ standard library (atomics, barriers, memory) - cudax: experimental features (memory resources, allocators) - Tuning policies: per-SM hardware-specific algorithm parameters Competition optimization vectors mapped to CCCL: - Output TPS (83% weight): warp_reduce, block_reduce, device_topk - Input TPS (14% weight): device_scan, block_load, prefetch - Cache TPS (3% weight): prefix caching strategy patterns - Memory (0.9 util): pooled/cached/buddy allocators Source: https://github.com/NVIDIA/cccl (shallow clone, HEAD only) License: Apache-2.0
28 lines
773 B
Plaintext
28 lines
773 B
Plaintext
#include <thrust/device_vector.h>
|
|
#include <thrust/iterator/permutation_iterator.h>
|
|
#include <thrust/reduce.h>
|
|
|
|
#include <iostream>
|
|
|
|
// this example fuses a gather operation with a reduction for
|
|
// greater efficiency than separate gather() and reduce() calls
|
|
|
|
int main()
|
|
{
|
|
// gather locations
|
|
thrust::device_vector<int> map = {3, 1, 0, 5};
|
|
|
|
// array to gather from
|
|
thrust::device_vector<int> source = {10, 20, 30, 40, 50, 60};
|
|
|
|
// fuse gather with reduction:
|
|
// sum = source[map[0]] + source[map[1]] + ...
|
|
int sum = thrust::reduce(thrust::make_permutation_iterator(source.begin(), map.begin()),
|
|
thrust::make_permutation_iterator(source.begin(), map.end()));
|
|
|
|
// print sum
|
|
std::cout << "sum is " << sum << '\n';
|
|
|
|
return 0;
|
|
}
|