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 evaluates monotone spline for a sequence of query points residing in the array z of size m
  15 */
  16 template<typename Tx, typename Ty>
  17 __global__ void d_MonSplineValue(Tx* z, int K, double* t, double * alpha, double * beta, double * gamma, int T, Ty *value)
  18 {
  19   int tid = threadIdx.x + blockIdx.x * blockDim.x;
  20   int mi=0, ma=T, i=0;
  21   Ty r;
  22   while(tid<K)
  23   {
  24      Bisection_device(z[tid], t, mi, ma,  &i);
  25      r= z[tid]-t[i];
  26      r= alpha[i] + r*(beta[i] + gamma[i]*r);
  27      value[tid]=r;
  28      tid += blockDim.x * gridDim.x;
  29    }
  30    __syncthreads();
  31 }
  32
  33 template<typename Tx, typename Ty>
  34 void MonotoneSplineValue(Tx *z, int K, double* t, double * alpha, double * beta, double * gamma, int T, Ty* result)
  35 {
  36   int blocks,threads=256;
  37   blocks=(K-1)/threads+1;
  38   d_MonSplineValue<<<blocks,threads>>>(z,K,t,alpha,beta,gamma,T,result);
  39 }