Pass cuda array to thrust::inclusive_scan

Question

I can use inclusive_scan for an array on the cpu, but is it possible to do it with an array on the gpu? (commentated is the way that i know works but that I don't need). Alternatively, are there any other easy methods to perform an inclusive scan on an array in device memory?

Code:

#include <stdio.h>
#include <stdlib.h> /* for rand() */
#include <unistd.h> /* for getpid() */
#include <time.h> /* for time() */
#include <math.h>
#include <assert.h>
#include <iostream>
#include <ctime>
  #include <thrust/scan.h>
#include <cuda.h>



#ifdef DOUBLE
 #define REAL double
 #define MAXT 256
#else
 #define REAL float
 #define MAXT 512
#endif

#ifndef MIN
#define MIN(x,y) ((x < y) ? x : y)
#endif

using namespace std;

bool errorAsk(const char *s="n/a")
{
cudaError_t err=cudaGetLastError();
if(err==cudaSuccess)
    return false;
printf("CUDA error [%s]: %s\n",s,cudaGetErrorString(err));
return true;
};

double *fillArray(double *c_idata,int N,double constant) {
    int n;
    for (n = 0; n < N; n++) {
            c_idata[n] = constant*floor(drand48()*10);

    }
return c_idata;
}

int main(int argc,char *argv[])
{
    int N,blocks,threads;
    N = 100;
    threads=MAXT;
    blocks=N/threads+(N%threads==0?0:1);

    double *c_data,*g_data;

    c_data = new double[N];
    c_data = fillArray(c_data,N,1);
    cudaMalloc(&g_data,N*sizeof(double));

    cudaMemcpy(g_data,c_data,N*sizeof(double),cudaMemcpyHostToDevice);
    thrust::inclusive_scan(g_data, g_data + N, g_data); // in-place scan
    cudaMemcpy(c_data,g_data,N*sizeof(double),cudaMemcpyDeviceToHost);

//        thrust::inclusive_scan(c_data, c_data + N, c_data); // in-place scan

    for(int i = 0; i < N; i++) {
            cout<<c_data[i]<<endl;
    }
}

Answer 1

If you read the thrust quick start guide you'll find one suggestion for handling "raw" device data: use a thrust::device_ptr:

You may wonder what happens when a "raw" pointer is used as an argument to a Thrust function. Like the STL, Thrust permits this usage and it will dispatch the host path of the algorithm. If the pointer in question is in fact a pointer to device memory then you'll need to wrap it with thrust::device_ptr before calling the function.

To fix your code, you would want to

#include <thrust/device_ptr.h>

and replace your existing call to thrust::inclusive_scan with the following 2 lines:

thrust::device_ptr<double> g_ptr = thrust::device_pointer_cast(g_data);
thrust::inclusive_scan(g_ptr, g_ptr + N, g_ptr); // in-place scan

Another approach would be to use thrust execution policies and modify your call like this:

thrust::inclusive_scan(thrust::device, g_data, g_data + N, g_data);

And there are various other possibilities as well.