eic-opticks/sysrap/strided_range.h

0001 /*
0002  * Copyright (c) 2019 Opticks Team. All Rights Reserved.
0003  *
0004  * This file is part of Opticks
0005  * (see https://bitbucket.org/simoncblyth/opticks).
0006  *
0007  * Licensed under the Apache License, Version 2.0 (the "License");
0008  * you may not use this file except in compliance with the License.
0009  * You may obtain a copy of the License at
0010  *
0011  *   http://www.apache.org/licenses/LICENSE-2.0
0012  *
0013  * Unless required by applicable law or agreed to in writing, software
0014  * distributed under the License is distributed on an "AS IS" BASIS,
0015  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
0016  * See the License for the specific language governing permissions and
0017  * limitations under the License.
0018  */
0019
0020 #pragma once
0021
0022 /**
0023 strided_range.h
0024 ==================
0025
0026
0027 Based on /usr/local/env/numerics/thrust/examples/strided_range.cu
0028
0029 This example illustrates how to make strided access to a range of values
0030 examples::
0031
0032    strided_range([0, 1, 2, 3, 4, 5, 6], 1) -> [0, 1, 2, 3, 4, 5, 6]
0033    strided_range([0, 1, 2, 3, 4, 5, 6], 2) -> [0, 2, 4, 6]
0034    strided_range([0, 1, 2, 3, 4, 5, 6], 3) -> [0, 3, 6]
0035    ...
0036
0037 This enables the plucking of photon counts from the GPU side
0038 genstep buffer, as used by seeding in okop-::
0039
0040     195 void OpSeeder::seedPhotonsFromGenstepsImp(const CBufSpec& s_gs, const CBufSpec& s_ox)
0041     196 {
0042     ...
0043     235     // src slice is plucking photon counts from each genstep
0044     237     // buffer size and num_bytes comes directly from CBufSpec
0045     238     CBufSlice src = tgs.slice(6*4,3,num_genstep_values) ;  // stride, begin, end
0046     ...
0047
0048
0049 **/
0050
0051
0052 #include <thrust/iterator/counting_iterator.h>
0053 #include <thrust/iterator/transform_iterator.h>
0054 #include <thrust/iterator/permutation_iterator.h>
0055 #include <thrust/functional.h>
0056 #include <thrust/device_vector.h>
0057
0058 template <typename Iterator>
0059 class strided_range
0060 {
0061     public:
0062
0063     typedef typename thrust::iterator_difference<Iterator>::type difference_type;
0064
0065     struct stride_functor
0066     {
0067         difference_type stride;
0068
0069         stride_functor(difference_type stride)
0070             : stride(stride) {}
0071
0072         __host__ __device__
0073         difference_type operator()(const difference_type& i) const
0074         {
0075             return stride * i;
0076         }
0077     };
0078
0079     typedef typename thrust::counting_iterator<difference_type>                   CountingIterator;
0080     typedef typename thrust::transform_iterator<stride_functor, CountingIterator> TransformIterator;
0081     typedef typename thrust::permutation_iterator<Iterator,TransformIterator>     PermutationIterator;
0082
0083     // type of the strided_range iterator
0084     typedef PermutationIterator iterator;
0085
0086     // construct strided_range for the range [first,last)
0087     strided_range(Iterator first, Iterator last, difference_type stride)
0088         :
0089         first(first),
0090         last(last),
0091         stride(stride)
0092     {
0093     }
0094
0095     iterator begin(void) const
0096     {
0097         return PermutationIterator(first, TransformIterator(CountingIterator(0), stride_functor(stride)));
0098     }
0099
0100     iterator end(void) const
0101     {
0102         return begin() + ((last - first) + (stride - 1)) / stride;
0103     }
0104
0105     protected:
0106     Iterator first;
0107     Iterator last;
0108     difference_type stride;
0109 };
0110
0111
0112
0113