1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013-2014 Kyle Lutz <kyle.r.lutz@gmail.com>
3 //
4 // Distributed under the Boost Software License, Version 1.0
5 // See accompanying file LICENSE_1_0.txt or copy at
6 // http://www.boost.org/LICENSE_1_0.txt
7 //
8 // See http://boostorg.github.com/compute for more information.
9 //---------------------------------------------------------------------------//
10
11 #include <cstdlib>
12 #include <iostream>
13
14 #include <boost/compute/command_queue.hpp>
15 #include <boost/compute/system.hpp>
16 #include <boost/compute/algorithm/copy_n.hpp>
17 #include <boost/compute/container/vector.hpp>
18 #include <boost/compute/utility/source.hpp>
19
20 namespace compute = boost::compute;
21
22 // return a random float between lo and hi
rand_float(float lo,float hi)23 float rand_float(float lo, float hi)
24 {
25 float x = (float) std::rand() / (float) RAND_MAX;
26
27 return (1.0f - x) * lo + x * hi;
28 }
29
30 // this example demostrates a black-scholes option pricing kernel.
main()31 int main()
32 {
33 // number of options
34 const int N = 4000000;
35
36 // black-scholes parameters
37 const float risk_free_rate = 0.02f;
38 const float volatility = 0.30f;
39
40 // get default device and setup context
41 compute::device gpu = compute::system::default_device();
42 compute::context context(gpu);
43 compute::command_queue queue(context, gpu);
44 std::cout << "device: " << gpu.name() << std::endl;
45
46 // initialize option data on host
47 std::vector<float> stock_price_data(N);
48 std::vector<float> option_strike_data(N);
49 std::vector<float> option_years_data(N);
50
51 std::srand(5347);
52 for(int i = 0; i < N; i++){
53 stock_price_data[i] = rand_float(5.0f, 30.0f);
54 option_strike_data[i] = rand_float(1.0f, 100.0f);
55 option_years_data[i] = rand_float(0.25f, 10.0f);
56 }
57
58 // create memory buffers on the device
59 compute::vector<float> call_result(N, context);
60 compute::vector<float> put_result(N, context);
61 compute::vector<float> stock_price(N, context);
62 compute::vector<float> option_strike(N, context);
63 compute::vector<float> option_years(N, context);
64
65 // copy initial values to the device
66 compute::copy_n(stock_price_data.begin(), N, stock_price.begin(), queue);
67 compute::copy_n(option_strike_data.begin(), N, option_strike.begin(), queue);
68 compute::copy_n(option_years_data.begin(), N, option_years.begin(), queue);
69
70 // source code for black-scholes program
71 const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
72 // approximation of the cumulative normal distribution function
73 static float cnd(float d)
74 {
75 const float A1 = 0.319381530f;
76 const float A2 = -0.356563782f;
77 const float A3 = 1.781477937f;
78 const float A4 = -1.821255978f;
79 const float A5 = 1.330274429f;
80 const float RSQRT2PI = 0.39894228040143267793994605993438f;
81
82 float K = 1.0f / (1.0f + 0.2316419f * fabs(d));
83 float cnd =
84 RSQRT2PI * exp(-0.5f * d * d) *
85 (K * (A1 + K * (A2 + K * (A3 + K * (A4 + K * A5)))));
86
87 if(d > 0){
88 cnd = 1.0f - cnd;
89 }
90
91 return cnd;
92 }
93
94 // black-scholes option pricing kernel
95 __kernel void black_scholes(__global float *call_result,
96 __global float *put_result,
97 __global const float *stock_price,
98 __global const float *option_strike,
99 __global const float *option_years,
100 float risk_free_rate,
101 float volatility)
102 {
103 const uint opt = get_global_id(0);
104
105 float S = stock_price[opt];
106 float X = option_strike[opt];
107 float T = option_years[opt];
108 float R = risk_free_rate;
109 float V = volatility;
110
111 float sqrtT = sqrt(T);
112 float d1 = (log(S / X) + (R + 0.5f * V * V) * T) / (V * sqrtT);
113 float d2 = d1 - V * sqrtT;
114 float CNDD1 = cnd(d1);
115 float CNDD2 = cnd(d2);
116
117 float expRT = exp(-R * T);
118 call_result[opt] = S * CNDD1 - X * expRT * CNDD2;
119 put_result[opt] = X * expRT * (1.0f - CNDD2) - S * (1.0f - CNDD1);
120 }
121 );
122
123 // build black-scholes program
124 compute::program program = compute::program::create_with_source(source, context);
125 program.build();
126
127 // setup black-scholes kernel
128 compute::kernel kernel(program, "black_scholes");
129 kernel.set_arg(0, call_result);
130 kernel.set_arg(1, put_result);
131 kernel.set_arg(2, stock_price);
132 kernel.set_arg(3, option_strike);
133 kernel.set_arg(4, option_years);
134 kernel.set_arg(5, risk_free_rate);
135 kernel.set_arg(6, volatility);
136
137 // execute black-scholes kernel
138 queue.enqueue_1d_range_kernel(kernel, 0, N, 0);
139
140 // print out the first option's put and call prices
141 float call0, put0;
142 compute::copy_n(put_result.begin(), 1, &put0, queue);
143 compute::copy_n(call_result.begin(), 1, &call0, queue);
144
145 std::cout << "option 0 call price: " << call0 << std::endl;
146 std::cout << "option 0 put price: " << put0 << std::endl;
147
148 // due to the differences in the random-number generators between Operating Systems
149 // and/or compilers, we will get different "expected" results for this example
150 #ifdef __APPLE__
151 double expected_call0 = 0.000249461;
152 double expected_put0 = 26.2798;
153 #elif _MSC_VER
154 double expected_call0 = 8.21412;
155 double expected_put0 = 2.25904;
156 #else
157 double expected_call0 = 0.0999f;
158 double expected_put0 = 43.0524f;
159 #endif
160
161 // check option prices
162 if(std::abs(call0 - expected_call0) > 1e-4 || std::abs(put0 - expected_put0) > 1e-4){
163 std::cerr << "error: option prices are wrong" << std::endl;
164 return -1;
165 }
166
167 return 0;
168 }
169