Last active
March 2, 2020 17:03
-
-
Save Robadob/b2b7704a36e2f679942e854a0f41082a to your computer and use it in GitHub Desktop.
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
/** | |
* 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.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
This is probably worth raising as a jitify issue next week once we have had a chat about it.