| 2478 | } |
| 2479 | |
| 2480 | inline nvrtcResult compile_kernel(std::string program_name, |
| 2481 | std::map<std::string, std::string> sources, |
| 2482 | std::vector<std::string> options, |
| 2483 | std::string instantiation = "", |
| 2484 | std::string* log = 0, std::string* ptx = 0, |
| 2485 | std::string* mangled_instantiation = 0) { |
| 2486 | std::string program_source = sources[program_name]; |
| 2487 | // Build arrays of header names and sources |
| 2488 | std::vector<const char*> header_names_c; |
| 2489 | std::vector<const char*> header_sources_c; |
| 2490 | int num_headers = (int)(sources.size() - 1); |
| 2491 | header_names_c.reserve(num_headers); |
| 2492 | header_sources_c.reserve(num_headers); |
| 2493 | typedef std::map<std::string, std::string> source_map; |
| 2494 | for (source_map::const_iterator iter = sources.begin(); iter != sources.end(); |
| 2495 | ++iter) { |
| 2496 | std::string const& name = iter->first; |
| 2497 | std::string const& code = iter->second; |
| 2498 | if (name == program_name) { |
| 2499 | continue; |
| 2500 | } |
| 2501 | header_names_c.push_back(name.c_str()); |
| 2502 | header_sources_c.push_back(code.c_str()); |
| 2503 | } |
| 2504 | |
| 2505 | // TODO: This WAR is expected to be unnecessary as of CUDA > 10.2. |
| 2506 | bool should_remove_unused_globals = |
| 2507 | detail::pop_remove_unused_globals_flag(&options); |
| 2508 | |
| 2509 | std::vector<const char*> options_c(options.size() + 2); |
| 2510 | options_c[0] = "--device-as-default-execution-space"; |
| 2511 | options_c[1] = "--pre-include=jitify_preinclude.h"; |
| 2512 | for (int i = 0; i < (int)options.size(); ++i) { |
| 2513 | options_c[i + 2] = options[i].c_str(); |
| 2514 | } |
| 2515 | |
| 2516 | #if CUDA_VERSION < 8000 |
| 2517 | std::string inst_dummy; |
| 2518 | if (!instantiation.empty()) { |
| 2519 | // WAR for no nvrtcAddNameExpression before CUDA 8.0 |
| 2520 | // Force template instantiation by adding dummy reference to kernel |
| 2521 | inst_dummy = "__jitify_instantiation"; |
| 2522 | program_source += |
| 2523 | "\nvoid* " + inst_dummy + " = (void*)" + instantiation + ";\n"; |
| 2524 | } |
| 2525 | #endif |
| 2526 | |
| 2527 | #define CHECK_NVRTC(call) \ |
| 2528 | do { \ |
| 2529 | nvrtcResult ret = call; \ |
| 2530 | if (ret != NVRTC_SUCCESS) { \ |
| 2531 | return ret; \ |
| 2532 | } \ |
| 2533 | } while (0) |
| 2534 | |
| 2535 | nvrtcProgram nvrtc_program; |
| 2536 | CHECK_NVRTC(nvrtcCreateProgram( |
| 2537 | &nvrtc_program, program_source.c_str(), program_name.c_str(), num_headers, |
no test coverage detected