BookGPU/Chapters/chapter11/code4.cu

   1 template<typename Tx>
   2 __device__ Tx Aver(Tx z,int i,int j, Tx *z) {return (z-z[j+1])/(j-i+1);}
   3
   4 template<typename Tx>
   5 __global__ void monotonizekernel(Tx *y, Tx *z, Tx *u, int *key, int n)
   6 { int i = threadIdx.x + blockIdx.x * blockDim.x;
   7    if(i<n) {
   8       int smallestJ = i;
   9       Tx curP, smallestP, curz=z[i];
  10       smallestP=Aver(curz,i,i,z);
  11       for(int j = i+1; j < n; j++) {
  12           curP=Aver(curz,i,j,z);
  13           if(smallestP>curP) {
  14                smallestJ = j;
  15                smallestP = curP;
  16           }
  17       }
  18       curP=y[i];
  19       if(curP > smallestP)
  20         t=smallestP;
  21       else
  22                                 smallestJ=i;
  23       key[i]=smallestJ;
  24       u[i]=t;
  25    }
  26 }
  27
  28 template< typename Tx >
  29 void MonotonizeData(Tx *y, int n, Tx *u) {
  30     thrust::less_equal<int> binary_pred;
  31     thrust::maximum<Tx>     binary_op2;
  32     thrust::device_vector<Tx> z_d(n+1);
  33     thrust::device_vector<int> keys_d(n);
  34     thrust::device_ptr<Tx> y_d(y), u_d(u);
  35     thrust::fill(u_d, u_d+n, -1e100);
  36     thrust::fill(keys_d.begin(), keys_d.end(), 0);
  37
  38     thrust::reverse_iterator< typename thrust::device_vector<Tx>::iterator >  y_reverse_b(y_d+n), y_reverse_end(y_d), z_reverse_b(z_d.end());
  39
  40     thrust::inclusive_scan(y_reverse_b, y_reverse_end, z_reverse_b+1);
  41
  42     monotonizekernel<<<grid, block>>>(y, thrust::raw_pointer_cast(&z_d[0]), u, thrust::raw_pointer_cast(&keys_d[0]), n );
  43
  44     thrust::sort(keys_d.begin(), keys_d.end());
  45     thrust::inclusive_scan_by_key(keys_d.begin(), keys_d.end(), u_d, u_d, binary_pred, binary_op2);
  46 }