create a map of __constant__ and __device__ variables in the ptx file mapping demangled to mangled name
| 1212 | // create a map of __constant__ and __device__ variables in the ptx file |
| 1213 | // mapping demangled to mangled name |
| 1214 | inline void create_global_variable_map() { |
| 1215 | size_t pos = 0; |
| 1216 | while (pos < _ptx.size()) { |
| 1217 | pos = std::min(_ptx.find(".const .align", pos), |
| 1218 | _ptx.find(".global .align", pos)); |
| 1219 | if (pos == std::string::npos) break; |
| 1220 | size_t end = _ptx.find_first_of(";=", pos); |
| 1221 | if (_ptx[end] == '=') --end; |
| 1222 | std::string line = _ptx.substr(pos, end - pos); |
| 1223 | pos = end; |
| 1224 | size_t symbol_start = line.find_last_of(" ") + 1; |
| 1225 | size_t symbol_end = line.find_last_of("["); |
| 1226 | std::string entry = line.substr(symbol_start, symbol_end - symbol_start); |
| 1227 | std::string key = detail::demangle_ptx_variable_name(entry.c_str()); |
| 1228 | // Skip unsupported mangled names. E.g., a static variable defined inside |
| 1229 | // a function (such variables are not directly addressable from outside |
| 1230 | // the function, so skipping them is the correct behavior). |
| 1231 | if (key == "") continue; |
| 1232 | _global_map[key] = entry; |
| 1233 | } |
| 1234 | } |
| 1235 | |
| 1236 | inline void set_linker_log() { |
| 1237 | #ifdef JITIFY_PRINT_LINKER_LOG |
no test coverage detected