utiltools/X86MAC64/cuda/samples/4_Finance/binomialOptions/binomialOptions_kernel.cuh

   1 /*
   2  * Copyright 1993-2013 NVIDIA Corporation.  All rights reserved.
   3  *
   4  * Please refer to the NVIDIA end user license agreement (EULA) associated
   5  * with this source code for terms and conditions that govern your use of
   6  * this software. Any use, reproduction, disclosure, or distribution of
   7  * this software and related documentation outside the terms of the EULA
   8  * is strictly prohibited.
   9  *
  10  */
  11
  12
  13
  14 ////////////////////////////////////////////////////////////////////////////////
  15 // Global types and parameters
  16 ////////////////////////////////////////////////////////////////////////////////
  17 #include <stdio.h>
  18 #include <stdlib.h>
  19 #include <helper_cuda.h>
  20 #include "realtype.h"
  21 #include "binomialOptions_common.h"
  22
  23
  24
  25 ////////////////////////////////////////////////////////////////////////////////
  26 // Internal GPU-side constants and data structures
  27 ////////////////////////////////////////////////////////////////////////////////
  28 #define  TIME_STEPS 16
  29 #define CACHE_DELTA (2 * TIME_STEPS)
  30 #define  CACHE_SIZE (256)
  31 #define  CACHE_STEP (CACHE_SIZE - CACHE_DELTA)
  32
  33 #if NUM_STEPS % CACHE_DELTA
  34 #error Bad constants
  35 #endif
  36
  37 //Preprocessed input option data
  38 typedef struct
  39 {
  40     real S;
  41     real X;
  42     real vDt;
  43     real puByDf;
  44     real pdByDf;
  45 } __TOptionData;
  46 static __constant__ __TOptionData d_OptionData[MAX_OPTIONS];
  47 static __device__           float d_CallValue[MAX_OPTIONS];
  48 static __device__            real d_CallBuffer[MAX_OPTIONS * (NUM_STEPS + 16)];
  49
  50
  51
  52 ////////////////////////////////////////////////////////////////////////////////
  53 // Overloaded shortcut functions for different precision modes
  54 ////////////////////////////////////////////////////////////////////////////////
  55 #ifndef DOUBLE_PRECISION
  56 __device__ inline float expiryCallValue(float S, float X, float vDt, int i)
  57 {
  58     real d = S * expf(vDt * (2.0f * i - NUM_STEPS)) - X;
  59     return (d > 0) ? d : 0;
  60 }
  61 #else
  62 __device__ inline double expiryCallValue(double S, double X, double vDt, int i)
  63 {
  64     double d = S * exp(vDt * (2.0 * i - NUM_STEPS)) - X;
  65     return (d > 0) ? d : 0;
  66 }
  67 #endif
  68
  69
  70 ////////////////////////////////////////////////////////////////////////////////
  71 // GPU kernel
  72 ////////////////////////////////////////////////////////////////////////////////
  73 static __global__ void binomialOptionsKernel()
  74 {
  75     __shared__ real callA[CACHE_SIZE+1];
  76     __shared__ real callB[CACHE_SIZE+1];
  77     //Global memory frame for current option (thread block)
  78     real *const d_Call = &d_CallBuffer[blockIdx.x * (NUM_STEPS + 16)];
  79
  80     const int       tid = threadIdx.x;
  81     const real      S = d_OptionData[blockIdx.x].S;
  82     const real      X = d_OptionData[blockIdx.x].X;
  83     const real    vDt = d_OptionData[blockIdx.x].vDt;
  84     const real puByDf = d_OptionData[blockIdx.x].puByDf;
  85     const real pdByDf = d_OptionData[blockIdx.x].pdByDf;
  86
  87     //Compute values at expiry date
  88     for (int i = tid; i <= NUM_STEPS; i += CACHE_SIZE)
  89     {
  90         d_Call[i] = expiryCallValue(S, X, vDt, i);
  91     }
  92
  93     //Walk down binomial tree
  94     //So double-buffer and synchronize to avoid read-after-write hazards.
  95     for (int i = NUM_STEPS; i > 0; i -= CACHE_DELTA)
  96         for (int c_base = 0; c_base < i; c_base += CACHE_STEP)
  97         {
  98             //Start and end positions within shared memory cache
  99             int c_start = min(CACHE_SIZE - 1, i - c_base);
 100             int c_end   = c_start - CACHE_DELTA;
 101
 102             //Read data(with apron) to shared memory
 103             __syncthreads();
 104
 105             if (tid <= c_start)
 106             {
 107                 callA[tid] = d_Call[c_base + tid];
 108             }
 109
 110             //Calculations within shared memory
 111             for (int k = c_start - 1; k >= c_end;)
 112             {
 113                 //Compute discounted expected value
 114                 __syncthreads();
 115                 callB[tid] = puByDf * callA[tid + 1] + pdByDf * callA[tid];
 116                 k--;
 117
 118                 //Compute discounted expected value
 119                 __syncthreads();
 120                 callA[tid] = puByDf * callB[tid + 1] + pdByDf * callB[tid];
 121                 k--;
 122             }
 123
 124             //Flush shared memory cache
 125             __syncthreads();
 126
 127             if (tid <= c_end)
 128             {
 129                 d_Call[c_base + tid] = callA[tid];
 130             }
 131         }
 132
 133     //Write the value at the top of the tree to destination buffer
 134     if (threadIdx.x == 0)
 135     {
 136         d_CallValue[blockIdx.x] = (float)callA[0];
 137     }
 138 }
 139
 140
 141
 142 ////////////////////////////////////////////////////////////////////////////////
 143 // Host-side interface to GPU binomialOptions
 144 ////////////////////////////////////////////////////////////////////////////////
 145 static void binomialOptionsGPU(
 146     float *callValue,
 147     TOptionData  *optionData,
 148     int optN
 149 )
 150 {
 151     __TOptionData h_OptionData[MAX_OPTIONS];
 152
 153     for (int i = 0; i < optN; i++)
 154     {
 155         const double      T = optionData[i].T;
 156         const double      R = optionData[i].R;
 157         const double      V = optionData[i].V;
 158
 159         const double     dt = T / (double)NUM_STEPS;
 160         const double    vDt = V * sqrt(dt);
 161         const double    rDt = R * dt;
 162         //Per-step interest and discount factors
 163         const double     If = exp(rDt);
 164         const double     Df = exp(-rDt);
 165         //Values and pseudoprobabilities of upward and downward moves
 166         const double      u = exp(vDt);
 167         const double      d = exp(-vDt);
 168         const double     pu = (If - d) / (u - d);
 169         const double     pd = 1.0 - pu;
 170         const double puByDf = pu * Df;
 171         const double pdByDf = pd * Df;
 172
 173         h_OptionData[i].S      = (real)optionData[i].S;
 174         h_OptionData[i].X      = (real)optionData[i].X;
 175         h_OptionData[i].vDt    = (real)vDt;
 176         h_OptionData[i].puByDf = (real)puByDf;
 177         h_OptionData[i].pdByDf = (real)pdByDf;
 178     }
 179
 180     checkCudaErrors(cudaMemcpyToSymbol(d_OptionData, h_OptionData, optN * sizeof(__TOptionData)));
 181     binomialOptionsKernel<<<optN, CACHE_SIZE>>>();
 182     getLastCudaError("binomialOptionsKernel() execution failed.\n");
 183     checkCudaErrors(cudaMemcpyFromSymbol(callValue, d_CallValue, optN *sizeof(float)));
 184 }