Archivo BlackScholes_kernel.cuh

• BlackScholes.cu
• BlackScholes_gold.cpp
• BlackScholes_kernel.cuh

1. Encuentre el ayudante del dispositivo cuda : se almacena en un helper_functions.h. Esta función simplemente verifica si hay una GPU NVIDIA disponible en la máquina.

findCudaDevice(argc, (const char **)argv);

h_CallResultCPU = (float *)malloc(OPT_SZ);

free(h_CallResultCPU);

checkCudaErrors(cudaMalloc((void **)&d_CallResult, OPT_SZ));

checkCudaErrors(cudaFree(d_CallResult));

checkCudaErrors(cudaMemcpy(d_StockPrice, h_StockPrice, OPT_SZ, cudaMemcpyHostToDevice));

BlackScholesGPU<<<DIV_UP((OPT_N/2), 128), 128>>>

#include <iostream>
using namespace std;
#define DIV_UP(a, b) ( ((a) + (b) — 1) / (b) )
int main()
{
 cout<<DIV_UP(128, 128) << endl;
 cout<<DIV_UP(400, 128) << endl;
 cout<<DIV_UP(1000, 128) << endl;
 
 return 0;
}

1
4
8

static double CND(double d)
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)
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)

El archivo consta del mismo código c ++ que en el archivo .cpp anterior con un ligero código de sintaxis adicional.

1. La función __global__ también se llama función "kernel". Es la función que puede llamar desde el lado del host usando la semántica de llamada del núcleo CUDA ( <<<...>>>). En nuestro caso, se define en el archivo anterior una función BlackScholesGPU.

__global__ void BlackScholesGPU

__launch_bounds__(128)

__device__ inline void BlackScholesBodyGPU

En cuanto al punto de partida, dejé los datos originales configurados por los desarrolladores de CUDA:

const int OPT_N = 4000000;
const int NUM_ITERATIONS = 512;
const int OPT_SZ = OPT_N * sizeof(float);
const float RISKFREE = 0.02f;
const float VOLATILITY = 0.30f;

GPU Device 0: “Turing” with compute capability 7.5
Initializing data…
…allocating CPU memory for options.
…allocating GPU memory for options.
…generating input data in CPU mem.
…copying input data to GPU mem.
Data init done.
Executing Black-Scholes GPU kernel (512 iterations)…
Options count : 8000000
BlackScholesGPU() time : 2.458114 msec 
Effective memory bandwidth: 32.545271 GB/s
Gigaoptions per second : 3.254527
BlackScholes, Throughput = 3.2545 GOptions/s, Time = 0.00246 s, Size = 8000000 options, NumDevsUsed = 1, Workgroup = 128
Reading back GPU results…
Checking the results…
…running CPU calculations.
Comparing the results…
L1 norm: 1.787766E-07
Max absolute error: 1.192093E-05
Shutting down…
…releasing GPU memory.
…releasing CPU memory.
Shutdown done.
[BlackScholes] — Test Summary
NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
Test passed

¡Todo lo mejor!