Fichier BlackScholes_kernel.cuh

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

1. Trouvez l'assistant de périphérique cuda - il est stocké dans un helper_functions.h. Cette fonction vérifie simplement si un GPU NVIDIA est disponible dans la machine.

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)

Le fichier se compose du même code C ++ que dans le fichier .cpp ci-dessus avec un léger code de syntaxe supplémentaire.

1. La fonction __global__ est également appelée fonction «noyau». C'est la fonction que vous pouvez appeler du côté hôte en utilisant la sémantique d'appel du noyau CUDA ( <<<...>>>). Dans notre cas, il est défini dans le fichier ci-dessus une fonction BlackScholesGPU.

__global__ void BlackScholesGPU

__launch_bounds__(128)

__device__ inline void BlackScholesBodyGPU

En ce qui concerne le point de départ, j'ai laissé les données d'origine configurées par les développeurs 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

Tous mes vœux!