| 2732 | } |
| 2733 | |
| 2734 | inline void instantiate_kernel( |
| 2735 | std::string const& program_name, |
| 2736 | std::map<std::string, std::string> const& program_sources, |
| 2737 | std::string const& instantiation, std::vector<std::string> const& options, |
| 2738 | std::string* log, std::string* ptx, std::string* mangled_instantiation, |
| 2739 | std::vector<std::string>* linker_files, |
| 2740 | std::vector<std::string>* linker_paths) { |
| 2741 | std::vector<std::string> compiler_options; |
| 2742 | detail::split_compiler_and_linker_options(options, &compiler_options, |
| 2743 | linker_files, linker_paths); |
| 2744 | |
| 2745 | std::cout << "ABout to compile kernel" << std::endl; |
| 2746 | nvrtcResult ret = |
| 2747 | detail::compile_kernel(program_name, program_sources, compiler_options, |
| 2748 | instantiation, log, ptx, mangled_instantiation); |
| 2749 | #if JITIFY_PRINT_LOG |
| 2750 | if (log->size() > 1) { |
| 2751 | detail::print_compile_log(program_name, *log); |
| 2752 | } |
| 2753 | #endif |
| 2754 | if (ret != NVRTC_SUCCESS) { |
| 2755 | throw std::runtime_error(std::string("NVRTC error: ") + |
| 2756 | nvrtcGetErrorString(ret)); |
| 2757 | } |
| 2758 | std::cout << "done compilling" << std::endl; |
| 2759 | |
| 2760 | #if JITIFY_PRINT_PTX |
| 2761 | std::cout << "---------------------------------------" << std::endl; |
| 2762 | std::cout << *mangled_instantiation << std::endl; |
| 2763 | std::cout << "---------------------------------------" << std::endl; |
| 2764 | std::cout << "--- PTX for " << mangled_instantiation << " in " << program_name |
| 2765 | << " ---" << std::endl; |
| 2766 | std::cout << "---------------------------------------" << std::endl; |
| 2767 | std::cout << *ptx << std::endl; |
| 2768 | std::cout << "---------------------------------------" << std::endl; |
| 2769 | #endif |
| 2770 | } |
| 2771 | |
| 2772 | inline void get_1d_max_occupancy(CUfunction func, |
| 2773 | CUoccupancyB2DSize smem_callback, |
no test coverage detected