University:
Boston University
Course:
MET CS 201 | Introduction to Programming
Academic year:

2023
Views:

70

Pages:

16
Author:

Damian Osborn

Thrust - Fancy iterators ‣ Behave like “normal’” iterators - Also they can be seen as pointers ‣ Examples ‐ constant_iterator ‐ counting_iterator ‐ transform_iterator ‐ permutation_iterator ‐ zip_iterator Thrust - Fancy iterators ‣ constant_iterator - Mimics an infinite array with constant values // create iterators constant_iterator begin(10); constant_iterator end = begin + 3; begin[0] // returns 10 begin[1] // returns 10 begin[100] // returns 10 // sum of (begin, end) reduce(begin, end); // returns 30 (i.e. 3 * 10) Thrust - Fancy iterators ‣ counting_iterator - Mimics an infinite array with sequential values // create iterators counting_iterator begin(10); counting_iterator end = begin + 3; begin[0] // returns 10 begin[1] // returns 11 begin[100] // returns 110 // sum of (begin, end) reduce(begin, end); // returns 33 (i.e. 10 + 11 + 12) Thrust - Fancy iterators ‣ transform_iterator - Allows us to fuse separate algorithms into one // initialize vector device_vector vec(3); vec[0] = 10; vec[1] = 20; vec[2] = 30; // create iterator (type omitted) begin = make_transform_iterator(vec.begin(), negate()); end = make_transform_iterator(vec.end(), negate()); begin[0] begin[1] begin[2] // returns ‐10 // returns ‐20 // returns ‐30 // sum of [begin, end) reduce(begin, end); // returns ‐60 (i.e. ‐10 + ‐20 + ‐30) Thrust - Fancy iterators ‣ permutation_iterator - Allows to fuse gather and scatter operations #include ... //gather locations thrust::device_vector map(4); map[0] = 3; map[1] = 1; map[2] = 0; map[3] = 5; //array to gather from thrust::device_vector source(6); source[0] = 10; source[1] = 20; source[2] = 30; source[3] = 40; source[4] = 50; source[5] = 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())); Thrust - Fancy iterators ‣ zip_iterator - Looks like an array of structs - Stored in structure of arrays // initialize vectors device_vector A(3); device_vector B(3); A[0] = 10; A[1] = 20; A[2] = 30; B[0] = ‘x’; B[1] = ‘y’; B[2] = ‘z’; // create iterator (type omitted) begin = make_zip_iterator(make_tuple(A.begin(), B.begin())); end = make_zip_iterator(make_tuple(A.end(), B.end())); begin[0] begin[1] begin[2] // returns tuple(10, ‘x’) // returns tuple(20, ‘y’) // returns tuple(30, ‘z’) // maximum of [begin, end) maximum< tuple > binary_op; reduce(begin, end, begin[0], binary_op); // returns tuple(30, ‘z’) Thrust - Best practices ‣ Fusion - Combine related operations together ‣ Structure of Arrays - Ensure memory coalescing ‣ Implicit Sequences - Eliminate memory accesses Thrust - Fusion ‣ Combine related operations together - Conserves memory bandwidth ‣ Example: norm of a vector - Square each element - Compute the sum of squares and take sqrt() Thrust - Fusion ‣ Unoptimized example // define transformation f(x) ‐> x^2 struct square { __host__ __device__ float operator()(float x) { return x * x; } }; float snrm2_slow(device_vector& x) { // without fusion device_vector temp(x.size()); transform(x.begin(), x.end(), temp.begin(), square()); return sqrt( reduce(temp.begin(), temp.end()) ); } Thrust - Fusion ‣ Optimized implementation (3.8x) // define transformation f(x) ‐> x^2 struct square { __host__ __device__ float operator()(float x) { return x * x; } }; float snrm2_fast(device_vector& x) { // with fusion return sqrt( transform_reduce(x.begin(), x.end(), square(), 0.0f, plus()); } Thrust - Structure of Arrays (SoA) ‣ Array of structures (AoS) - Often does not obey coalescing rules device_vector ‣ Structure of arrays (SoA) - Obeys coealescing rules - Components stored in separate arrays device_vector x,y,z; ‣ Example: rotate 3D vectors Thrust - Structure of Arrays (SoA) struct rotate_float3 { __host__ __device__ float3 operator()(float3 v) { float x = v.x; float y = v.y; float z = v.z; float rx = 0.36f*x + 0.48f*y + ‐0.80f*z; float ry =‐0.80f*x + 0.60f*y + 0.00f*z; float rz = 0.48f*x + 0.64f*y + 0.60f*z; return make_float3(rx, ry, rz); } }; ... device_vector vec(N); transform(vec.begin(), vec.end, vec.begin(), rotate_float3()); Thrust - Structure of Arrays (SoA) struct rotate_tuple { __host__ __device__ tuple operator()(tuple v) { float x = get<0>(v); float y = get<1>(v); float z = get<2>(v); float rx = 0.36f*x + 0.48f*y + ‐0.80f*z; float ry =‐0.80f*x + 0.60f*y + 0.00f*z; float rz = 0.48f*x + 0.64f*y + 0.60f*z; return make_tuple(rx, ry, rz); } }; ... device_vector x(N), y(N), z(N); transform(make_zip_iterator(make_tuple(x.begin(), y.begin(),z.begin())), make_zip_iterator(make_tuple(x.end(), y.end(), z.end())), make_zip_iterator(make_tuple(x.begin(), y.begin(),z.begin())), rotate_tuple()); Thrust - Implicit sequences ‣ Avoid storing sequences explicitly - Constant sequences - Incrementing sequences ‣ Implicit sequences require no storage constant_iterator counting_iterator ‣ Example - Index of the smallest element Thrust - Implicit sequences // return the smaller of two tuples struct smaller_tuple { tuple operator()(tuple a, tuple b) { if (a < b) return a; else return b; } }; int min_index(device_vector& vec) { // create explicit index sequence [0, 1, 2, ... ) device_vector indices(vec.size()); sequence(indices.begin(), indices.end()); tuple init(vec[0],0); tuple smallest; smallest = reduce(make_zip_iterator(make_tuple(vec.begin(), indices.begin())), make_zip_iterator(make_tuple(vec.end(), indices.end())), init, smaller_tuple()); return get<1>(smallest); } Thrust - Implicit sequences // return the smaller of two tuples struct smaller_tuple { tuple operator()(tuple a, tuple b) { if (a < b) return a; else return b; } }; int min_index(device_vector& vec) { // create implicit index sequence [0, 1, 2, ... ) counting_iterator begin(0); counting_iterator end(vec.size()); tuple init(vec[0],0); tuple smallest; smallest = reduce(make_zip_iterator(make_tuple(vec.begin(), begin)), make_zip_iterator(make_tuple(vec.end(), end)), init, smaller_tuple()); return get<1>(small); }