BookGPU/Chapters/chapter11/code3.cu

   1 template<typename T>
   2 __device__ void Bisection_device(T z, T* t, int mi,int ma,int* l)
   3 {
   4   int i; ma--;
   5   while(1) {
   6     i=(mi+ma)/2;
   7     if(z >= t[i]) mi=i+1;
   8              else ma=i;
   9     if(mi>=ma) break;
  10   }
  11   *l = mi-1;
  12 }
  13
  14 // Kernel to evaluate monotone spline for a sequence of query points
  15 // residing in the array z of size m
  16
  17 template<typename Tx, typename Ty>
  18 __global__ void d_MonSplineValue(Tx* z, int K, double* t, double * alpha, double * beta, double * gamma, int T, Ty *value)
  19 {
  20   int tid = threadIdx.x + blockIdx.x * blockDim.x;
  21   int mi=0, ma=T, i=0;
  22   Ty r;
  23   while(tid<K)
  24   {
  25      Bisection_device(z[tid], t, mi, ma,  &i);
  26      r= z[tid]-t[i];
  27      r= alpha[i] + r*(beta[i] + gamma[i]*r);
  28      value[tid]=r;
  29      tid += blockDim.x * gridDim.x;
  30    }
  31    __syncthreads();
  32 }
  33
  34 template<typename Tx, typename Ty>
  35 void MonotoneSplineValue(Tx *z, int K, double* t, double * alpha, double * beta, double * gamma, int T, Ty* result)
  36 {
  37   int blocks,threads=256;
  38   blocks=(K-1)/threads+1;
  39   d_MonSplineValue<<<blocks,threads>>>(z,K,t,alpha,beta,gamma,T,result);
  40 }