Getting calling a __host__ function from a global function is not allowed when using cudaMallocManaged

Question

I have a written code that I am trying to modify in order to make it use CUDA and I am having plenty of troubles, currently, I was trying to make the functions I want to be kernel functions to be void and I got some errors

Here is the list of errors I am getting:

black_scholes.cu(54): error: calling a __host__ function("cudaMallocManaged<double> ") from a __global__ function("black_scholes_iterate") is not allowed

black_scholes.cu(54): error: identifier "cudaMallocManaged<double> " is undefined in device code

black_scholes.cu(56): error: calling a __host__ function("init_gaussrand_state") from a __global__ function("black_scholes_iterate") is not allowed

black_scholes.cu(56): error: identifier "init_gaussrand_state" is undefined in device code

black_scholes.cu(65): error: calling a __host__ function("spawn_prng_stream") from a __global__ function("black_scholes_iterate") is not allowed

black_scholes.cu(65): error: identifier "spawn_prng_stream" is undefined in device code

black_scholes.cu(66): error: calling a __host__ function("gaussrand1") from a __global__ function("black_scholes_iterate") is not allowed

black_scholes.cu(66): error: identifier "gaussrand1" is undefined in device code

black_scholes.cu(66): error: identifier "uniform_random_double" is undefined in device code

black_scholes.cu(73): error: calling a __host__ function("free_prng_stream") from a __global__ function("black_scholes_iterate") is not allowed

black_scholes.cu(73): error: identifier "free_prng_stream" is undefined in device code

black_scholes.cu(74): error: calling a __host__ function("cudaFree") from a __global__ function("black_scholes_iterate") is not allowed

black_scholes.cu(74): error: identifier "cudaFree" is undefined in device code

I am particularly posting concerning the first 2 errors as while learning CUDA via an Nvidia Introductory course, it was common to call cudaMallocManaged inside a __global__ function and I don't get what is different here

Here is my .cu code :

#include "black_scholes.h"
#include "gaussian.h"
#include "random.h" 
#include "util.h"
#include <assert.h>
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

__managed__ double stddev;

__global__ void black_scholes_stddev (void* the_args)
{

  black_scholes_args_t* args = (black_scholes_args_t*) the_args;
  const double mean = args->mean;
  const int M = args->M;
  double variance = 0.0;
  int k = blockIdx.x * blockDim.x + threadIdx.x;

  if(k<M)
  {
   const double diff = args->trials[k] - mean;
   variance += diff * diff / (double) M;
  }

  args->variance = variance;
  stddev=sqrt(variance);

}


__global__ void black_scholes_iterate (void* the_args)
{

  black_scholes_args_t* args = (black_scholes_args_t*) the_args;

  const int S = args->S;
  const int E = args->E;
  const int M = args->M;
  const double r = args->r;
  const double sigma = args->sigma;
  const double T = args->T;

  double* trials = args->trials;
  double mean = 0.0;

  gaussrand_state_t gaussrand_state;
  void* prng_stream = NULL; 

double *randnumbs;
cudaMallocManaged(&randnumbs, M * sizeof (double));

init_gaussrand_state (&gaussrand_state);

int i = blockIdx.x * blockDim.x + threadIdx.x;
int k = blockIdx.x * blockDim.x + threadIdx.x;


//for (int i = 0; i < M; i++)
if(i<M)
{
  prng_stream = spawn_prng_stream(i%4);
  const double gaussian_random_number = gaussrand1 (&uniform_random_double, prng_stream, &gaussrand_state);
  randnumbs[i]=gaussian_random_number;
  const double current_value = S * exp ( (r - (sigma*sigma) / 2.0) * T + sigma * sqrt (T) * randnumbs[k]);
  trials[k] = exp (-r * T) * ((current_value - E < 0.0) ? 0.0 : current_value - E);
   mean += trials[k] / (double) M;//needs to be shared
  args->mean = mean;
}
  free_prng_stream (prng_stream);
  cudaFree(randnumbs);
}



void black_scholes (confidence_interval_t* interval,
           const double S,
           const double E,
           const double r,
           const double sigma,
           const double T,
           const int M,
         const int n)
{
  black_scholes_args_t args;
  double mean = 0.0;
  double conf_width = 0.0;
  double* trials = NULL;

  assert (M > 0);
  trials = (double*) malloc (M * sizeof (double));
  assert (trials != NULL);

  args.S = S;
  args.E = E;
  args.r = r;
  args.sigma = sigma;
  args.T = T;
  args.M = M;
  args.trials = trials;
  args.mean = 0.0;
  args.variance = 0.0;

  (void)black_scholes_iterate<<<1,1>>>(&args);
  mean = args.mean;
  black_scholes_stddev<<<1,1>>> (&args);
  cudaDeviceSynchronize();

  conf_width = 1.96 * stddev / sqrt ((double) M);
  interval->min = mean - conf_width;
  interval->max = mean + conf_width;

  deinit_black_scholes_args (&args);
}


void deinit_black_scholes_args (black_scholes_args_t* args)
{
  if (args != NULL)
    if (args->trials != NULL)
      {
    free (args->trials);
    args->trials = NULL;
      }
}

Any help in understanding what is going on would be appreciated, it seems to be a recurrent theme.

Answer 1

Currently, it's not possible to call cudaMallocManaged in CUDA device code. It has never been possible. I don't believe there is NVIDIA training material that demonstrates using cudaMallocManaged in device code.

If you wish to make an in-kernel allocation, I suggest using the methods described in the programming guide. Also, new and delete work similarly to malloc() and free(), for in-kernel usage.