ErrorNoBinaryForGPU when doing Black Sholes

Oct 5, 2012 at 7:03 PM

Strange issue I've run into:

CUDA.NET exception: ErrorNoBinaryForGPU (Ensure that compiled architecture version is suitable for device).

From either this (the double way to do this)

dCND = (

float)(1.0 - 1.0 / Math.Sqrt(2 * Math.PI) * Math.Exp(-L * L / 2.0) * (a1 * K + a2 * K * K + a3 * Math .Pow(K, 3.0) + a4 *Math.Pow(K, 4.0) + a5 * Math .Pow(K, 5.0)));

And This code snippit (which is converted to single

dCND = 1.0f - 1.0f /

 

It looks like it's when I use GMath.Exp or Pow

Any ideas?

GMath.Sqrt(2f * GMath.PI) * GMath.Exp(-L * L / 2.0f) * (a1 * K + a2 * K * K + a3 * GMath.Pow(K, 3.0f) );

Oct 5, 2012 at 7:25 PM

Actually just this causes the error in both Cudafy 1.9 and 1.10

dCND =

GMath .Pow(K, 3.0f);

Coordinator
Oct 6, 2012 at 5:32 PM

Not able to reproduce this, Dan. What architecture are you specifying in your Cudafy call and what type of GPU do you have?

Oct 8, 2012 at 6:22 PM

It's probably my GPU, I'll check it out on a Fermi shortly.  I'm using 1.2 architecture on NVS-3100M (yea embarrassing).

Coordinator
Oct 8, 2012 at 6:29 PM

1.2 does not support double floating point.  When you use doubles in your code for 1.2 then they should automagically revert to single.  It could be that Exp and Pow do not do this.  This could be a CUDA issue or the cudafy translation does not work right for these.  Can you post your generated .cu code?

Oct 8, 2012 at 8:25 PM
Edited Oct 9, 2012 at 1:52 PM

For anyone reading, the CND function is not true CND, it's abbreviated to show this problem. 

 

struct BlackScholesParameters
{
 __device__ BlackScholesParameters()
 {
 }
 float StockPrice;
 float StrikePrice;
 float TimeToExpiration;
 float RiskFreeInterestRate;
 float Volatility;
 float CallBlackScholes;
 float PutBlackScholes;
 __device__ BlackScholesParameters(float stockPrice, float strikePrice, float timeToExpiration, float riskFreeInterestRate, float volatility)
 {
  StockPrice = stockPrice;
  StrikePrice = strikePrice;
  TimeToExpiration = timeToExpiration;
  RiskFreeInterestRate = riskFreeInterestRate;
  Volatility = volatility;
  CallBlackScholes = 0;
  PutBlackScholes = 0;
 }
};


// Calculations.BlackSholesGpu
__device__ float CND(float X);
// Calculations.BlackSholesGpu
__device__ float BlackScholesGpu(int isCall, float S, float X, float T, float r, float v);
// Calculations.BlackSholesGpu
extern "C" __global__ void GpuExecute(BlackScholesParameters* options, int optionsLen0);

// Calculations.BlackSholesGpu
__device__ float CND(float X)
{
 float num = fabsf(X);
 float x = 1 / (1 + 0.2316419 * num);
 float num2 = powf(x, 3);
 float result;
 if (X < 0)
 {
  result = 1 - num2;
 }
 else
 {
  result = num2;
 }
 return result;
}
// Calculations.BlackSholesGpu
__device__ float BlackScholesGpu(int isCall, float S, float X, float T, float r, float v)
{
 float result = 0;
 float num = (logf(S / X) + (r + v * v / 2) * T) / (v * sqrtf(T));
 float num2 = num - v * sqrtf(T);
 if (isCall == 1)
 {
  result = S * CND(num) - X * expf(-r * T) * CND(num2);
 }
 else
 {
  if (isCall == 0)
  {
   result = X * expf(-r * T) * CND(-num2) - S * CND(-num);
  }
 }
 return result;
}
// Calculations.BlackSholesGpu
extern "C" __global__ void GpuExecute(BlackScholesParameters* options, int optionsLen0)
{
 int num = blockIdx.x * blockDim.x + threadIdx.x;
 if (optionsLen0 < num)
 {
  BlackScholesGpu(1, options[(num)].StockPrice, options[(num)].StrikePrice, options[(num)].TimeToExpiration, options[(num)].RiskFreeInterestRate, options[(num)].Volatility);
  BlackScholesGpu(0, options[(num)].StockPrice, options[(num)].StrikePrice, options[(num)].TimeToExpiration, options[(num)].RiskFreeInterestRate, options[(num)].Volatility);
 }
}