summaryrefslogtreecommitdiff
path: root/benchmarks/CUDA/BlackScholes
diff options
context:
space:
mode:
Diffstat (limited to 'benchmarks/CUDA/BlackScholes')
-rw-r--r--benchmarks/CUDA/BlackScholes/BlackScholes.cu268
-rw-r--r--benchmarks/CUDA/BlackScholes/BlackScholes_gold.cpp117
-rw-r--r--benchmarks/CUDA/BlackScholes/BlackScholes_kernel.cuh122
-rw-r--r--benchmarks/CUDA/BlackScholes/Makefile51
4 files changed, 558 insertions, 0 deletions
diff --git a/benchmarks/CUDA/BlackScholes/BlackScholes.cu b/benchmarks/CUDA/BlackScholes/BlackScholes.cu
new file mode 100644
index 0000000..0c4ab7f
--- /dev/null
+++ b/benchmarks/CUDA/BlackScholes/BlackScholes.cu
@@ -0,0 +1,268 @@
+/*
+ * 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
new file mode 100644
index 0000000..2405fea
--- /dev/null
+++ b/benchmarks/CUDA/BlackScholes/BlackScholes_gold.cpp
@@ -0,0 +1,117 @@
+/*
+ * 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
new file mode 100644
index 0000000..d99f69a
--- /dev/null
+++ b/benchmarks/CUDA/BlackScholes/BlackScholes_kernel.cuh
@@ -0,0 +1,122 @@
+/*
+ * 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
new file mode 100644
index 0000000..0a628ac
--- /dev/null
+++ b/benchmarks/CUDA/BlackScholes/Makefile
@@ -0,0 +1,51 @@
+################################################################################
+#
+# 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
+
+include ../../../common/common.mk
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 23:48:17 +0000