diff options
Diffstat (limited to 'benchmarks/CUDA/BlackScholes')
| -rw-r--r-- | benchmarks/CUDA/BlackScholes/BlackScholes.cu | 268 | ||||
| -rw-r--r-- | benchmarks/CUDA/BlackScholes/BlackScholes_gold.cpp | 117 | ||||
| -rw-r--r-- | benchmarks/CUDA/BlackScholes/BlackScholes_kernel.cuh | 122 | ||||
| -rw-r--r-- | benchmarks/CUDA/BlackScholes/Makefile | 54 |
4 files changed, 0 insertions, 561 deletions
diff --git a/benchmarks/CUDA/BlackScholes/BlackScholes.cu b/benchmarks/CUDA/BlackScholes/BlackScholes.cu deleted file mode 100644 index 0c4ab7f..0000000 --- a/benchmarks/CUDA/BlackScholes/BlackScholes.cu +++ /dev/null @@ -1,268 +0,0 @@ -/* - * Copyright 1993-2007 NVIDIA Corporation. All rights reserved. - * - * NOTICE TO USER: - * - * This source code is subject to NVIDIA ownership rights under U.S. and - * international Copyright laws. Users and possessors of this source code - * are hereby granted a nonexclusive, royalty-free license to use this code - * in individual and commercial software. - * - * NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE - * CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR - * IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH - * REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF - * MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. - * IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, - * OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS - * OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE - * OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE - * OR PERFORMANCE OF THIS SOURCE CODE. - * - * U.S. Government End Users. This source code is a "commercial item" as - * that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of - * "commercial computer software" and "commercial computer software - * documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) - * and is provided to the U.S. Government only as a commercial end item. - * Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through - * 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the - * source code with only those rights set forth herein. - * - * Any use of this source code in individual and commercial software must - * include, in the user documentation and internal comments to the code, - * the above Disclaimer and U.S. Government End Users Notice. - */ - -/* - * This sample evaluates fair call and put prices for a - * given set of European options by Black-Scholes formula. - * See supplied whitepaper for more explanations. - */ - - - -#include <stdlib.h> -#include <stdio.h> -#include <time.h> -#include <string.h> -#include <cutil_inline.h> - - - -//////////////////////////////////////////////////////////////////////////////// -// Process an array of optN options on CPU -//////////////////////////////////////////////////////////////////////////////// -extern "C" void BlackScholesCPU( - float *h_CallResult, - float *h_PutResult, - float *h_StockPrice, - float *h_OptionStrike, - float *h_OptionYears, - float Riskfree, - float Volatility, - int optN -); - - - -//////////////////////////////////////////////////////////////////////////////// -// Process an array of OptN options on GPU -//////////////////////////////////////////////////////////////////////////////// -#include "BlackScholes_kernel.cuh" - - - -//////////////////////////////////////////////////////////////////////////////// -// Helper function, returning uniformly distributed -// random float in [low, high] range -//////////////////////////////////////////////////////////////////////////////// -float RandFloat(float low, float high){ - float t = (float)rand() / (float)RAND_MAX; - return (1.0f - t) * low + t * high; -} - - - -//////////////////////////////////////////////////////////////////////////////// -// Data configuration -//////////////////////////////////////////////////////////////////////////////// -const int OPT_N = 2000000; // 4000000; - -#ifdef __DEVICE_EMULATION__ -const int NUM_ITERATIONS = 1; -#else -const int NUM_ITERATIONS = 1; // 512; -#endif - - -const int OPT_SZ = OPT_N * sizeof(float); -const float RISKFREE = 0.02f; -const float VOLATILITY = 0.30f; - - -//////////////////////////////////////////////////////////////////////////////// -// Main program -//////////////////////////////////////////////////////////////////////////////// -int main(int argc, char **argv){ - //'h_' prefix - CPU (host) memory space - float - //Results calculated by CPU for reference - *h_CallResultCPU, - *h_PutResultCPU, - //CPU copy of GPU results - *h_CallResultGPU, - *h_PutResultGPU, - //CPU instance of input data - *h_StockPrice, - *h_OptionStrike, - *h_OptionYears; - - //'d_' prefix - GPU (device) memory space - float - //Results calculated by GPU - *d_CallResult, - *d_PutResult, - //GPU instance of input data - *d_StockPrice, - *d_OptionStrike, - *d_OptionYears; - - double - delta, ref, sum_delta, sum_ref, max_delta, L1norm, gpuTime; - - unsigned int hTimer; - int i; - - - if( cutCheckCmdLineFlag(argc, (const char**)argv, "device") ) - cutilDeviceInit(argc, argv); - else - cudaSetDevice( cutGetMaxGflopsDeviceId() ); - - cutilCheckError( cutCreateTimer(&hTimer) ); - - printf("Initializing data...\n"); - printf("...allocating CPU memory for options.\n"); - h_CallResultCPU = (float *)malloc(OPT_SZ); - h_PutResultCPU = (float *)malloc(OPT_SZ); - h_CallResultGPU = (float *)malloc(OPT_SZ); - h_PutResultGPU = (float *)malloc(OPT_SZ); - h_StockPrice = (float *)malloc(OPT_SZ); - h_OptionStrike = (float *)malloc(OPT_SZ); - h_OptionYears = (float *)malloc(OPT_SZ); - - printf("...allocating GPU memory for options.\n"); - cutilSafeCall( cudaMalloc((void **)&d_CallResult, OPT_SZ) ); - cutilSafeCall( cudaMalloc((void **)&d_PutResult, OPT_SZ) ); - cutilSafeCall( cudaMalloc((void **)&d_StockPrice, OPT_SZ) ); - cutilSafeCall( cudaMalloc((void **)&d_OptionStrike, OPT_SZ) ); - cutilSafeCall( cudaMalloc((void **)&d_OptionYears, OPT_SZ) ); - - printf("...generating input data in CPU mem.\n"); - srand(5347); - //Generate options set - for(i = 0; i < OPT_N; i++){ - h_CallResultCPU[i] = 0.0f; - h_PutResultCPU[i] = -1.0f; - h_StockPrice[i] = RandFloat(5.0f, 30.0f); - h_OptionStrike[i] = RandFloat(1.0f, 100.0f); - h_OptionYears[i] = RandFloat(0.25f, 10.0f); - } - - printf("...copying input data to GPU mem.\n"); - //Copy options data to GPU memory for further processing - cutilSafeCall( cudaMemcpy(d_StockPrice, h_StockPrice, OPT_SZ, cudaMemcpyHostToDevice) ); - cutilSafeCall( cudaMemcpy(d_OptionStrike, h_OptionStrike, OPT_SZ, cudaMemcpyHostToDevice) ); - cutilSafeCall( cudaMemcpy(d_OptionYears, h_OptionYears, OPT_SZ, cudaMemcpyHostToDevice) ); - printf("Data init done.\n"); - - - printf("Executing Black-Scholes GPU kernel (%i iterations)...\n", NUM_ITERATIONS); - cutilSafeCall( cudaThreadSynchronize() ); - cutilCheckError( cutResetTimer(hTimer) ); - cutilCheckError( cutStartTimer(hTimer) ); - for(i = 0; i < NUM_ITERATIONS; i++){ - BlackScholesGPU<<<480, 128>>>( - d_CallResult, - d_PutResult, - d_StockPrice, - d_OptionStrike, - d_OptionYears, - RISKFREE, - VOLATILITY, - OPT_N - ); - cutilCheckMsg("BlackScholesGPU() execution failed\n"); - } - cutilSafeCall( cudaThreadSynchronize() ); - cutilCheckError( cutStopTimer(hTimer) ); - gpuTime = cutGetTimerValue(hTimer) / NUM_ITERATIONS; - //Both call and put is calculated - printf("Options count : %i \n", 2 * OPT_N); - printf("BlackScholesGPU() time : %f msec\n", gpuTime); - printf("Effective memory bandwidth: %f GB/s\n", ((double)(5 * OPT_N * sizeof(float)) * 1E-9) / (gpuTime * 1E-3)); - printf("Gigaoptions per second : %f \n", ((double)(2 * OPT_N) * 1E-9) / (gpuTime * 1E-3)); - - - printf("Reading back GPU results...\n"); - //Read back GPU results to compare them to CPU results - cutilSafeCall( cudaMemcpy(h_CallResultGPU, d_CallResult, OPT_SZ, cudaMemcpyDeviceToHost) ); - cutilSafeCall( cudaMemcpy(h_PutResultGPU, d_PutResult, OPT_SZ, cudaMemcpyDeviceToHost) ); - - - printf("Checking the results...\n"); - printf("...running CPU calculations.\n"); - //Calculate options values on CPU - BlackScholesCPU( - h_CallResultCPU, - h_PutResultCPU, - h_StockPrice, - h_OptionStrike, - h_OptionYears, - RISKFREE, - VOLATILITY, - OPT_N - ); - - printf("Comparing the results...\n"); - //Calculate max absolute difference and L1 distance - //between CPU and GPU results - sum_delta = 0; - sum_ref = 0; - max_delta = 0; - for(i = 0; i < OPT_N; i++){ - ref = h_CallResultCPU[i]; - delta = fabs(h_CallResultCPU[i] - h_CallResultGPU[i]); - if(delta > max_delta) max_delta = delta; - sum_delta += delta; - sum_ref += fabs(ref); - } - L1norm = sum_delta / sum_ref; - printf("L1 norm: %E\n", L1norm); - printf("Max absolute error: %E\n", max_delta); - printf((L1norm < 1e-6) ? "TEST PASSED\n" : "TEST FAILED\n"); - - - printf("Shutting down...\n"); - printf("...releasing GPU memory.\n"); - cutilSafeCall( cudaFree(d_OptionYears) ); - cutilSafeCall( cudaFree(d_OptionStrike) ); - cutilSafeCall( cudaFree(d_StockPrice) ); - cutilSafeCall( cudaFree(d_PutResult) ); - cutilSafeCall( cudaFree(d_CallResult) ); - - printf("...releasing CPU memory.\n"); - free(h_OptionYears); - free(h_OptionStrike); - free(h_StockPrice); - free(h_PutResultGPU); - free(h_CallResultGPU); - free(h_PutResultCPU); - free(h_CallResultCPU); - cutilCheckError( cutDeleteTimer(hTimer) ); - printf("Shutdown done.\n"); - - cudaThreadExit(); - - cutilExit(argc, argv); -} diff --git a/benchmarks/CUDA/BlackScholes/BlackScholes_gold.cpp b/benchmarks/CUDA/BlackScholes/BlackScholes_gold.cpp deleted file mode 100644 index 2405fea..0000000 --- a/benchmarks/CUDA/BlackScholes/BlackScholes_gold.cpp +++ /dev/null @@ -1,117 +0,0 @@ -/* - * Copyright 1993-2007 NVIDIA Corporation. All rights reserved. - * - * NOTICE TO USER: - * - * This source code is subject to NVIDIA ownership rights under U.S. and - * international Copyright laws. Users and possessors of this source code - * are hereby granted a nonexclusive, royalty-free license to use this code - * in individual and commercial software. - * - * NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE - * CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR - * IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH - * REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF - * MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. - * IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, - * OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS - * OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE - * OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE - * OR PERFORMANCE OF THIS SOURCE CODE. - * - * U.S. Government End Users. This source code is a "commercial item" as - * that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of - * "commercial computer software" and "commercial computer software - * documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) - * and is provided to the U.S. Government only as a commercial end item. - * Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through - * 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the - * source code with only those rights set forth herein. - * - * Any use of this source code in individual and commercial software must - * include, in the user documentation and internal comments to the code, - * the above Disclaimer and U.S. Government End Users Notice. - */ - - - -#include <math.h> - - - -/////////////////////////////////////////////////////////////////////////////// -// Polynomial approximation of cumulative normal distribution function -/////////////////////////////////////////////////////////////////////////////// -static double CND(double d){ - const double A1 = 0.31938153; - const double A2 = -0.356563782; - const double A3 = 1.781477937; - const double A4 = -1.821255978; - const double A5 = 1.330274429; - const double RSQRT2PI = 0.39894228040143267793994605993438; - - double - K = 1.0 / (1.0 + 0.2316419 * fabs(d)); - - double - cnd = RSQRT2PI * exp(- 0.5 * d * d) * - (K * (A1 + K * (A2 + K * (A3 + K * (A4 + K * A5))))); - - if(d > 0) - cnd = 1.0 - cnd; - - return cnd; -} - - -/////////////////////////////////////////////////////////////////////////////// -// Black-Scholes formula for both call and put -/////////////////////////////////////////////////////////////////////////////// -static void BlackScholesBodyCPU( - float& callResult, - float& putResult, - float Sf, //Stock price - float Xf, //Option strike - float Tf, //Option years - float Rf, //Riskless rate - float Vf //Volatility rate -){ - double S = Sf, X = Xf, T = Tf, R = Rf, V = Vf; - - double sqrtT = sqrt(T); - double d1 = (log(S / X) + (R + 0.5 * V * V) * T) / (V * sqrtT); - double d2 = d1 - V * sqrtT; - double CNDD1 = CND(d1); - double CNDD2 = CND(d2); - - //Calculate Call and Put simultaneously - double expRT = exp(- R * T); - callResult = (float)(S * CNDD1 - X * expRT * CNDD2); - putResult = (float)(X * expRT * (1.0 - CNDD2) - S * (1.0 - CNDD1)); -} - - -//////////////////////////////////////////////////////////////////////////////// -// Process an array of optN options -//////////////////////////////////////////////////////////////////////////////// -extern "C" void BlackScholesCPU( - float *h_CallResult, - float *h_PutResult, - float *h_StockPrice, - float *h_OptionStrike, - float *h_OptionYears, - float Riskfree, - float Volatility, - int optN -){ - for(int opt = 0; opt < optN; opt++) - BlackScholesBodyCPU( - h_CallResult[opt], - h_PutResult[opt], - h_StockPrice[opt], - h_OptionStrike[opt], - h_OptionYears[opt], - Riskfree, - Volatility - ); -} diff --git a/benchmarks/CUDA/BlackScholes/BlackScholes_kernel.cuh b/benchmarks/CUDA/BlackScholes/BlackScholes_kernel.cuh deleted file mode 100644 index d99f69a..0000000 --- a/benchmarks/CUDA/BlackScholes/BlackScholes_kernel.cuh +++ /dev/null @@ -1,122 +0,0 @@ -/* - * Copyright 1993-2007 NVIDIA Corporation. All rights reserved. - * - * NOTICE TO USER: - * - * This source code is subject to NVIDIA ownership rights under U.S. and - * international Copyright laws. Users and possessors of this source code - * are hereby granted a nonexclusive, royalty-free license to use this code - * in individual and commercial software. - * - * NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE - * CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR - * IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH - * REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF - * MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. - * IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, - * OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS - * OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE - * OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE - * OR PERFORMANCE OF THIS SOURCE CODE. - * - * U.S. Government End Users. This source code is a "commercial item" as - * that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of - * "commercial computer software" and "commercial computer software - * documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) - * and is provided to the U.S. Government only as a commercial end item. - * Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through - * 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the - * source code with only those rights set forth herein. - * - * Any use of this source code in individual and commercial software must - * include, in the user documentation and internal comments to the code, - * the above Disclaimer and U.S. Government End Users Notice. - */ - - - -/////////////////////////////////////////////////////////////////////////////// -// Polynomial approximation of cumulative normal distribution function -/////////////////////////////////////////////////////////////////////////////// -__device__ inline float cndGPU(float d){ - const float A1 = 0.31938153f; - const float A2 = -0.356563782f; - const float A3 = 1.781477937f; - const float A4 = -1.821255978f; - const float A5 = 1.330274429f; - const float RSQRT2PI = 0.39894228040143267793994605993438f; - - float - K = 1.0f / (1.0f + 0.2316419f * fabsf(d)); - - float - cnd = RSQRT2PI * __expf(- 0.5f * d * d) * - (K * (A1 + K * (A2 + K * (A3 + K * (A4 + K * A5))))); - - if(d > 0) - cnd = 1.0f - cnd; - - return cnd; -} - - -/////////////////////////////////////////////////////////////////////////////// -// Black-Scholes formula for both call and put -/////////////////////////////////////////////////////////////////////////////// -__device__ inline void BlackScholesBodyGPU( - float& CallResult, - float& PutResult, - float S, //Stock price - float X, //Option strike - float T, //Option years - float R, //Riskless rate - float V //Volatility rate -){ - float sqrtT, expRT; - float d1, d2, CNDD1, CNDD2; - - sqrtT = sqrtf(T); - d1 = (__logf(S / X) + (R + 0.5f * V * V) * T) / (V * sqrtT); - d2 = d1 - V * sqrtT; - - CNDD1 = cndGPU(d1); - CNDD2 = cndGPU(d2); - - //Calculate Call and Put simultaneously - expRT = __expf(- R * T); - CallResult = S * CNDD1 - X * expRT * CNDD2; - PutResult = X * expRT * (1.0f - CNDD2) - S * (1.0f - CNDD1); -} - - -//////////////////////////////////////////////////////////////////////////////// -//Process an array of optN options on GPU -//////////////////////////////////////////////////////////////////////////////// -__global__ void BlackScholesGPU( - float *d_CallResult, - float *d_PutResult, - float *d_StockPrice, - float *d_OptionStrike, - float *d_OptionYears, - float Riskfree, - float Volatility, - int optN -){ - //Thread index - const int tid = blockDim.x * blockIdx.x + threadIdx.x; - //Total number of threads in execution grid - const int THREAD_N = blockDim.x * gridDim.x; - - //No matter how small is execution grid or how large OptN is, - //exactly OptN indices will be processed with perfect memory coalescing - for(int opt = tid; opt < optN; opt += THREAD_N) - BlackScholesBodyGPU( - d_CallResult[opt], - d_PutResult[opt], - d_StockPrice[opt], - d_OptionStrike[opt], - d_OptionYears[opt], - Riskfree, - Volatility - ); -} diff --git a/benchmarks/CUDA/BlackScholes/Makefile b/benchmarks/CUDA/BlackScholes/Makefile deleted file mode 100644 index 36e244c..0000000 --- a/benchmarks/CUDA/BlackScholes/Makefile +++ /dev/null @@ -1,54 +0,0 @@ -################################################################################ -# -# Copyright 1993-2006 NVIDIA Corporation. All rights reserved. -# -# NOTICE TO USER: -# -# This source code is subject to NVIDIA ownership rights under U.S. and -# international Copyright laws. -# -# NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE -# CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR -# IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH -# REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF -# MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. -# IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, -# OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS -# OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE -# OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE -# OR PERFORMANCE OF THIS SOURCE CODE. -# -# U.S. Government End Users. This source code is a "commercial item" as -# that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of -# "commercial computer software" and "commercial computer software -# documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) -# and is provided to the U.S. Government only as a commercial end item. -# Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through -# 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the -# source code with only those rights set forth herein. -# -################################################################################ -# -# Build script for project -# -################################################################################ - -# Add source files here -EXECUTABLE := BlackScholes -# CUDA source files (compiled with cudacc) -CUFILES := BlackScholes.cu -# CUDA dependency files -CU_DEPS := BlackScholes_kernel.cuh -# C/C++ source files (compiled with gcc / c++) -CCFILES := BlackScholes_gold.cpp -# Need good occupancy -CUDACCFLAGS := -po maxrregcount=16 - - -################################################################################ -# Rules and targets - -ROOTDIR=$(NVIDIA_CUDA_SDK_LOCATION)/src -BINDIR=. -include $(NVIDIA_CUDA_SDK_LOCATION)/common/common.mk - |
