NVRTC_link_with_executing_module.cu

/**
 * Build this with -rdc=true
 * You will need to link against : nvrtc, cuda, cudart_static, dl
 */
#include <string>
#include <vector>
#include <memory>
#include <cassert>
#include <cuda_runtime.h>
#include <nvrtc.h>
#include <cctype>  // Required to build Jitify header under VS2015
#include "jitify/jitify.hpp"

__device__ int TEST_DEVICE_SYMBOL;
__constant__ int TEST_CONSTANT_SYMBOL;

int main() {
    const char* rtc_kernel = R"###(my_program
// Extern as these symbols are defined elsewhere
extern __device__ int TEST_DEVICE_SYMBOL;
extern __constant__ int TEST_CONSTANT_SYMBOL;
__global__ void simple_test(const int * const TEST_DEVICE_POINTER) {
    printf("TEST_DEVICE_POINTER: %s\n",  (*TEST_DEVICE_POINTER == 12) ? "Success" : "Failure");
    printf("TEST_DEVICE_SYMBOL: %s\n",   (TEST_DEVICE_SYMBOL == 12)   ? "Success" : "Failure");
    printf("TEST_CONSTANT_SYMBOL: %s\n", (TEST_CONSTANT_SYMBOL == 12) ? "Success" : "Failure");
})###";

    // Set TEST_SYMBOL
    const int t_TEST_SYMBOL = 12;
    assert(cudaMemcpyToSymbol(TEST_DEVICE_SYMBOL, &t_TEST_SYMBOL, sizeof(int)) == CUDA_SUCCESS);
    assert(cudaMemcpyToSymbol(TEST_CONSTANT_SYMBOL, &t_TEST_SYMBOL, sizeof(int)) == CUDA_SUCCESS);
    int *d_TEST_DEVICE_POINTER = nullptr;
    assert(cudaMalloc(&d_TEST_DEVICE_POINTER, sizeof(int)) == CUDA_SUCCESS);
    assert(cudaMemcpy(d_TEST_DEVICE_POINTER, &t_TEST_SYMBOL, sizeof(int), cudaMemcpyHostToDevice) == CUDA_SUCCESS);

    // Build RTC kernel
    std::vector<std::string> options;
    // cuda path
    std::string include_cuda;
    include_cuda = "-I" + std::string(std::getenv("CUDA_PATH")) + "\\include";
    options.push_back(include_cuda);
    // rdc
    std::string rdc;
    rdc = "-rdc=true";
    options.push_back(rdc);
#ifndef EXPERIMENTAL
    static jitify::JitCache jit_cache;
    auto program = std::make_shared<jitify::Program>(jit_cache, rtc_kernel, 0, options);
#else
    auto program = std::make_shared<jitify::experimental::Program>(rtc_kernel, std::vector<std::string>(), options);
#endif
    // Launch RTC kernel
    dim3 grid(1);
    dim3 block(1);
    auto kernelLauncher = program->kernel("simple_test").instantiate().configure(grid, block);
    assert(kernelLauncher.launch({ reinterpret_cast<void*>(&d_TEST_DEVICE_POINTER) }) == CUresult::CUDA_SUCCESS);

    // Cleanup
    assert(cudaDeviceSynchronize() == CUDA_SUCCESS);
    assert(cudaFree(d_TEST_DEVICE_POINTER) == CUDA_SUCCESS);

#ifdef _WIN32
    // Prevent window closing as probably launched via visual studio debugger
    printf("Hit enter to exit.\n");
    getchar();
#endif
    return EXIT_SUCCESS;
}

Try adding

cudaFree(0);
cuInit(0);
cuDeviceGet(&cuDevice, 0);
CUcontext context;
cuDevicePrimaryCtxRetain(&context, cuDevice);

before any jitify stuff.