use customised stream in thrust APIs

[lgarithm]

@xiaoming-qxm @baoleai we can use thrust::cuda::par.on(stream).

e.g.

#include <cuda_runtime.h>

#include <thrust/device_vector.h>
#include <thrust/transform.h>

template <typename T>
class cap_by
{
    const T cap;

  public:
    cap_by(const T cap) : cap(cap) {}

    __host__ __device__ T operator()(T x) const
    {
        if (x > cap) { return cap; }
        return x;
    }
};

void f(cudaStream_t stream)
{
    int n = 1 << 10;
    using T = int;
    thrust::device_vector<T> xs(n);
    thrust::device_vector<T> ys(n);

    thrust::sort(thrust::cuda::par.on(stream), xs.begin(), xs.end());
    int k = 5;
    thrust::transform(thrust::cuda::par.on(stream), xs.begin(), xs.end(),
                      ys.begin(), cap_by<T>(k));
}

int main()
{
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    f(stream);
    cudaStreamDestroy(stream);
    return 0;
}

[xiaoming-qxm]

Awesome! So we can port pytorch_quiver to actor framework easily.

[xiaoming-qxm]

But it seems that thrust APIs doesn't support asynchronous behaviors, unavoidable synchronization after every algorithm call is a performance killer.

Reference: NVIDIA/thrust#961

[lgarithm]

• async algorithms #17