Purpose and usage of counting_iterators in CUDA Th

Question

I have trouble understanding counting_iterator in thrust library for CUDA. What is its purpose and how is it used? Is it available in other programming languages such as C++ also?

Answer 1

As a further example, which is a slight modification of @talonmies answer, you can emulate Matlab's linspace command by the following code

#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/constant_iterator.h>
#include <cstdio>

void main()
{ 
    const int N = 20;

    float a     = 3.87f;
    float b     = 7.11f;

    float Dx    = (b-a)/(float)(N-1);

    thrust::device_vector<float> myvector(N);

    thrust::transform(thrust::make_counting_iterator(a/Dx),
                  thrust::make_counting_iterator((b+1.f)/Dx),
                  thrust::make_constant_iterator(Dx),
                  myvector.begin(),
                  thrust::multiplies<float>());

    for(int i=0; i<N; i++) {
        float val = myvector[i];
        printf("%d %f\n", i, val);
    }

    getchar();

}

The above code will return

0 3.870000
1 4.040526
2 4.211052
3 4.381579
4 4.552105
5 4.722631
6 4.893158
7 5.063684
8 5.234210
9 5.404737
10 5.575263
11 5.745790
12 5.916316
13 6.086842
14 6.257369
15 6.427895
16 6.598421
17 6.768948
18 6.939474
19 7.110000

exacltly as linspace(3.87,7.11,20).

Answer 2

A counting iterator is just an iterator which returns the next value from a sequence which is advanced each time the iterator is incremented. The simplest possible example is something like this:

#include <iostream>
#include <thrust/iterator/counting_iterator.h>

int main(void)
{
    int n = 10;

    thrust::counting_iterator<int>x(1);

    for(int i=0; i<n; ++i, ++x) {
        std::cout << *x << std::endl;
    }

    return 0;
}

which does this when compiled and run:

$ /usr/local/cuda/bin/nvcc counting.cc 
$ ./a.out
1
2
3
4
5
6
7
8
9
10

ie. the counting iterator was initialised with a value of one, and each time the iterator is incremented, we get the next value in the counting sequence, ie. 1,2,3,4,5....

Where counting increments come in handy in thrust is any time you need a sequence to fill a vector or manipulate in a transform iterator or functor. The counting iterator removes the need to explicitly create and fill a vector with the sequence you require. For example (from my answer to this Stack Overflow question):

#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/constant_iterator.h>
#include <cstdio>

int main(void)
{
    const int N = 18, M = 3;
    thrust::device_vector<int> myvector(N);

    thrust::transform(  thrust::make_counting_iterator(0),
                        thrust::make_counting_iterator(N),
                        thrust::make_constant_iterator(M),
                        myvector.begin(),
                        thrust::divides<int>() );

    for(int i=0; i<N; i++) {
        int val = myvector[i];
        printf("%d %d\n", i, val);
    }
    return 0;
}

which produces a device vector filled with the sequence 0, 0, 0, 1, 1, 1, 2, 2 ,2, 3, 3, 3.