| 18 | |
| 19 | namespace { |
| 20 | std::string NVRTCCompile(const std::string& code, int cap_major, int cap_minor) { |
| 21 | static std::vector<std::string> cuda_include_opts = get_cuda_include_opts(); |
| 22 | |
| 23 | auto arch_opt = ssprintf("--gpu-architecture=compute_%d%d", cap_major, cap_minor); |
| 24 | std::vector<const char*> opts; |
| 25 | opts.push_back(arch_opt.c_str()); |
| 26 | for (auto& inc_path : cuda_include_opts) |
| 27 | opts.push_back(inc_path.c_str()); |
| 28 | nvrtcProgram prog; |
| 29 | MGB_NVRTC_CHECK( |
| 30 | nvrtcCreateProgram(&prog, code.c_str(), nullptr, 0, nullptr, nullptr)); |
| 31 | std::unique_ptr<nvrtcProgram, void (*)(nvrtcProgram*)> prog_release{ |
| 32 | &prog, [](nvrtcProgram* p) { MGB_NVRTC_CHECK(nvrtcDestroyProgram(p)); }}; |
| 33 | nvrtcResult compile_res = nvrtcCompileProgram(prog, opts.size(), opts.data()); |
| 34 | size_t log_size; |
| 35 | MGB_NVRTC_CHECK(nvrtcGetProgramLogSize(prog, &log_size)); |
| 36 | std::string log; |
| 37 | log.resize(log_size); |
| 38 | MGB_NVRTC_CHECK(nvrtcGetProgramLog(prog, &log[0])); |
| 39 | mgb_throw_if( |
| 40 | compile_res != NVRTC_SUCCESS, SystemError, |
| 41 | "nvrtc compile error: %s\n========= source code\n%s", log.c_str(), |
| 42 | code.c_str()); |
| 43 | size_t ptx_size; |
| 44 | MGB_NVRTC_CHECK(nvrtcGetPTXSize(prog, &ptx_size)); |
| 45 | std::string ptx; |
| 46 | ptx.resize(ptx_size); |
| 47 | MGB_NVRTC_CHECK(nvrtcGetPTX(prog, &ptx[0])); |
| 48 | return ptx; |
| 49 | } |
| 50 | |
| 51 | void make_fastdiv(Uint32Fastdiv& fdiv, uint32_t d) { |
| 52 | mgb_assert(d); |