NVRTC uses __host__ as the default execution space, and it generates an error if it encounters any host code in the input. That is, if the input contains entities with explicit __host__ annotations or no execution space annotation, NVRTC will emit an error. __host__ __device__ functions are treated as device functions.

NVRTC provides a compile option, --device-as-default-execution-space, that enables an alternative compilation mode, in which entities with no execution space annotations are treated as __device__ entities.

NVRTC itself does not provide any linker. Users can, however, use cuLinkAddData in the CUDA Driver API to link the generated relocatable PTX code with other relocatable code. To generate relocatable PTX code, the compile option --relocatable-device-code=true or --device-c is required.

Example: Dynamic Parallelism provides a simple example.

Different operating systems define integer type sizes differently. Linux x86_64 and Mac OS X implement LP64, and Windows x86_64 implements LLP64.

NVRTC implements LP64 on Linux and Mac OS X, and LLP64 on Windows.

Builtin functions provided by the CUDA Runtime headers when compiling offline with nvcc are available, unless otherwise noted.

NVRTC will mangle __global__ function names as specified by the IA64 ABI. If the generated PTX is being loaded using the CUDA Driver API, the kernel function must be looked up by name, but this is hard to do when the name has been mangled. To address this problem, NVRTC provides API functions that map source level __global__ function/template instantiation names to the mangled names present in the generated PTX.

The two API functions nvrtcAddNameExpression and nvrtcGetLoweredName work together to provide this functionality. First, a string denoting the source level 'name expression' for the __global__ function/template instantiation is provided to nvrtcAddNameExpression. Then, the program is compiled with nvrtcCompileProgram. During compilation, NVRTC will parse the name expression string as a C++ constant expression at the end of the user program. The constant expression must provide the address of a __global__ function/template instantiation. Finally, the function nvrtcGetLoweredName is called with the original name expression and it returns a pointer to the lowered name. The lowered name can be used to refer to the kernel in the CUDA Driver API.

NVRTC guarantees that any __global__ function/template instantiation referenced in a call to nvrtcAddNameExpression will be present in the generated PTX (if the definition is available in the input source code).

Example: Using Lowered Name lists a complete runnable example. Some relevant snippets:

1. The GPU source code ('gpu_program') contains definitions of various __global__ functions and function templates:

   const char *gpu_program = " \n\
   static __global__ void f1(int *result) { *result = 10; } \n\
   namespace N1 { \n\
   namespace N2 { \n\
   __global__ void f2(int *result) { *result = 20; } \n\
   } \n\
   } \n\
   template<typename T> \n\
   __global__ void f3(int *result) { *result = sizeof(T); } \n\
   \n";
   
   

2. The host source code invokes nvrtcAddNameExpression with various name expressions referring to __global__ functions and function template instantiations:

   name_vec.push_back("&f1");
   ..
   name_vec.push_back("N1::N2::f2");
   ..
   name_vec.push_back("f3<int>");
   ..
   name_vec.push_back("f3<double>");
   
   // add name expressions to NVRTC. Note this must be done before
   // the program is compiled.
   for (size_t i = 0; i < name_vec.size(); ++i)
   NVRTC_SAFE_CALL(nvrtcAddNameExpression(prog, name_vec[i].c_str()));
   
   

3. The GPU program is then compiled with nvrtcCompileProgram. The generated PTX is loaded on the GPU. The mangled names of the __global__ function and template instantiations are looked up:

   // note: this call must be made after NVRTC program has been
   // compiled and before it has been destroyed.
   NVRTC_SAFE_CALL(nvrtcGetLoweredName(
   prog,
   name_vec[i].c_str(), // name expression
   &name // lowered name
   ));
   
   

4. The mangled name is then used to launch the kernel using the CUDA Driver API:

   CUfunction kernel;
   CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, name));
   ...
   CUDA_SAFE_CALL(
   cuLaunchKernel(kernel,
   1, 1, 1, // grid dim
   1, 1, 1, // block dim
   0, NULL, // shared mem and stream
   args, 0));    
     

1. Sequence of calls: All name expressions must be added using nvrtcAddNameExpression before the NVRTC program is compiled with nvrtcCompileProgram. This is required because the name expressions are parsed at the end of the user program, and may trigger template instantiations. The lowered names must be looked up by calling nvrtcGetLoweredName only after the NVRTC program has been compiled, and before it has been destroyed. The pointer returned by nvrtcGetLoweredName points to memory owned by NVRTC, and this memory is freed when the NVRTC program has been destroyed (nvrtcDestroyProgram). Thus the correct sequence of calls is : nvrtcAddNameExpression, nvrtcCompileProgram, nvrtcGetLoweredName, nvrtcDestroyProgram.
2. Identical Name Expressions: The name expression string passed to nvrtcAddNameExpression and nvrtcGetLoweredName must have identical characters. For example, "foo" and "foo " are not identical strings, even though semantically they refer to the same entity (foo), because the second string has a extra whitespace character.
3. Constant Expressions: The characters in the name expression string are parsed as a C++ constant expression at the end of the user program. Any errors during parsing will cause compilation failure and compiler diagnostics will be generated in the compilation log. The constant expression must refer to the address of a __global__ function or function template instantiation.
4. Address of overloaded function: If the NVRTC source code has multiple overloaded __global__ functions, then the name expression must use a cast operation to disambiguate. However, casts are not allowed in constant expression for C++ dialects before C++11. If using such name expressions, please compile the code in C++11 or later dialect using the '-std' command line flag. Example: Consider that the GPU code string contains:

   __global__ void foo(int) { }
   __global__ void foo(char) { }
   

   The name expression '(void(*)(int))foo' correctly disambiguates 'foo(int)', but the program must be compiled in C++11 or later dialect (e.g. '-std=c++11') because casts are not allowed in pre-C++11 constant expressions.

In some scenarios, it is useful to instantiate __global__ function templates in device code based on template arguments in host code. The NVRTC helper function nvrtcGetTypeName can be used to extract the source level name of a type in host code, and this string can be used to instantiate a __global__ function template and get the mangled name of the instantiation using the nvrtcAddNameExpression and nvrtcGetLoweredName functions.

nvrtcGetTypeName is defined inline in the NVRTC header file, and is available when the macro NVRTC_GET_TYPE_NAME is defined with a non-zero value. It uses the abi::__cxa_demangle and UnDecorateSymbolName host code functions when using gcc/clang and cl.exe compilers, respectively. Users may need to specify additional header paths and libraries to find the host functions used (abi::__cxa_demangle / UnDecorateSymbolName). See the build instructions for the example below for reference (Build Instruction).

Example: Using nvrtcGetTypeName lists a complete runnable example. Some relevant snippets:

1. The GPU source code ('gpu_program') contains definitions of a __global__ function template:

   const char *gpu_program = " \n\
   namespace N1 { struct S1_t { int i; double d; }; } \n\
   template<typename T> \n\
   __global__ void f3(int *result) { *result = sizeof(T); } \n\
   \n";
   
   
   

2. The host code function getKernelNameForType creates the name expression for a __global__ function template instantiation based on the host template type T. The name of the type T is extracted using nvrtcGetTypeName:

   template <typename T>
   std::string getKernelNameForType(void)
   {
   // Look up the source level name string for the type "T" using
   // nvrtcGetTypeName() and use it to create the kernel name
   std::string type_name;
   NVRTC_SAFE_CALL(nvrtcGetTypeName<T>(&type_name));
   return std::string("f3<") + type_name + ">";
   }
   
   
   

3. The name expressions are presented to NVRTC using the nvrtcAddNameExpression function:

   name_vec.push_back(getKernelNameForType<int>());
   ..
   name_vec.push_back(getKernelNameForType<double>());
   ..
   name_vec.push_back(getKernelNameForType<N1::S1_t>());
   ..
   for (size_t i = 0; i < name_vec.size(); ++i)
   NVRTC_SAFE_CALL(nvrtcAddNameExpression(prog, name_vec[i].c_str()));
   
   
   

4. The GPU program is then compiled with nvrtcCompileProgram. The generated PTX is loaded on the GPU. The mangled names of the __global__ function template instantiations are looked up:

   // note: this call must be made after NVRTC program has been
   // compiled and before it has been destroyed.
   NVRTC_SAFE_CALL(nvrtcGetLoweredName(
   prog,
   name_vec[i].c_str(), // name expression
   &name // lowered name
   ));
   
   

5. The mangled name is then used to launch the kernel using the CUDA Driver API:

   CUfunction kernel;
   CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, name));
   ...
   CUDA_SAFE_CALL(
   cuLaunchKernel(kernel,
   1, 1, 1, // grid dim
   1, 1, 1, // block dim
   0, NULL, // shared mem and stream
   args, 0));    
     

#include <nvrtc.h>
#include <cuda.h>
#include <iostream>

#define NUM_THREADS 128
#define NUM_BLOCKS 32
#define NVRTC_SAFE_CALL(x)                                        \
  do {                                                            \
    nvrtcResult result = x;                                       \
    if (result != NVRTC_SUCCESS) {                                \
      std::cerr << "\nerror: " #x " failed with error "           \
                << nvrtcGetErrorString(result) << '\n';           \
      exit(1);                                                    \
    }                                                             \
  } while(0)
#define CUDA_SAFE_CALL(x)                                         \
  do {                                                            \
    CUresult result = x;                                          \
    if (result != CUDA_SUCCESS) {                                 \
      const char *msg;                                            \
      cuGetErrorName(result, &msg);                               \
      std::cerr << "\nerror: " #x " failed with error "           \
                << msg << '\n';                                   \
      exit(1);                                                    \
    }                                                             \
  } while(0)

const char *saxpy = "                                           \n\
extern \"C\" __global__                                         \n\
void saxpy(float a, float *x, float *y, float *out, size_t n)   \n\
{                                                               \n\
  size_t tid = blockIdx.x * blockDim.x + threadIdx.x;           \n\
  if (tid < n) {                                                \n\
    out[tid] = a * x[tid] + y[tid];                             \n\
  }                                                             \n\
}                                                               \n";

int main()
{
  // Create an instance of nvrtcProgram with the SAXPY code string.
  nvrtcProgram prog;

  NVRTC_SAFE_CALL(
    nvrtcCreateProgram(&prog,         // prog
                       saxpy,         // buffer
                       "saxpy.cu",    // name
                       0,             // numHeaders
                       NULL,          // headers
                       NULL));        // includeNames
  // Compile the program for compute_30 with fmad disabled.
  const char *opts[] = {"--gpu-architecture=compute_30",
                        "--fmad=false"};
  nvrtcResult compileResult = nvrtcCompileProgram(prog,  // prog
                                                  2,     // numOptions
                                                  opts); // options
  // Obtain compilation log from the program.
  size_t logSize;
  NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
  char *log = new char[logSize];
  NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
  std::cout << log << '\n';
  delete[] log;
  if (compileResult != NVRTC_SUCCESS) {
    exit(1);
  }
  // Obtain PTX from the program.
  size_t ptxSize;
  NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptxSize));
  char *ptx = new char[ptxSize];
  NVRTC_SAFE_CALL(nvrtcGetPTX(prog, ptx));
  // Destroy the program.
  NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));
  // Load the generated PTX and get a handle to the SAXPY kernel.
  CUdevice cuDevice;
  CUcontext context;
  CUmodule module;
  CUfunction kernel;
  CUDA_SAFE_CALL(cuInit(0));
  CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
  CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));
  CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, 0, 0));
  CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, "saxpy"));
  // Generate input for execution, and create output buffers.
  size_t n = NUM_THREADS * NUM_BLOCKS;
  size_t bufferSize = n * sizeof(float);
  float a = 5.1f;
  float *hX = new float[n], *hY = new float[n], *hOut = new float[n];
  for (size_t i = 0; i < n; ++i) {
    hX[i] = static_cast<float>(i);
    hY[i] = static_cast<float>(i * 2);
  }
  CUdeviceptr dX, dY, dOut;
  CUDA_SAFE_CALL(cuMemAlloc(&dX, bufferSize));
  CUDA_SAFE_CALL(cuMemAlloc(&dY, bufferSize));
  CUDA_SAFE_CALL(cuMemAlloc(&dOut, bufferSize));
  CUDA_SAFE_CALL(cuMemcpyHtoD(dX, hX, bufferSize));
  CUDA_SAFE_CALL(cuMemcpyHtoD(dY, hY, bufferSize));
  // Execute SAXPY.
  void *args[] = { &a, &dX, &dY, &dOut, &n };
  CUDA_SAFE_CALL(
    cuLaunchKernel(kernel,
                   NUM_BLOCKS, 1, 1,    // grid dim
                   NUM_THREADS, 1, 1,   // block dim
                   0, NULL,             // shared mem and stream
                   args, 0));           // arguments
  CUDA_SAFE_CALL(cuCtxSynchronize());
  // Retrieve and print output.
  CUDA_SAFE_CALL(cuMemcpyDtoH(hOut, dOut, bufferSize));

  for (size_t i = 0; i < n; ++i) {
    std::cout << a << " * " << hX[i] << " + " << hY[i]
              << " = " << hOut[i] << '\n';
  }
  // Release resources.
  CUDA_SAFE_CALL(cuMemFree(dX));
  CUDA_SAFE_CALL(cuMemFree(dY));
  CUDA_SAFE_CALL(cuMemFree(dOut));
  CUDA_SAFE_CALL(cuModuleUnload(module));
  CUDA_SAFE_CALL(cuCtxDestroy(context));
  delete[] hX;
  delete[] hY;
  delete[] hOut;
  return 0;
}

#include <nvrtc.h>
#include <cuda.h>
#include <iostream>
#include <vector>
#include <string>

#define NVRTC_SAFE_CALL(x)                                        \
  do {                                                            \
    nvrtcResult result = x;                                       \
    if (result != NVRTC_SUCCESS) {                                \
      std::cerr << "\nerror: " #x " failed with error "           \
                << nvrtcGetErrorString(result) << '\n';           \
      exit(1);                                                    \
    }                                                             \
  } while(0)
#define CUDA_SAFE_CALL(x)                                         \
  do {                                                            \
    CUresult result = x;                                          \
    if (result != CUDA_SUCCESS) {                                 \
      const char *msg;                                            \
      cuGetErrorName(result, &msg);                               \
      std::cerr << "\nerror: " #x " failed with error "           \
                << msg << '\n';                                   \
      exit(1);                                                    \
    }                                                             \
  } while(0)

const char *gpu_program = "                                     \n\
static __global__ void f1(int *result) { *result = 10; }        \n\
namespace N1 {                                                  \n\
  namespace N2 {                                                \n\
    __global__ void f2(int *result) { *result = 20; }           \n\
  }                                                             \n\
}                                                               \n\
template<typename T>                                            \n\
__global__ void f3(int *result) { *result = sizeof(T); }        \n\
                                                                \n";
int main()
{
  // Create an instance of nvrtcProgram
  nvrtcProgram prog;

  NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog,         // prog
                                     gpu_program,   // buffer
                                     "prog.cu",     // name
                                     0,             // numHeaders
                                     NULL,          // headers
                                     NULL));        // includeNames

  // add all name expressions for kernels
  std::vector<std::string> name_vec;
  std::vector<int> expected_result;
  
  // note the name expressions are parsed as constant expressions
  name_vec.push_back("&f1");
  expected_result.push_back(10);
  
  name_vec.push_back("N1::N2::f2");
  expected_result.push_back(20);
    
  name_vec.push_back("f3<int>");
  expected_result.push_back(sizeof(int));
  
  name_vec.push_back("f3<double>");
  expected_result.push_back(sizeof(double));
  
  // add name expressions to NVRTC. Note this must be done before
  // the program is compiled.
  for (size_t i = 0; i < name_vec.size(); ++i)
    NVRTC_SAFE_CALL(nvrtcAddNameExpression(prog, name_vec[i].c_str()));
  
  nvrtcResult compileResult = nvrtcCompileProgram(prog,  // prog
                                                  0,     // numOptions
                                                  NULL); // options
  // Obtain compilation log from the program.
  size_t logSize;
  NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
  char *log = new char[logSize];
  NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
  std::cout << log << '\n';
  delete[] log;
  if (compileResult != NVRTC_SUCCESS) {
    exit(1);
  }
  // Obtain PTX from the program.
  size_t ptxSize;
  NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptxSize));
  char *ptx = new char[ptxSize];
  NVRTC_SAFE_CALL(nvrtcGetPTX(prog, ptx));

  // Load the generated PTX
  CUdevice cuDevice;
  CUcontext context;
  CUmodule module;
  
  CUDA_SAFE_CALL(cuInit(0));
  CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
  CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));
  CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, 0, 0));
  
  CUdeviceptr dResult;
  int hResult = 0;
  CUDA_SAFE_CALL(cuMemAlloc(&dResult, sizeof(hResult)));
  CUDA_SAFE_CALL(cuMemcpyHtoD(dResult, &hResult, sizeof(hResult)));
  

  // for each of the name expressions previously provided to NVRTC,
  // extract the lowered name for corresponding __global__ function,
  // and launch it.
  
  for (size_t i = 0; i < name_vec.size(); ++i) {
    const char *name;
    
    // note: this call must be made after NVRTC program has been 
    // compiled and before it has been destroyed.
    NVRTC_SAFE_CALL(nvrtcGetLoweredName(
                          prog, 
			  name_vec[i].c_str(), // name expression
			  &name                // lowered name
                                        ));
    
    // get pointer to kernel from loaded PTX
    CUfunction kernel;
    CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, name));
    
    // launch the kernel
    std::cout << "\nlaunching " << name << " ("
	      << name_vec[i] << ")" << std::endl;
    
    void *args[] = { &dResult };
    CUDA_SAFE_CALL(
      cuLaunchKernel(kernel,
		     1, 1, 1,             // grid dim
		     1, 1, 1,             // block dim
		     0, NULL,             // shared mem and stream
		     args, 0));           // arguments
    CUDA_SAFE_CALL(cuCtxSynchronize());
    
    // Retrieve the result
    CUDA_SAFE_CALL(cuMemcpyDtoH(&hResult, dResult, sizeof(hResult)));
    
    // check against expected value
    if (expected_result[i] != hResult) {
      std::cout << "\n Error: expected result = " << expected_result[i]
		<< " , actual result = " << hResult << std::endl;
      exit(1);
    }
  }  // for
    
  // Release resources.
  CUDA_SAFE_CALL(cuMemFree(dResult));
  CUDA_SAFE_CALL(cuModuleUnload(module));
  CUDA_SAFE_CALL(cuCtxDestroy(context));
  
  // Destroy the program. 
  NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));

  return 0;
}

#include <nvrtc.h>
#include <cuda.h>
#include <iostream>
#include <vector>
#include <string>

#define NVRTC_SAFE_CALL(x)                                        \
  do {                                                            \
    nvrtcResult result = x;                                       \
    if (result != NVRTC_SUCCESS) {                                \
      std::cerr << "\nerror: " #x " failed with error "           \
                << nvrtcGetErrorString(result) << '\n';           \
      exit(1);                                                    \
    }                                                             \
  } while(0)
#define CUDA_SAFE_CALL(x)                                         \
  do {                                                            \
    CUresult result = x;                                          \
    if (result != CUDA_SUCCESS) {                                 \
      const char *msg;                                            \
      cuGetErrorName(result, &msg);                               \
      std::cerr << "\nerror: " #x " failed with error "           \
                << msg << '\n';                                   \
      exit(1);                                                    \
    }                                                             \
  } while(0)

const char *gpu_program = "                                     \n\
namespace N1 { struct S1_t { int i; double d; }; }              \n\
template<typename T>                                            \n\
__global__ void f3(int *result) { *result = sizeof(T); }        \n\
                                                                \n";


// note: this structure is also defined in GPU code string. Should ideally
// be in a header file included by both GPU code string and by CPU code.
namespace N1 { struct S1_t { int i; double d; }; };

template <typename T>
std::string getKernelNameForType(void)
{
   // Look up the source level name string for the type "T" using
   // nvrtcGetTypeName() and use it to create the kernel name 
   std::string type_name;
   NVRTC_SAFE_CALL(nvrtcGetTypeName<T>(&type_name));
   return std::string("f3<") + type_name + ">";
}

int main()
{
  // Create an instance of nvrtcProgram
  nvrtcProgram prog;
  NVRTC_SAFE_CALL(
    nvrtcCreateProgram(&prog,         // prog
                       gpu_program,   // buffer
                       "gpu_program.cu",    // name
                       0,             // numHeaders
                       NULL,          // headers
                       NULL));        // includeNames

  // add all name expressions for kernels
  std::vector<std::string> name_vec;
  std::vector<int> expected_result;
  
  // note the name expressions are parsed as constant expressions
  name_vec.push_back(getKernelNameForType<int>());
  expected_result.push_back(sizeof(int));
    
  name_vec.push_back(getKernelNameForType<double>());
  expected_result.push_back(sizeof(double));
  
  name_vec.push_back(getKernelNameForType<N1::S1_t>());
  expected_result.push_back(sizeof(N1::S1_t));
  
  
  // add name expressions to NVRTC. Note this must be done before
  // the program is compiled.
  for (size_t i = 0; i < name_vec.size(); ++i)
    NVRTC_SAFE_CALL(nvrtcAddNameExpression(prog, name_vec[i].c_str()));
  
  nvrtcResult compileResult = nvrtcCompileProgram(prog,  // prog
                                                  0,     // numOptions
                                                  NULL); // options
  // Obtain compilation log from the program.
  size_t logSize;
  NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
  char *log = new char[logSize];
  NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
  std::cout << log << '\n';
  delete[] log;
  if (compileResult != NVRTC_SUCCESS) {
    exit(1);
  }
  // Obtain PTX from the program.
  size_t ptxSize;
  NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptxSize));
  char *ptx = new char[ptxSize];
  NVRTC_SAFE_CALL(nvrtcGetPTX(prog, ptx));
  

  // Load the generated PTX
  CUdevice cuDevice;
  CUcontext context;
  CUmodule module;
  
  CUDA_SAFE_CALL(cuInit(0));
  CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
  CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));
  CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, 0, 0));
  
  CUdeviceptr dResult;
  int hResult = 0;
  CUDA_SAFE_CALL(cuMemAlloc(&dResult, sizeof(hResult)));
  CUDA_SAFE_CALL(cuMemcpyHtoD(dResult, &hResult, sizeof(hResult)));
  
  // for each of the name expressions previously provided to NVRTC,
  // extract the lowered name for corresponding __global__ function,
  // and launch it.
  
  for (size_t i = 0; i < name_vec.size(); ++i) {
    const char *name;
    
    // note: this call must be made after NVRTC program has been 
    // compiled and before it has been destroyed.
    NVRTC_SAFE_CALL(nvrtcGetLoweredName(
                          prog, 
			  name_vec[i].c_str(), // name expression
			  &name                // lowered name
                                        ));
    
    // get pointer to kernel from loaded PTX
    CUfunction kernel;
    CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, name));
    
    // launch the kernel
    std::cout << "\nlaunching " << name << " ("
	      << name_vec[i] << ")" << std::endl;
    
    void *args[] = { &dResult };
    CUDA_SAFE_CALL(
      cuLaunchKernel(kernel,
		     1, 1, 1,             // grid dim
		     1, 1, 1,             // block dim
		     0, NULL,             // shared mem and stream
		     args, 0));           // arguments
    CUDA_SAFE_CALL(cuCtxSynchronize());
    
    // Retrieve the result
    CUDA_SAFE_CALL(cuMemcpyDtoH(&hResult, dResult, sizeof(hResult)));
    
    // check against expected value
    if (expected_result[i] != hResult) {
      std::cout << "\n Error: expected result = " << expected_result[i]
		<< " , actual result = " << hResult << std::endl;
      exit(1);
    }
  }  // for
    
  // Release resources.
  CUDA_SAFE_CALL(cuMemFree(dResult));
  CUDA_SAFE_CALL(cuModuleUnload(module));
  CUDA_SAFE_CALL(cuCtxDestroy(context));
  
  // Destroy the program. 
  NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));

  return 0;
}

#include <nvrtc.h>
#include <cuda.h>
#include <iostream>

#define NVRTC_SAFE_CALL(x)                                        \
  do {                                                            \
    nvrtcResult result = x;                                       \
    if (result != NVRTC_SUCCESS) {                                \
      std::cerr << "\nerror: " #x " failed with error "           \
                << nvrtcGetErrorString(result) << '\n';           \
      exit(1);                                                    \
    }                                                             \
  } while(0)
#define CUDA_SAFE_CALL(x)                                         \
  do {                                                            \
    CUresult result = x;                                          \
    if (result != CUDA_SUCCESS) {                                 \
      const char *msg;                                            \
      cuGetErrorName(result, &msg);                               \
      std::cerr << "\nerror: " #x " failed with error "           \
                << msg << '\n';                                   \
      exit(1);                                                    \
    }                                                             \
  } while(0)

const char *dynamic_parallelism = "                             \n\
extern \"C\" __global__                                         \n\
void child(float *out, size_t n)                                \n\
{                                                               \n\
  size_t tid = blockIdx.x * blockDim.x + threadIdx.x;           \n\
  if (tid < n) {                                                \n\
    out[tid] = tid;                                             \n\
  }                                                             \n\
}                                                               \n\
                                                                \n\
extern \"C\" __global__                                         \n\
void parent(float *out, size_t n,                               \n\
              size_t numBlocks, size_t numThreads)              \n\
{                                                               \n\
  child<<<numBlocks, numThreads>>>(out, n);                     \n\
  cudaDeviceSynchronize();                                      \n\
}                                                               \n";

int main(int argc, char *argv[])
{
  if (argc < 2) {
    std::cout << "Usage: dynamic-parallelism <path to cudadevrt library>\n\n"
              << "<path to cudadevrt library> must include the cudadevrt\n"
              << "library name itself, e.g., Z:\\path\\to\\cudadevrt.lib on \n"
              << "Windows and /path/to/libcudadevrt.a on Linux and Mac OS X.\n";
    exit(1);
  }
  size_t numBlocks = 32;
  size_t numThreads = 128;
  // Create an instance of nvrtcProgram with the code string.
  nvrtcProgram prog;
  NVRTC_SAFE_CALL(
    nvrtcCreateProgram(&prog,                       // prog
                       dynamic_parallelism,         // buffer
                       "dynamic_parallelism.cu",    // name
                       0,                           // numHeaders
                       NULL,                        // headers
                       NULL));                      // includeNames
  // Compile the program for compute_35 with rdc enabled.
  const char *opts[] = {"--gpu-architecture=compute_35",
                        "--relocatable-device-code=true"};
  nvrtcResult compileResult = nvrtcCompileProgram(prog,  // prog
                                                  2,     // numOptions
                                                  opts); // options
  // Obtain compilation log from the program.
  size_t logSize;
  NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
  char *log = new char[logSize];
  NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
  std::cout << log << '\n';
  delete[] log;
  if (compileResult != NVRTC_SUCCESS) {
    exit(1);
  }
  // Obtain PTX from the program.
  size_t ptxSize;
  NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptxSize));
  char *ptx = new char[ptxSize];
  NVRTC_SAFE_CALL(nvrtcGetPTX(prog, ptx));
  // Destroy the program.
  NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));
  // Load the generated PTX and get a handle to the parent kernel.
  CUdevice cuDevice;
  CUcontext context;
  CUlinkState linkState;
  CUmodule module;
  CUfunction kernel;
  CUDA_SAFE_CALL(cuInit(0));
  CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
  CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));
  CUDA_SAFE_CALL(cuLinkCreate(0, 0, 0, &linkState));
  CUDA_SAFE_CALL(cuLinkAddFile(linkState, CU_JIT_INPUT_LIBRARY, argv[1],
                               0, 0, 0));
  CUDA_SAFE_CALL(cuLinkAddData(linkState, CU_JIT_INPUT_PTX,
                               (void *)ptx, ptxSize, "dynamic_parallelism.ptx",
                               0, 0, 0));
  size_t cubinSize;
  void *cubin;
  CUDA_SAFE_CALL(cuLinkComplete(linkState, &cubin, &cubinSize));
  CUDA_SAFE_CALL(cuModuleLoadData(&module, cubin));
  CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, "parent"));

  // Generate input for execution, and create output buffers.
  size_t n = numBlocks * numThreads;
  size_t bufferSize = n * sizeof(float);
  float *hOut = new float[n];
  CUdeviceptr dX, dY, dOut;
  CUDA_SAFE_CALL(cuMemAlloc(&dOut, bufferSize));
  // Execute parent kernel.
  void *args[] = { &dOut, &n, &numBlocks, &numThreads };
  CUDA_SAFE_CALL(
    cuLaunchKernel(kernel,
                   1, 1, 1,    // grid dim
                   1, 1, 1,    // block dim
                   0, NULL,    // shared mem and stream
                   args, 0));  // arguments
  CUDA_SAFE_CALL(cuCtxSynchronize());
  // Retrieve and print output.
  CUDA_SAFE_CALL(cuMemcpyDtoH(hOut, dOut, bufferSize));

  for (size_t i = 0; i < n; ++i) {
    std::cout << hOut[i] << '\n';
  }
  // Release resources.
  CUDA_SAFE_CALL(cuMemFree(dOut));
  CUDA_SAFE_CALL(cuModuleUnload(module));
  CUDA_SAFE_CALL(cuLinkDestroy(linkState));
  CUDA_SAFE_CALL(cuCtxDestroy(context));
  delete[] hOut;
  return 0;
}