diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp index 53f0dc748a..fb35b030dc 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp @@ -656,6 +656,16 @@ typedef hipError_t (*t_hipModuleLoadData)(hipModule_t* module, const void* image typedef hipError_t (*t_hipModuleLoadDataEx)(hipModule_t* module, const void* image, unsigned int numOptions, hipJitOption* options, void** optionValues); +typedef hipError_t (*t_hipLinkAddData)(hipLinkState_t state, hipJitInputType type, void* data, + size_t size, const char* name, unsigned int numOptions, + hipJitOption* options, void** optionValues); +typedef hipError_t (*t_hipLinkAddFile)(hipLinkState_t state, hipJitInputType type, const char* path, + unsigned int numOptions, hipJitOption* options, void** optionValues); +typedef hipError_t (*t_hipLinkComplete)(hipLinkState_t state, void** hipBinOut, size_t* sizeOut); +typedef hipError_t (*t_hipLinkCreate)(unsigned int numOptions, hipJitOption* options, + void** optionValues, hipLinkState_t* stateOut); +typedef hipError_t (*t_hipLinkDestroy)(hipLinkState_t state); + typedef hipError_t (*t_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor)( int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk); typedef hipError_t (*t_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags)( @@ -1554,8 +1564,15 @@ struct HipDispatchTable { // HIP_RUNTIME_API_TABLE_STEP_VERSION == 9 t_hipEventRecordWithFlags hipEventRecordWithFlags_fn; - // DO NOT EDIT ABOVE! // HIP_RUNTIME_API_TABLE_STEP_VERSION == 10 + t_hipLinkAddData hipLinkAddData_fn; + t_hipLinkAddFile hipLinkAddFile_fn; + t_hipLinkComplete hipLinkComplete_fn; + t_hipLinkCreate hipLinkCreate_fn; + t_hipLinkDestroy hipLinkDestroy_fn; + + // DO NOT EDIT ABOVE! + // HIP_RUNTIME_API_TABLE_STEP_VERSION == 11 // ******************************************************************************************* // // diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index 9f083b7482..e1116ebdde 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -431,7 +431,12 @@ enum hip_api_id_t { HIP_API_ID_hipGraphBatchMemOpNodeSetParams = 411, HIP_API_ID_hipGraphExecBatchMemOpNodeSetParams = 412, HIP_API_ID_hipEventRecordWithFlags = 413, - HIP_API_ID_LAST = 413, + HIP_API_ID_hipLinkAddData = 414, + HIP_API_ID_hipLinkAddFile = 415, + HIP_API_ID_hipLinkComplete = 416, + HIP_API_ID_hipLinkCreate = 417, + HIP_API_ID_hipLinkDestroy = 418, + HIP_API_ID_LAST = 418, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -699,6 +704,11 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipLaunchCooperativeKernelMultiDevice: return "hipLaunchCooperativeKernelMultiDevice"; case HIP_API_ID_hipLaunchHostFunc: return "hipLaunchHostFunc"; case HIP_API_ID_hipLaunchKernel: return "hipLaunchKernel"; + case HIP_API_ID_hipLinkAddData: return "hipLinkAddData"; + case HIP_API_ID_hipLinkAddFile: return "hipLinkAddFile"; + case HIP_API_ID_hipLinkComplete: return "hipLinkComplete"; + case HIP_API_ID_hipLinkCreate: return "hipLinkCreate"; + case HIP_API_ID_hipLinkDestroy: return "hipLinkDestroy"; case HIP_API_ID_hipMalloc: return "hipMalloc"; case HIP_API_ID_hipMalloc3D: return "hipMalloc3D"; case HIP_API_ID_hipMalloc3DArray: return "hipMalloc3DArray"; @@ -1107,6 +1117,11 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipLaunchCooperativeKernelMultiDevice", name) == 0) return HIP_API_ID_hipLaunchCooperativeKernelMultiDevice; if (strcmp("hipLaunchHostFunc", name) == 0) return HIP_API_ID_hipLaunchHostFunc; if (strcmp("hipLaunchKernel", name) == 0) return HIP_API_ID_hipLaunchKernel; + if (strcmp("hipLinkAddData", name) == 0) return HIP_API_ID_hipLinkAddData; + if (strcmp("hipLinkAddFile", name) == 0) return HIP_API_ID_hipLinkAddFile; + if (strcmp("hipLinkComplete", name) == 0) return HIP_API_ID_hipLinkComplete; + if (strcmp("hipLinkCreate", name) == 0) return HIP_API_ID_hipLinkCreate; + if (strcmp("hipLinkDestroy", name) == 0) return HIP_API_ID_hipLinkDestroy; if (strcmp("hipMalloc", name) == 0) return HIP_API_ID_hipMalloc; if (strcmp("hipMalloc3D", name) == 0) return HIP_API_ID_hipMalloc3D; if (strcmp("hipMalloc3DArray", name) == 0) return HIP_API_ID_hipMalloc3DArray; @@ -2592,6 +2607,49 @@ typedef struct hip_api_data_s { size_t sharedMemBytes; hipStream_t stream; } hipLaunchKernel; + struct { + hipLinkState_t state; + hipJitInputType type; + void* data; + size_t size; + const char* name; + char name__val; + unsigned int numOptions; + hipJitOption* options; + hipJitOption options__val; + void** optionValues; + void* optionValues__val; + } hipLinkAddData; + struct { + hipLinkState_t state; + hipJitInputType type; + const char* path; + char path__val; + unsigned int numOptions; + hipJitOption* options; + hipJitOption options__val; + void** optionValues; + void* optionValues__val; + } hipLinkAddFile; + struct { + hipLinkState_t state; + void** hipBinOut; + void* hipBinOut__val; + size_t* sizeOut; + size_t sizeOut__val; + } hipLinkComplete; + struct { + unsigned int numOptions; + hipJitOption* options; + hipJitOption options__val; + void** optionValues; + void* optionValues__val; + hipLinkState_t* stateOut; + hipLinkState_t stateOut__val; + } hipLinkCreate; + struct { + hipLinkState_t state; + } hipLinkDestroy; struct { void** ptr; void* ptr__val; @@ -4991,6 +5049,43 @@ typedef struct hip_api_data_s { cb_data.args.hipLaunchKernel.sharedMemBytes = (size_t)sharedMemBytes; \ cb_data.args.hipLaunchKernel.stream = (hipStream_t)stream; \ }; +// hipLinkAddData[('hipLinkState_t', 'state'), ('hipJitInputType', 'type'), ('void*', 'data'), ('size_t', 'size'), ('const char*', 'name'), ('unsigned int', 'numOptions'), ('hipJitOption*', 'options'), ('void**', 'optionValues')] +#define INIT_hipLinkAddData_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLinkAddData.state = (hipLinkState_t)hip_link_state; \ + cb_data.args.hipLinkAddData.type = (hipJitInputType)input_type; \ + cb_data.args.hipLinkAddData.data = (void*)image; \ + cb_data.args.hipLinkAddData.size = (size_t)image_size; \ + cb_data.args.hipLinkAddData.name = (name) ? strdup(name) : NULL; \ + cb_data.args.hipLinkAddData.numOptions = (unsigned int)num_options; \ + cb_data.args.hipLinkAddData.options = (hipJitOption*)options_ptr; \ + cb_data.args.hipLinkAddData.optionValues = (void**)option_values; \ +}; +// hipLinkAddFile[('hipLinkState_t', 'state'), ('hipJitInputType', 'type'), ('const char*', 'path'), ('unsigned int', 'numOptions'), ('hipJitOption*', 'options'), ('void**', 'optionValues')] +#define INIT_hipLinkAddFile_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLinkAddFile.state = (hipLinkState_t)hip_link_state; \ + cb_data.args.hipLinkAddFile.type = (hipJitInputType)input_type; \ + cb_data.args.hipLinkAddFile.path = (file_path) ? strdup(file_path) : NULL; \ + cb_data.args.hipLinkAddFile.numOptions = (unsigned int)num_options; \ + cb_data.args.hipLinkAddFile.options = (hipJitOption*)options_ptr; \ + cb_data.args.hipLinkAddFile.optionValues = (void**)option_values; \ +}; +// hipLinkComplete[('hipLinkState_t', 'state'), ('void**', 'hipBinOut'), ('size_t*', 'sizeOut')] +#define INIT_hipLinkComplete_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLinkComplete.state = (hipLinkState_t)hip_link_state; \ + cb_data.args.hipLinkComplete.hipBinOut = (void**)bin_out; \ + cb_data.args.hipLinkComplete.sizeOut = (size_t*)size_out; \ +}; +// hipLinkCreate[('unsigned int', 'numOptions'), ('hipJitOption*', 'options'), ('void**', 'optionValues'), ('hipLinkState_t*', 'stateOut')] +#define INIT_hipLinkCreate_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLinkCreate.numOptions = (unsigned int)num_options; \ + cb_data.args.hipLinkCreate.options = (hipJitOption*)options_ptr; \ + cb_data.args.hipLinkCreate.optionValues = (void**)options_vals_pptr; \ + cb_data.args.hipLinkCreate.stateOut = (hipLinkState_t*)hip_link_state_ptr; \ +}; +// hipLinkDestroy[('hipLinkState_t', 'state')] +#define INIT_hipLinkDestroy_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLinkDestroy.state = (hipLinkState_t)hip_link_state; \ +}; // hipMalloc[('void**', 'ptr'), ('size_t', 'size')] #define INIT_hipMalloc_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMalloc.ptr = (void**)ptr; \ @@ -7082,6 +7177,32 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipLaunchKernel: if (data->args.hipLaunchKernel.args) data->args.hipLaunchKernel.args__val = *(data->args.hipLaunchKernel.args); break; +// hipLinkAddData[('hipLinkState_t', 'state'), ('hipJitInputType', 'type'), ('void*', 'data'), ('size_t', 'size'), ('const char*', 'name'), ('unsigned int', 'numOptions'), ('hipJitOption*', 'options'), ('void**', 'optionValues')] + case HIP_API_ID_hipLinkAddData: + if (data->args.hipLinkAddData.name) data->args.hipLinkAddData.name__val = *(data->args.hipLinkAddData.name); + if (data->args.hipLinkAddData.options) data->args.hipLinkAddData.options__val = *(data->args.hipLinkAddData.options); + if (data->args.hipLinkAddData.optionValues) data->args.hipLinkAddData.optionValues__val = *(data->args.hipLinkAddData.optionValues); + break; +// hipLinkAddFile[('hipLinkState_t', 'state'), ('hipJitInputType', 'type'), ('const char*', 'path'), ('unsigned int', 'numOptions'), ('hipJitOption*', 'options'), ('void**', 'optionValues')] + case HIP_API_ID_hipLinkAddFile: + if (data->args.hipLinkAddFile.path) data->args.hipLinkAddFile.path__val = *(data->args.hipLinkAddFile.path); + if (data->args.hipLinkAddFile.options) data->args.hipLinkAddFile.options__val = *(data->args.hipLinkAddFile.options); + if (data->args.hipLinkAddFile.optionValues) data->args.hipLinkAddFile.optionValues__val = *(data->args.hipLinkAddFile.optionValues); + break; +// hipLinkComplete[('hipLinkState_t', 'state'), ('void**', 'hipBinOut'), ('size_t*', 'sizeOut')] + case HIP_API_ID_hipLinkComplete: + if (data->args.hipLinkComplete.hipBinOut) data->args.hipLinkComplete.hipBinOut__val = *(data->args.hipLinkComplete.hipBinOut); + if (data->args.hipLinkComplete.sizeOut) data->args.hipLinkComplete.sizeOut__val = *(data->args.hipLinkComplete.sizeOut); + break; +// hipLinkCreate[('unsigned int', 'numOptions'), ('hipJitOption*', 'options'), ('void**', 'optionValues'), ('hipLinkState_t*', 'stateOut')] + case HIP_API_ID_hipLinkCreate: + if (data->args.hipLinkCreate.options) data->args.hipLinkCreate.options__val = *(data->args.hipLinkCreate.options); + if (data->args.hipLinkCreate.optionValues) data->args.hipLinkCreate.optionValues__val = *(data->args.hipLinkCreate.optionValues); + if (data->args.hipLinkCreate.stateOut) data->args.hipLinkCreate.stateOut__val = *(data->args.hipLinkCreate.stateOut); + break; +// hipLinkDestroy[('hipLinkState_t', 'state')] + case HIP_API_ID_hipLinkDestroy: + break; // hipMalloc[('void**', 'ptr'), ('size_t', 'size')] case HIP_API_ID_hipMalloc: if (data->args.hipMalloc.ptr) data->args.hipMalloc.ptr__val = *(data->args.hipMalloc.ptr); @@ -9525,6 +9646,59 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLaunchKernel.stream); oss << ")"; break; + case HIP_API_ID_hipLinkAddData: + oss << "hipLinkAddData("; + oss << "state="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddData.state); + oss << ", type="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddData.type); + oss << ", data="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddData.data); + oss << ", size="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddData.size); + if (data->args.hipLinkAddData.name == NULL) oss << ", name=NULL"; + else { oss << ", name="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddData.name__val); } + oss << ", numOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddData.numOptions); + if (data->args.hipLinkAddData.options == NULL) oss << ", options=NULL"; + else { oss << ", options="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddData.options__val); } + if (data->args.hipLinkAddData.optionValues == NULL) oss << ", optionValues=NULL"; + else { oss << ", optionValues="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddData.optionValues__val); } + oss << ")"; + break; + case HIP_API_ID_hipLinkAddFile: + oss << "hipLinkAddFile("; + oss << "state="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddFile.state); + oss << ", type="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddFile.type); + if (data->args.hipLinkAddFile.path == NULL) oss << ", path=NULL"; + else { oss << ", path="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddFile.path__val); } + oss << ", numOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddFile.numOptions); + if (data->args.hipLinkAddFile.options == NULL) oss << ", options=NULL"; + else { oss << ", options="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddFile.options__val); } + if (data->args.hipLinkAddFile.optionValues == NULL) oss << ", optionValues=NULL"; + else { oss << ", optionValues="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddFile.optionValues__val); } + oss << ")"; + break; + case HIP_API_ID_hipLinkComplete: + oss << "hipLinkComplete("; + oss << "state="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkComplete.state); + if (data->args.hipLinkComplete.hipBinOut == NULL) oss << ", hipBinOut=NULL"; + else { oss << ", hipBinOut="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkComplete.hipBinOut__val); } + if (data->args.hipLinkComplete.sizeOut == NULL) oss << ", sizeOut=NULL"; + else { oss << ", sizeOut="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkComplete.sizeOut__val); } + oss << ")"; + break; + case HIP_API_ID_hipLinkCreate: + oss << "hipLinkCreate("; + oss << "numOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkCreate.numOptions); + if (data->args.hipLinkCreate.options == NULL) oss << ", options=NULL"; + else { oss << ", options="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkCreate.options__val); } + if (data->args.hipLinkCreate.optionValues == NULL) oss << ", optionValues=NULL"; + else { oss << ", optionValues="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkCreate.optionValues__val); } + if (data->args.hipLinkCreate.stateOut == NULL) oss << ", stateOut=NULL"; + else { oss << ", stateOut="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkCreate.stateOut__val); } + oss << ")"; + break; + case HIP_API_ID_hipLinkDestroy: + oss << "hipLinkDestroy("; + oss << "state="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkDestroy.state); + oss << ")"; + break; case HIP_API_ID_hipMalloc: oss << "hipMalloc("; if (data->args.hipMalloc.ptr == NULL) oss << "ptr=NULL"; diff --git a/projects/clr/hipamd/src/CMakeLists.txt b/projects/clr/hipamd/src/CMakeLists.txt index a276588825..5f6736d8dd 100644 --- a/projects/clr/hipamd/src/CMakeLists.txt +++ b/projects/clr/hipamd/src/CMakeLists.txt @@ -132,11 +132,12 @@ target_sources(amdhip64 PRIVATE hip_vm.cpp hip_api_trace.cpp hip_table_interface.cpp - hip_table_interface_c.cpp) + hip_table_interface_c.cpp + hip_comgr_helper.cpp) if(WIN32) target_sources(amdhip64 PRIVATE - hip_runtime.cpp) + hip_runtime.cpp hiprtc/hiprtcInternal.cpp) endif() if(BUILD_SHARED_LIBS) @@ -208,16 +209,14 @@ set(HIPRTC_OBJECTS) # Add hiprtc add_subdirectory(hiprtc) -if(NOT WIN32) - target_compile_definitions(amdhip64 PRIVATE __HIP_ENABLE_RTC) - if(BUILD_SHARED_LIBS) - target_link_libraries(amdhip64 PRIVATE ${HIPRTC_OBJECTS}) - add_dependencies(amdhip64 hiprtc-builtins) - INSTALL(TARGETS hiprtc-builtins - RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} - LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} - ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}) - endif() +target_compile_definitions(amdhip64 PRIVATE __HIP_ENABLE_RTC) +if(BUILD_SHARED_LIBS) + target_link_libraries(amdhip64 PRIVATE ${HIPRTC_OBJECTS}) + add_dependencies(amdhip64 hiprtc-builtins) + INSTALL(TARGETS hiprtc-builtins + RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} + LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} + ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}) endif() ############################# diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 28d8d7be9f..f5d8dbc465 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -485,3 +485,8 @@ hipGraphBatchMemOpNodeGetParams hipGraphBatchMemOpNodeSetParams hipGraphExecBatchMemOpNodeSetParams hipEventRecordWithFlags +hipLinkAddData +hipLinkAddFile +hipLinkComplete +hipLinkCreate +hipLinkDestroy \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index 73b352a891..9a5bb9151d 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -537,6 +537,15 @@ hipError_t hipModuleLoad(hipModule_t* module, const char* fname); hipError_t hipModuleLoadData(hipModule_t* module, const void* image); hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned int numOptions, hipJitOption* options, void** optionValues); +hipError_t hipLinkAddData(hipLinkState_t state, hipJitInputType type, void* data, size_t size, + const char* name, unsigned int numOptions, hipJitOption* options, + void** optionValues); +hipError_t hipLinkAddFile(hipLinkState_t state, hipJitInputType type, const char* path, unsigned int numOptions, + hipJitOption* options, void** optionValues); +hipError_t hipLinkComplete(hipLinkState_t state, void** hipBinOut, size_t* sizeOut); +hipError_t hipLinkCreate(unsigned int numOptions, hipJitOption* options, + void** optionValues, hipLinkState_t* stateOut); +hipError_t hipLinkDestroy(hipLinkState_t state); hipError_t hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk); @@ -1166,6 +1175,11 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipModuleLoad_fn = hip::hipModuleLoad; ptrDispatchTable->hipModuleLoadData_fn = hip::hipModuleLoadData; ptrDispatchTable->hipModuleLoadDataEx_fn = hip::hipModuleLoadDataEx; + ptrDispatchTable->hipLinkAddData_fn = hip::hipLinkAddData; + ptrDispatchTable->hipLinkAddFile_fn = hip::hipLinkAddFile; + ptrDispatchTable->hipLinkComplete_fn = hip::hipLinkComplete; + ptrDispatchTable->hipLinkCreate_fn = hip::hipLinkCreate; + ptrDispatchTable->hipLinkDestroy_fn = hip::hipLinkDestroy; ptrDispatchTable->hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_fn = hip::hipModuleOccupancyMaxActiveBlocksPerMultiprocessor; ptrDispatchTable->hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_fn = @@ -1952,13 +1966,19 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipGraphBatchMemOpNodeSetParams_fn, 466); HIP_ENFORCE_ABI(HipDispatchTable, hipGraphExecBatchMemOpNodeSetParams_fn, 467); // HIP_RUNTIME_API_TABLE_STEP_VERSION == 9 HIP_ENFORCE_ABI(HipDispatchTable, hipEventRecordWithFlags_fn, 468) +// HIP_RUNTIME_API_TABLE_STEP_VERSION == 10 +HIP_ENFORCE_ABI(HipDispatchTable, hipLinkAddData_fn , 469) +HIP_ENFORCE_ABI(HipDispatchTable, hipLinkAddFile_fn , 470) +HIP_ENFORCE_ABI(HipDispatchTable, hipLinkComplete_fn , 471) +HIP_ENFORCE_ABI(HipDispatchTable, hipLinkCreate_fn , 472) +HIP_ENFORCE_ABI(HipDispatchTable, hipLinkDestroy_fn , 473) // if HIP_ENFORCE_ABI entries are added for each new function pointer in the table, the number below // will be +1 of the number in the last HIP_ENFORCE_ABI line. E.g.: // // HIP_ENFORCE_ABI(, , 8) // // HIP_ENFORCE_ABI_VERSIONING(
, 9) <- 8 + 1 = 9 -HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 469) +HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 474) static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 9, "If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function " diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp b/projects/clr/hipamd/src/hip_comgr_helper.cpp similarity index 65% rename from projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp rename to projects/clr/hipamd/src/hip_comgr_helper.cpp index 5d0dd74bf9..5e336df75e 100644 --- a/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.cpp +++ b/projects/clr/hipamd/src/hip_comgr_helper.cpp @@ -20,14 +20,16 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "hiprtcComgrHelper.hpp" +#include "hip_comgr_helper.hpp" #if defined(_WIN32) #include #endif -#include "../amd_hsa_elf.hpp" +#include "../src/amd_hsa_elf.hpp" -namespace hiprtc { +namespace hip { + +std::unordered_set LinkProgram::linker_set_; namespace helpers { @@ -40,6 +42,7 @@ constexpr char const* OFFLOAD_KIND_HIP = "hip"; constexpr char const* OFFLOAD_KIND_HIPV4 = "hipv4"; constexpr char const* OFFLOAD_KIND_HCC = "hcc"; constexpr char const* AMDGCN_TARGET_TRIPLE = "amdgcn-amd-amdhsa-"; +constexpr char const* SPIRV_BUNDLE_ENTRY_ID = "hip-spirv64-amd-amdhsa-amdgcnspirv"; static constexpr size_t bundle_magic_string_size = strLiteralLength(CLANG_OFFLOAD_BUNDLER_MAGIC_STR); @@ -819,13 +822,86 @@ bool compileToBitCode(const amd_comgr_data_set_t compileInputs, const std::strin return true; } +bool CheckIfBundled(std::vector& llvm_bitcode) { + std::string magic(llvm_bitcode.begin(), + llvm_bitcode.begin() + bundle_magic_string_size); + + if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) == 0) { + return true; + } + // File is not bundled + return false; + +} +// Unbundle Bitcode using COMGR action +// Supports only 1 Bundle Entry ID for now +bool UnbundleUsingComgr(std::vector& source, const std::string& isa, + std::vector& linkOptions, std::string& buildLog, + std::vector& unbundled_bitcode, const char *bundleEntryIDs[], + size_t bundleEntryIDsCount) { + amd_comgr_data_set_t linkinput; + if (amd::Comgr::create_data_set(&linkinput) != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + std::string name = "UnbundleCode.bc"; + if (!helpers::addCodeObjData(linkinput, source, name, AMD_COMGR_DATA_KIND_BC_BUNDLE)) { + return false; + } + + amd_comgr_action_info_t action; + if (createAction(action, linkOptions, isa, AMD_COMGR_LANGUAGE_NONE) != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + if (bundleEntryIDsCount > 1) { + LogError("Error in hip Linker : bundleEntryID count > 1"); + return false; + } + + if(amd::Comgr::action_info_set_bundle_entry_ids(action, bundleEntryIDs, bundleEntryIDsCount) != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return false; + } + + amd_comgr_data_set_t output; + if (amd::Comgr::create_data_set(&output) != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return false; + } + + if (auto res = + amd::Comgr::do_action(AMD_COMGR_ACTION_UNBUNDLE, action, linkinput, output); + res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return false; + } + + if (!extractBuildLog(output, buildLog)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return false; + } + + if (!extractByteCodeBinary(output, AMD_COMGR_DATA_KIND_BC, unbundled_bitcode)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return false; + } + + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + amd::Comgr::destroy_data_set(linkinput); + return true; +} + bool linkLLVMBitcode(const amd_comgr_data_set_t linkInputs, const std::string& isa, std::vector& linkOptions, std::string& buildLog, std::vector& LinkedLLVMBitcode) { - amd_comgr_language_t lang = AMD_COMGR_LANGUAGE_HIP; + const amd_comgr_language_t lang = AMD_COMGR_LANGUAGE_HIP; amd_comgr_action_info_t action; - if (auto res = createAction(action, linkOptions, isa, AMD_COMGR_LANGUAGE_HIP); + if (auto res = createAction(action, linkOptions, isa, lang); res != AMD_COMGR_STATUS_SUCCESS) { return false; } @@ -860,15 +936,66 @@ bool linkLLVMBitcode(const amd_comgr_data_set_t linkInputs, const std::string& i return true; } +bool convertSPIRVToLLVMBC(const amd_comgr_data_set_t linkInputs, const std::string& isa, + std::vector& linkOptions, std::string& buildLog, + std::vector& LinkedLLVMBitcode) { + amd_comgr_action_info_t action; + + if (auto res = createAction(action, linkOptions, isa, AMD_COMGR_LANGUAGE_NONE); + res != AMD_COMGR_STATUS_SUCCESS) { + return false; + } + + amd_comgr_data_set_t output; + if (auto res = amd::Comgr::create_data_set(&output); res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + return false; + } + + if (auto res = + amd::Comgr::do_action(AMD_COMGR_ACTION_TRANSLATE_SPIRV_TO_BC, action, linkInputs, output); + res != AMD_COMGR_STATUS_SUCCESS) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return false; + } + + if (!extractBuildLog(output, buildLog)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return false; + } + + if (!extractByteCodeBinary(output, AMD_COMGR_DATA_KIND_BC, LinkedLLVMBitcode)) { + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return false; + } + + amd::Comgr::destroy_action_info(action); + amd::Comgr::destroy_data_set(output); + return true; +} + bool createExecutable(const amd_comgr_data_set_t linkInputs, const std::string& isa, std::vector& exeOptions, std::string& buildLog, - std::vector& executable) { + std::vector& executable, bool spirv_bc /* default false */) { amd_comgr_action_info_t action; if (auto res = createAction(action, exeOptions, isa); res != AMD_COMGR_STATUS_SUCCESS) { return false; } + // If SPIRV bitcode was processed, make sure we link device libs to it + if (spirv_bc) { + if (auto res = amd::Comgr::action_info_set_device_lib_linking(action, true); + res != AMD_COMGR_STATUS_SUCCESS) { + LogError("Can not link device libs to action"); + amd::Comgr::destroy_action_info(action); + return false; + } + } + amd_comgr_data_set_t relocatableData; if (auto res = amd::Comgr::create_data_set(&relocatableData); res != AMD_COMGR_STATUS_SUCCESS) { amd::Comgr::destroy_action_info(action); @@ -985,7 +1112,7 @@ bool dumpIsaFromBC(const amd_comgr_data_set_t isaInputs, const std::string& isa, if (name.size() == 0) { // Generate a unique name if the program name is not specified by the user - name = std::string("hiprtcXXXXXX"); + name = std::string("hipXXXXXX"); GenerateUniqueFileName(name); } std::string isaName = isa; @@ -1128,4 +1255,422 @@ bool fillMangledNames(std::vector& dataVec, std::map getLinkOptions(const LinkArguments& args) { + std::vector res; + + { // process optimization level + std::string opt("-O"); + opt += std::to_string(args.optimization_level_); + res.push_back(opt); + } + + const auto irArgCount = args.linker_ir2isa_args_count_; + if (irArgCount > 0) { + res.reserve(irArgCount); + const auto irArg = args.linker_ir2isa_args_; + for (size_t i = 0; i < irArgCount; i++) { + res.emplace_back(std::string(irArg[i])); + } + } + return res; +} + + +// RTC Program Member Functions +RTCProgram::RTCProgram(std::string name) : name_(name) { + constexpr bool kComgrVersioned = true; + std::call_once(amd::Comgr::initialized, amd::Comgr::LoadLib, kComgrVersioned); + if (amd::Comgr::create_data_set(&exec_input_) != AMD_COMGR_STATUS_SUCCESS) { + guarantee(false, "Failed to allocate internal hiprtc structure"); + } +} + +bool RTCProgram::findIsa() { + +#ifdef BUILD_SHARED_LIBS + const char* libName; +#ifdef _WIN32 + std::string dll_name = std::string("amdhip64_" + std::to_string(HIP_VERSION_MAJOR) + ".dll"); + libName = dll_name.c_str(); +#else + libName = "libamdhip64.so"; +#endif + + void* handle = amd::Os::loadLibrary(libName); + + if (!handle) { + LogInfo("hip runtime failed to load using dlopen"); + build_log_ += + "hip runtime failed to load.\n" + "Error: Please provide architecture for which code is to be " + "generated.\n"; + return false; + } + + void* sym_hipGetDevice = amd::Os::getSymbol(handle, "hipGetDevice"); + void* sym_hipGetDeviceProperties = + amd::Os::getSymbol(handle, "hipGetDevicePropertiesR0600"); // Try to find the new symbol + if (sym_hipGetDeviceProperties == nullptr) { + sym_hipGetDeviceProperties = + amd::Os::getSymbol(handle, "hipGetDeviceProperties"); // Fall back to old one + } + + if (sym_hipGetDevice == nullptr || sym_hipGetDeviceProperties == nullptr) { + LogInfo("ISA cannot be found to dlsym failure"); + build_log_ += + "ISA cannot be found from hip runtime.\n" + "Error: Please provide architecture for which code is to be " + "generated.\n"; + return false; + } + + hipError_t (*dyn_hipGetDevice)(int*) = reinterpret_cast(sym_hipGetDevice); + + hipError_t (*dyn_hipGetDeviceProperties)(hipDeviceProp_t*, int) = + reinterpret_cast(sym_hipGetDeviceProperties); + + int device; + hipError_t status = dyn_hipGetDevice(&device); + if (status != hipSuccess) { + return false; + } + hipDeviceProp_t props; + status = dyn_hipGetDeviceProperties(&props, device); + if (status != hipSuccess) { + return false; + } + isa_ = "amdgcn-amd-amdhsa--"; + isa_.append(props.gcnArchName); + + amd::Os::unloadLibrary(handle); + return true; + +#else + int device; + hipError_t status = hipGetDevice(&device); + if (status != hipSuccess) { + return false; + } + hipDeviceProp_t props; + status = hipGetDeviceProperties(&props, device); + if (status != hipSuccess) { + return false; + } + isa_ = "amdgcn-amd-amdhsa--"; + isa_.append(props.gcnArchName); + + return true; +#endif +} + +// RTC Program Member Functions +void RTCProgram::AppendOptions(const std::string app_env_var, std::vector* options) { + if (options == nullptr) { + LogError("Append options passed is nullptr."); + return; + } + + std::stringstream ss(app_env_var); + std::istream_iterator begin{ss}, end; + options->insert(options->end(), begin, end); +} + +// HIPRTC Program lock +amd::Monitor RTCProgram::lock_(true); + +LinkProgram::LinkProgram(std::string name) : RTCProgram(name) { + if (amd::Comgr::create_data_set(&link_input_) != AMD_COMGR_STATUS_SUCCESS) { + guarantee(false, "Failed to allocate internal comgr structure"); + } + amd::ScopedLock lock(lock_); + linker_set_.insert(this); +} + +bool LinkProgram::isLinkerValid(LinkProgram* link_program) { + amd::ScopedLock lock(lock_); + if (linker_set_.find(link_program) == linker_set_.end()) { + return false; + } + return true; +} + +bool LinkProgram::AddLinkerOptions(unsigned int num_options, hipJitOption* options_ptr, + void** options_vals_ptr) { + for (size_t opt_idx = 0; opt_idx < num_options; ++opt_idx) { + if ( options_vals_ptr[opt_idx] == nullptr) { + LogError("Options value can not be nullptr"); + return false; + } + switch (options_ptr[opt_idx]) { + case hipJitOptionMaxRegisters: + link_args_.max_registers_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionThreadsPerBlock: + link_args_.threads_per_block_ = + *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionWallTime: + link_args_.wall_time_ = *(reinterpret_cast(options_vals_ptr[opt_idx])); + break; + case hipJitOptionInfoLogBuffer: { + link_args_.info_log_ = (reinterpret_cast(options_vals_ptr[opt_idx])); + break; + } + case hipJitOptionInfoLogBufferSizeBytes: + link_args_.info_log_size_ = (reinterpret_cast(options_vals_ptr[opt_idx])); + break; + case hipJitOptionErrorLogBuffer: { + link_args_.error_log_ = reinterpret_cast(options_vals_ptr[opt_idx]); + break; + } + case hipJitOptionErrorLogBufferSizeBytes: + link_args_.error_log_size_ = (reinterpret_cast(options_vals_ptr[opt_idx])); + break; + case hipJitOptionOptimizationLevel: + link_args_.optimization_level_ = + *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionTargetFromContext: + link_args_.target_from_hip_context_ = + *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionTarget: + link_args_.jit_target_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionFallbackStrategy: + link_args_.fallback_strategy_ = + *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionGenerateDebugInfo: + link_args_.generate_debug_info_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionLogVerbose: + link_args_.log_verbose_ = reinterpret_cast(options_vals_ptr[opt_idx]); + break; + case hipJitOptionGenerateLineInfo: + link_args_.generate_line_info_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionCacheMode: + link_args_.cache_mode_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionSm3xOpt: + link_args_.sm3x_opt_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionFastCompile: + link_args_.fast_compile_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionGlobalSymbolNames: { + link_args_.global_symbol_names_ = reinterpret_cast(options_vals_ptr[opt_idx]); + break; + } + case hipJitOptionGlobalSymbolAddresses: { + link_args_.global_symbol_addresses_ = reinterpret_cast(options_vals_ptr[opt_idx]); + break; + } + case hipJitOptionGlobalSymbolCount: + link_args_.global_symbol_count_ = + *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionLto: + link_args_.lto_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionFtz: + link_args_.ftz_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionPrecDiv: + link_args_.prec_div_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionPrecSqrt: + link_args_.prec_sqrt_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionFma: + link_args_.fma_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + break; + case hipJitOptionPositionIndependentCode: + link_args_.pic_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + case hipJitOptionMinCTAPerSM: + link_args_.min_cta_per_sm_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + case hipJitOptionMaxThreadsPerBlock: + link_args_.max_threads_per_block_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + case hipJitOptionOverrideDirectiveValues: + link_args_.override_directive_values_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); + case hipJitOptionIRtoISAOptExt: { + link_args_.linker_ir2isa_args_ = reinterpret_cast(options_vals_ptr[opt_idx]); + break; + } + case hipJitOptionIRtoISAOptCountExt: + link_args_.linker_ir2isa_args_count_ = reinterpret_cast(options_vals_ptr[opt_idx]); + break; + default: + break; + } + } + + return true; +} + + + +amd_comgr_data_kind_t LinkProgram::GetCOMGRDataKind(hipJitInputType input_type) { + amd_comgr_data_kind_t data_kind = AMD_COMGR_DATA_KIND_UNDEF; + + // Map the hiprtc input type to comgr data kind + switch (input_type) { + case hipJitInputLLVMBitcode: + data_kind = AMD_COMGR_DATA_KIND_BC; + break; + case hipJitInputLLVMBundledBitcode: + data_kind = + HIPRTC_USE_RUNTIME_UNBUNDLER ? AMD_COMGR_DATA_KIND_BC : AMD_COMGR_DATA_KIND_BC_BUNDLE; + break; + case hipJitInputLLVMArchivesOfBundledBitcode: + data_kind = AMD_COMGR_DATA_KIND_AR_BUNDLE; + break; + case hipJitInputSpirv: + data_kind = AMD_COMGR_DATA_KIND_SPIRV; + break; + default: + LogError("hip link : Cannot find the corresponding comgr data kind"); + break; + } + + return data_kind; +} + + +bool LinkProgram::AddLinkerDataImpl(std::vector& link_data, hipJitInputType input_type, + std::string& link_file_name) { + std::vector llvm_code_object; + is_bundled_ = helpers::CheckIfBundled(link_data); + + if (HIPRTC_USE_RUNTIME_UNBUNDLER && input_type == hipJitInputLLVMBundledBitcode) { + if (!findIsa()) { + return false; + } + + size_t co_offset = 0; + size_t co_size = 0; + if (!helpers::UnbundleBitCode(link_data, isa_, co_offset, co_size)) { + LogError("Error in hip Linker: unable to unbundle the llvm bitcode"); + return false; + } + + llvm_code_object.assign(link_data.begin() + co_offset, link_data.begin() + co_offset + co_size); + } else if (is_bundled_ && input_type == hipJitInputSpirv) { + const char* bundleEntryIDs[] = { helpers::SPIRV_BUNDLE_ENTRY_ID }; + size_t bundleEntryIDsCount = sizeof(bundleEntryIDs) / sizeof(bundleEntryIDs[0]); + if(!helpers::UnbundleUsingComgr(link_data, isa_, link_options_, build_log_, llvm_code_object, + bundleEntryIDs, bundleEntryIDsCount)) { + LogError("Error in hip Linker: Unable to unbundle SPIRV Bitcode"); + return false; + } + } else { + llvm_code_object.assign(link_data.begin(), link_data.end()); + } + + if ((data_kind_ = GetCOMGRDataKind(input_type)) == AMD_COMGR_DATA_KIND_UNDEF) { + LogError("Cannot find the correct COMGR data kind"); + return false; + } + + if (!helpers::addCodeObjData(link_input_, llvm_code_object, link_file_name, data_kind_)) { + LogError("Error in hip Linker: unable to add linked code object"); + return false; + } + + return true; +} + + +bool LinkProgram::AddLinkerFile(std::string file_path, hipJitInputType input_type) { + std::ifstream file_stream{file_path, std::ios_base::in | std::ios_base::binary}; + if (!file_stream.good()) { + return false; + } + + file_stream.seekg(0, std::ios::end); + std::streampos file_size = file_stream.tellg(); + file_stream.seekg(0, std::ios::beg); + + // Read the file contents + std::vector link_file_info(file_size); + file_stream.read(link_file_info.data(), file_size); + + file_stream.close(); + + std::string link_file_name("LinkerProgram"); + + return AddLinkerDataImpl(link_file_info, input_type, link_file_name); +} + +bool LinkProgram::AddLinkerData(void* image_ptr, size_t image_size, std::string link_file_name, + hipJitInputType input_type) { + char* image_char_buf = reinterpret_cast(image_ptr); + std::vector llvm_code_object(image_char_buf, image_char_buf + image_size); + + return AddLinkerDataImpl(llvm_code_object, input_type, link_file_name); +} + +bool LinkProgram::LinkComplete(void** bin_out, size_t* size_out) { + if (!findIsa()) { + return false; + } + + // If the data kind is SPIRV, convert it beforehand and pass it on to subsequent machinery + // TODO I think this can be simplified a bit, we are basically reading and writing into comgr data + // structures, do we need to do that? This might cause some errors, so adding this to come back to + // it. + amd_comgr_data_set_t link_input = link_input_; + if (data_kind_ == AMD_COMGR_DATA_KIND_SPIRV) { + // Convert SPIRV Unbundled code object to LLVM Bitcode + std::vector llvmbc_from_spirv; + if (!helpers::convertSPIRVToLLVMBC(link_input_, isa_, link_options_, build_log_, llvmbc_from_spirv)) { + LogError("Error in hip Linker: unable to convert SPIRV to BC"); + return false; + } + + std::string linkedFileName = "LLVMBitcodeFromSPIRV.bc"; + if (!helpers::addCodeObjData(link_input, llvmbc_from_spirv, linkedFileName, AMD_COMGR_DATA_KIND_BC)) { + LogError("Error in hip Linker: unable to add linked LLVM bitcode"); + return false; + } + } + + std::vector llvm_bitcode; + if (!helpers::linkLLVMBitcode(link_input, isa_, link_options_, build_log_, llvm_bitcode)) { + LogError("Error in hip linker: unable to add device libs to linked bitcode"); + return false; + } + + std::string linkedFileName = "LLVMBitcode.bc"; + if (!helpers::addCodeObjData(exec_input_, llvm_bitcode, linkedFileName, AMD_COMGR_DATA_KIND_BC)) { + LogError("Error in hip linker: unable to add linked bitcode"); + return false; + } + + std::vector exe_options = getLinkOptions(link_args_); + LogPrintfInfo("Exe options forwarded to compiler: %s", + [&]() { + std::string ret; + for (const auto& i : exe_options) { + ret += i; + ret += " "; + } + return ret; + }() + .c_str()); + if (!helpers::createExecutable(exec_input_, isa_, exe_options, build_log_, executable_, + data_kind_ == AMD_COMGR_DATA_KIND_SPIRV)) { + LogPrintfInfo("Error in hip linker: unable to create exectuable: %s", build_log_.c_str()); + return false; + } + + *size_out = executable_.size(); + *bin_out = executable_.data(); + + return true; +} + + +} // namespace hip diff --git a/projects/clr/hipamd/src/hip_comgr_helper.hpp b/projects/clr/hipamd/src/hip_comgr_helper.hpp new file mode 100644 index 0000000000..98fc72bb7f --- /dev/null +++ b/projects/clr/hipamd/src/hip_comgr_helper.hpp @@ -0,0 +1,210 @@ +/* +Copyright (c) 2022 - Present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include +#include + + +#include "vdi_common.hpp" +#include "rocclr/utils/debug.hpp" +#include "rocclr/utils/flags.hpp" +#include "device/comgrctx.hpp" +#include "hip/hip_runtime_api.h" + +namespace hip { +namespace helpers { +bool UnbundleBitCode(const std::vector& bundled_bit_code, const std::string& isa, + size_t& co_offset, size_t& co_size); +bool addCodeObjData(amd_comgr_data_set_t& input, const std::vector& source, + const std::string& name, const amd_comgr_data_kind_t type); +bool extractBuildLog(amd_comgr_data_set_t dataSet, std::string& buildLog); +bool extractByteCodeBinary(const amd_comgr_data_set_t inDataSet, + const amd_comgr_data_kind_t dataKind, std::vector& bin); +bool createAction(amd_comgr_action_info_t& action, std::vector& options, + const std::string& isa, + const amd_comgr_language_t lang = AMD_COMGR_LANGUAGE_NONE); +bool compileToExecutable(const amd_comgr_data_set_t compileInputs, const std::string& isa, + std::vector& compileOptions, + std::vector& linkOptions, std::string& buildLog, + std::vector& exe); +bool compileToBitCode(const amd_comgr_data_set_t compileInputs, const std::string& isa, + std::vector& compileOptions, std::string& buildLog, + std::vector& LLVMBitcode); +bool linkLLVMBitcode(const amd_comgr_data_set_t linkInputs, const std::string& isa, + std::vector& linkOptions, std::string& buildLog, + std::vector& LinkedLLVMBitcode); +bool createExecutable(const amd_comgr_data_set_t linkInputs, const std::string& isa, + std::vector& exeOptions, std::string& buildLog, + std::vector& executable, bool spirv_bc = false); +bool convertSPIRVToLLVMBC(const amd_comgr_data_set_t linkInputs, const std::string& isa, + std::vector& linkOptions, std::string& buildLog, + std::vector& linkedSPIRVBitcode); +bool dumpIsaFromBC(const amd_comgr_data_set_t isaInputs, const std::string& isa, + std::vector& exeOptions, std::string name, std::string& buildLog); +bool demangleName(const std::string& mangledName, std::string& demangledName); +std::string handleMangledName(std::string loweredName); +bool fillMangledNames(std::vector& executable, + std::map& mangledNames, bool isBitcode); +void GenerateUniqueFileName(std::string& name); + +bool CheckIfBundled(std::vector& llvm_bitcode); + +bool UnbundleUsingComgr(std::vector& source, const std::string& isa, + std::vector& linkOptions, std::string& buildLog, + std::vector& unbundled_spirv_bitcode, const char* bundleEntryIDs, + size_t bundleEntryIDsCount); +} // namespace helpers + +struct LinkArguments { + unsigned int max_registers_; + unsigned int threads_per_block_; + float wall_time_; + size_t info_log_size_; + char* info_log_; + size_t error_log_size_; + char* error_log_; + unsigned int optimization_level_; + unsigned int target_from_hip_context_; + unsigned int jit_target_; + unsigned int fallback_strategy_; + int generate_debug_info_; + long log_verbose_; + int generate_line_info_; + unsigned int cache_mode_; + bool sm3x_opt_; + bool fast_compile_; + const char** global_symbol_names_; + void** global_symbol_addresses_; + unsigned int global_symbol_count_; + int lto_; + int ftz_; + int prec_div_; + int prec_sqrt_; + int fma_; + int pic_; + int min_cta_per_sm_; + int max_threads_per_block_; + int override_directive_values_; + const char** linker_ir2isa_args_; + size_t linker_ir2isa_args_count_; + + LinkArguments() + : max_registers_{0}, + threads_per_block_{0}, + wall_time_{0.0f}, + info_log_size_{0}, + info_log_{nullptr}, + error_log_size_{0}, + error_log_{nullptr}, + optimization_level_{3}, + target_from_hip_context_{0}, + jit_target_{0}, + fallback_strategy_{0}, + generate_debug_info_{0}, + log_verbose_{0}, + generate_line_info_{0}, + cache_mode_{0}, + sm3x_opt_{false}, + fast_compile_{false}, + global_symbol_names_{nullptr}, + global_symbol_addresses_{nullptr}, + global_symbol_count_{0}, + lto_{0}, + ftz_{0}, + prec_div_{1}, + prec_sqrt_{1}, + fma_{1}, + pic_{0}, + min_cta_per_sm_{0}, + max_threads_per_block_{0}, + override_directive_values_{0}, + linker_ir2isa_args_{nullptr}, + linker_ir2isa_args_count_{0} {} +}; + +class RTCProgram { + protected: + // Lock and control variables + static amd::Monitor lock_; + static std::once_flag initialized_; + + RTCProgram(std::string name); + ~RTCProgram() { amd::Comgr::destroy_data_set(exec_input_); } + + // Member Functions + bool findIsa(); + static void AppendOptions(std::string app_env_var, std::vector* options); + + // Data Members + std::string name_; + std::string isa_; + std::string build_log_; + std::vector executable_; + + amd_comgr_data_set_t exec_input_; +}; + +class LinkProgram : public RTCProgram { + // Private Member Functions (forbid these function calls) + LinkProgram() = delete; + LinkProgram(LinkProgram&) = delete; + LinkProgram& operator=(LinkProgram&) = delete; + + amd_comgr_data_kind_t data_kind_; + amd_comgr_data_kind_t GetCOMGRDataKind(hipJitInputType input_type); + + // Linker Argumenets at hipLinkCreate + LinkArguments link_args_; + + // Spirv is bundled + bool is_bundled_ = false; + + // Private Data Members + amd_comgr_data_set_t link_input_; + std::vector link_options_; + static std::unordered_set linker_set_; + + bool AddLinkerDataImpl(std::vector& link_data, hipJitInputType input_type, + std::string& link_file_name); + + public: + LinkProgram(std::string name); + ~LinkProgram() { + amd::ScopedLock lock(lock_); + linker_set_.erase(this); + amd::Comgr::destroy_data_set(link_input_); + } + // Public Member Functions + bool AddLinkerOptions(unsigned int num_options, hipJitOption* options_ptr, + void** options_vals_ptr); + bool AddLinkerFile(std::string file_path, hipJitInputType input_type); + bool AddLinkerData(void* image_ptr, size_t image_size, std::string link_file_name, + hipJitInputType input_type); + bool LinkComplete(void** bin_out, size_t* size_out); + void AppendLinkerOptions() { AppendOptions(HIPRTC_LINK_OPTIONS_APPEND, &link_options_); } + static bool isLinkerValid(LinkProgram* link_program); +}; + + +} // namespace hip diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index 1ba5a6f89a..a5fef2b23c 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -588,6 +588,11 @@ global: hipGraphBatchMemOpNodeSetParams; hipGraphExecBatchMemOpNodeSetParams; hipEventRecordWithFlags; + hipLinkAddData; + hipLinkAddFile; + hipLinkComplete; + hipLinkCreate; + hipLinkDestroy; local: *; -} hip_6.2; +} hip_6.2; \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index ba775bdf17..5e6e136755 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -26,8 +26,10 @@ #include "platform/program.hpp" #include "hip_event.hpp" #include "hip_platform.hpp" +#include "hip_comgr_helper.hpp" namespace hip { + hipError_t ihipModuleLoadData(hipModule_t* module, const void* mmap_ptr, size_t mmap_size); extern hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim, @@ -944,4 +946,135 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const HIP_RETURN(err); } + + +hipError_t hipLinkAddData(hipLinkState_t hip_link_state, hipJitInputType input_type, void* image, + size_t image_size, const char* name, unsigned int num_options, + hipJitOption* options_ptr, void** option_values) { + + HIP_INIT_API(hipLinkAddData, hip_link_state, image, image_size, name, num_options, options_ptr, option_values); + + if (image == nullptr || image_size <= 0) { + HIP_RETURN(hipErrorInvalidImage); + } + + if (input_type == hipJitInputCubin || input_type == hipJitInputPtx || + input_type == hipJitInputFatBinary || input_type == hipJitInputObject || + input_type == hipJitInputLibrary || input_type == hipJitInputNvvm || + input_type == hipJitInputLLVMBitcode || input_type == hipJitInputLLVMBundledBitcode || + input_type == hipJitInputLLVMArchivesOfBundledBitcode ) { + HIP_RETURN(hipErrorInvalidValue); + } + + std::string input_name; + if (name) { + input_name = name; + } + + LinkProgram* hip_link_prog_ptr = + reinterpret_cast(hip_link_state); + + if (!LinkProgram::isLinkerValid(hip_link_prog_ptr)) { + HIP_RETURN(hipErrorInvalidHandle); + } + + if (!hip_link_prog_ptr->AddLinkerData(image, image_size, input_name, input_type)) { + HIP_RETURN(hipErrorInvalidConfiguration); + } + + HIP_RETURN(hipSuccess); +} + +hipError_t hipLinkAddFile(hipLinkState_t hip_link_state, hipJitInputType input_type, const char* file_path, + unsigned int num_options, hipJitOption* options_ptr, void** option_values) { + HIP_INIT_API(hipLinkAddFile, hip_link_state, input_type, file_path, num_options, options_ptr, option_values); + + if (hip_link_state == nullptr) { + HIP_RETURN(hipErrorInvalidHandle); + } + + if (input_type == hipJitInputCubin || input_type == hipJitInputPtx || + input_type == hipJitInputFatBinary || input_type == hipJitInputObject || + input_type == hipJitInputLibrary || input_type == hipJitInputNvvm || + input_type == hipJitInputLLVMBitcode || input_type == hipJitInputLLVMBundledBitcode || + input_type == hipJitInputLLVMArchivesOfBundledBitcode ) { + HIP_RETURN(hipErrorInvalidValue); + } + + LinkProgram* hip_link_prog_ptr = + reinterpret_cast(hip_link_state); + + if (!LinkProgram::isLinkerValid(hip_link_prog_ptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + if (!hip_link_prog_ptr->AddLinkerFile(std::string(file_path), input_type)) { + HIP_RETURN(hipErrorInvalidConfiguration); + } + + HIP_RETURN(hipSuccess); +} + +hipError_t hipLinkCreate(unsigned int num_options, hipJitOption* options_ptr, + void** options_vals_pptr, hipLinkState_t* hip_link_state_ptr) { + HIP_INIT_API(hipLinkCreate, num_options, options_ptr, options_vals_pptr, hip_link_state_ptr); + + if (hip_link_state_ptr == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + if (num_options != 0) { + for (int i = 0; i < num_options; i++) { + if (options_ptr == nullptr || options_vals_pptr == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + } + } + + std::string name("LinkerProgram"); + LinkProgram* hip_link_prog_ptr = new LinkProgram(name); + if (!hip_link_prog_ptr->AddLinkerOptions(num_options, options_ptr, options_vals_pptr)) { + HIP_RETURN(hipErrorInvalidConfiguration); + } + + *hip_link_state_ptr = reinterpret_cast(hip_link_prog_ptr); + + HIP_RETURN(hipSuccess); +} + +hipError_t hipLinkComplete(hipLinkState_t hip_link_state, void** bin_out, size_t* size_out) { + HIP_INIT_API(hipLinkComplete, hip_link_state, bin_out, size_out); + + if (bin_out == nullptr || size_out == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + LinkProgram* hip_link_prog_ptr = + reinterpret_cast(hip_link_state); + + if (!LinkProgram::isLinkerValid(hip_link_prog_ptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + if (!hip_link_prog_ptr->LinkComplete(bin_out, size_out)) { + HIP_RETURN(hipErrorInvalidConfiguration); + } + + HIP_RETURN(hipSuccess); +} + +hipError_t hipLinkDestroy(hipLinkState_t hip_link_state) { + HIP_INIT_API(hipLinkDestroy, hip_link_state); + + LinkProgram* hip_link_prog_ptr = + reinterpret_cast(hip_link_state); + + if (!LinkProgram::isLinkerValid(hip_link_prog_ptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + + delete hip_link_prog_ptr; + HIP_RETURN(hipSuccess); +} + } // namespace hip diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index cc34a40804..4d96eaec06 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -1261,6 +1261,31 @@ hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned return hip::GetHipDispatchTable()->hipModuleLoadDataEx_fn(module, image, numOptions, options, optionValues); } + +hipError_t hipLinkAddData(hipLinkState_t state, hipJitInputType type, void* data, size_t size, const char* name, + unsigned int numOptions, hipJitOption* options, void** optionValues) { + return hip::GetHipDispatchTable()->hipLinkAddData_fn(state, type, data, size, name, numOptions, + options, optionValues); +} + +hipError_t hipLinkAddFile(hipLinkState_t state, hipJitInputType type, const char* path, + unsigned int numOptions, hipJitOption* options, void** optionValues) { + return hip::GetHipDispatchTable()->hipLinkAddFile_fn(state, type, path, numOptions, options, + optionValues); +} + +hipError_t hipLinkComplete(hipLinkState_t state, void** hipBinOut, size_t* sizeOut) { + return hip::GetHipDispatchTable()->hipLinkComplete_fn(state, hipBinOut, sizeOut); +} + +hipError_t hipLinkCreate(unsigned int numOptions, hipJitOption* options, void** optionValues, hipLinkState_t* stateOut) { + return hip::GetHipDispatchTable()->hipLinkCreate_fn(numOptions, options, optionValues, stateOut); +} + +hipError_t hipLinkDestroy(hipLinkState_t state) { + return hip::GetHipDispatchTable()->hipLinkDestroy_fn(state); +} + extern "C" hipError_t hipModuleOccupancyMaxActiveBlocksPerMultiprocessor( int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk) { return hip::GetHipDispatchTable()->hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_fn( diff --git a/projects/clr/hipamd/src/hiprtc/CMakeLists.txt b/projects/clr/hipamd/src/hiprtc/CMakeLists.txt index 49e7672994..18f266b658 100644 --- a/projects/clr/hipamd/src/hiprtc/CMakeLists.txt +++ b/projects/clr/hipamd/src/hiprtc/CMakeLists.txt @@ -73,7 +73,7 @@ if(BUILD_SHARED_LIBS) endif() endif() -target_sources(hiprtc PRIVATE hiprtc.cpp hiprtcComgrHelper.cpp hiprtcInternal.cpp) +target_sources(hiprtc PRIVATE hiprtc.cpp ../hip_comgr_helper.cpp hiprtcInternal.cpp) set_target_properties(hiprtc PROPERTIES CXX_STANDARD 17 @@ -243,7 +243,7 @@ target_compile_definitions(hiprtc PRIVATE __HIP_ENABLE_RTC) if(NOT WIN32) if (BUILD_SHARED_LIBS) - target_sources(amdhip64 PRIVATE hiprtc.cpp hiprtcComgrHelper.cpp hiprtcInternal.cpp) + target_sources(amdhip64 PRIVATE hiprtc.cpp ../hip_comgr_helper.cpp hiprtcInternal.cpp) endif() endif() diff --git a/projects/clr/hipamd/src/hiprtc/hiprtc.cpp b/projects/clr/hipamd/src/hiprtc/hiprtc.cpp index 7ef2aff23a..daa2ce9a12 100644 --- a/projects/clr/hipamd/src/hiprtc/hiprtc.cpp +++ b/projects/clr/hipamd/src/hiprtc/hiprtc.cpp @@ -20,6 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include "../hip_comgr_helper.hpp" #include #include "hiprtcInternal.hpp" @@ -90,7 +91,7 @@ hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, const cha if (name == nullptr || strlen(name) == 0) { progName = "CompileSourceXXXXXX"; - hiprtc::helpers::GenerateUniqueFileName(progName); + hip::helpers::GenerateUniqueFileName(progName); } if (!rtcProgram->addSource(std::string(src), progName)) { @@ -306,8 +307,9 @@ hiprtcResult hiprtcLinkCreate(unsigned int num_options, hiprtcJIT_option* option } std::string name("LinkerProgram"); - hiprtc::RTCLinkProgram* rtc_link_prog_ptr = new hiprtc::RTCLinkProgram(name); - if (!rtc_link_prog_ptr->AddLinkerOptions(num_options, options_ptr, options_vals_pptr)) { + hip::LinkProgram* rtc_link_prog_ptr = new hip::LinkProgram(name); + if (!rtc_link_prog_ptr->AddLinkerOptions(num_options, options_ptr, + options_vals_pptr)) { HIPRTC_RETURN(HIPRTC_ERROR_INVALID_OPTION); } @@ -327,14 +329,15 @@ hiprtcResult hiprtcLinkAddFile(hiprtcLinkState hip_link_state, hiprtcJITInputTyp if (input_type == HIPRTC_JIT_INPUT_CUBIN || input_type == HIPRTC_JIT_INPUT_PTX || input_type == HIPRTC_JIT_INPUT_FATBINARY || input_type == HIPRTC_JIT_INPUT_OBJECT || - input_type == HIPRTC_JIT_INPUT_LIBRARY || input_type == HIPRTC_JIT_INPUT_NVVM) { + input_type == HIPRTC_JIT_INPUT_LIBRARY || input_type == HIPRTC_JIT_INPUT_NVVM || + input_type == HIPRTC_JIT_INPUT_SPIRV ) { HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); } - hiprtc::RTCLinkProgram* rtc_link_prog_ptr = - reinterpret_cast(hip_link_state); + hip::LinkProgram* rtc_link_prog_ptr = + reinterpret_cast(hip_link_state); - if (!hiprtc::RTCLinkProgram::isLinkerValid(rtc_link_prog_ptr)) { + if (!hip::LinkProgram::isLinkerValid(rtc_link_prog_ptr)) { HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); } @@ -357,7 +360,8 @@ hiprtcResult hiprtcLinkAddData(hiprtcLinkState hip_link_state, hiprtcJITInputTyp if (input_type == HIPRTC_JIT_INPUT_CUBIN || input_type == HIPRTC_JIT_INPUT_PTX || input_type == HIPRTC_JIT_INPUT_FATBINARY || input_type == HIPRTC_JIT_INPUT_OBJECT || - input_type == HIPRTC_JIT_INPUT_LIBRARY || input_type == HIPRTC_JIT_INPUT_NVVM) { + input_type == HIPRTC_JIT_INPUT_LIBRARY || input_type == HIPRTC_JIT_INPUT_NVVM || + input_type == HIPRTC_JIT_INPUT_SPIRV) { HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); } @@ -366,10 +370,10 @@ hiprtcResult hiprtcLinkAddData(hiprtcLinkState hip_link_state, hiprtcJITInputTyp input_name = name; } - hiprtc::RTCLinkProgram* rtc_link_prog_ptr = - reinterpret_cast(hip_link_state); + hip::LinkProgram* rtc_link_prog_ptr = + reinterpret_cast(hip_link_state); - if (!hiprtc::RTCLinkProgram::isLinkerValid(rtc_link_prog_ptr)) { + if (!hip::LinkProgram::isLinkerValid(rtc_link_prog_ptr)) { HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); } @@ -387,10 +391,10 @@ hiprtcResult hiprtcLinkComplete(hiprtcLinkState hip_link_state, void** bin_out, HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); } - hiprtc::RTCLinkProgram* rtc_link_prog_ptr = - reinterpret_cast(hip_link_state); + hip::LinkProgram* rtc_link_prog_ptr = + reinterpret_cast(hip_link_state); - if (!hiprtc::RTCLinkProgram::isLinkerValid(rtc_link_prog_ptr)) { + if (!hip::LinkProgram::isLinkerValid(rtc_link_prog_ptr)) { HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); } @@ -404,10 +408,10 @@ hiprtcResult hiprtcLinkComplete(hiprtcLinkState hip_link_state, void** bin_out, hiprtcResult hiprtcLinkDestroy(hiprtcLinkState hip_link_state) { HIPRTC_INIT_API(hip_link_state); - hiprtc::RTCLinkProgram* rtc_link_prog_ptr = - reinterpret_cast(hip_link_state); + hip::LinkProgram* rtc_link_prog_ptr = + reinterpret_cast(hip_link_state); - if (!hiprtc::RTCLinkProgram::isLinkerValid(rtc_link_prog_ptr)) { + if (!hip::LinkProgram::isLinkerValid(rtc_link_prog_ptr)) { HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT); } diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp b/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp deleted file mode 100644 index 31b616cc0f..0000000000 --- a/projects/clr/hipamd/src/hiprtc/hiprtcComgrHelper.hpp +++ /dev/null @@ -1,65 +0,0 @@ -/* -Copyright (c) 2022 - Present Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -#pragma once - -#include -#include - -#include "vdi_common.hpp" -#include "rocclr/utils/debug.hpp" -#include "device/comgrctx.hpp" - -namespace hiprtc { -namespace helpers { -bool UnbundleBitCode(const std::vector& bundled_bit_code, const std::string& isa, - size_t& co_offset, size_t& co_size); -bool addCodeObjData(amd_comgr_data_set_t& input, const std::vector& source, - const std::string& name, const amd_comgr_data_kind_t type); -bool extractBuildLog(amd_comgr_data_set_t dataSet, std::string& buildLog); -bool extractByteCodeBinary(const amd_comgr_data_set_t inDataSet, - const amd_comgr_data_kind_t dataKind, std::vector& bin); -bool createAction(amd_comgr_action_info_t& action, std::vector& options, - const std::string& isa, - const amd_comgr_language_t lang = AMD_COMGR_LANGUAGE_NONE); -bool compileToExecutable(const amd_comgr_data_set_t compileInputs, const std::string& isa, - std::vector& compileOptions, - std::vector& linkOptions, std::string& buildLog, - std::vector& exe); -bool compileToBitCode(const amd_comgr_data_set_t compileInputs, const std::string& isa, - std::vector& compileOptions, std::string& buildLog, - std::vector& LLVMBitcode); -bool linkLLVMBitcode(const amd_comgr_data_set_t linkInputs, const std::string& isa, - std::vector& linkOptions, std::string& buildLog, - std::vector& LinkedLLVMBitcode); -bool createExecutable(const amd_comgr_data_set_t linkInputs, const std::string& isa, - std::vector& exeOptions, std::string& buildLog, - std::vector& executable); -bool dumpIsaFromBC(const amd_comgr_data_set_t isaInputs, const std::string& isa, - std::vector& exeOptions, std::string name, std::string& buildLog); -bool demangleName(const std::string& mangledName, std::string& demangledName); -std::string handleMangledName(std::string loweredName); -bool fillMangledNames(std::vector& executable, - std::map& mangledNames, bool isBitcode); -void GenerateUniqueFileName(std::string& name); -} // namespace helpers -} // namespace hiprtc diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp index 7427b9c1eb..95537d3446 100644 --- a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp +++ b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.cpp @@ -31,131 +31,12 @@ THE SOFTWARE. #include "vdi_common.hpp" #include "rocclr/utils/flags.hpp" +#include "../hip_comgr_helper.hpp" + namespace hiprtc { -using namespace helpers; -std::unordered_setRTCLinkProgram::linker_set_; - -std::vector getLinkOptions(const LinkArguments& args) { - std::vector res; - - { // process optimization level - std::string opt("-O"); - opt += std::to_string(args.optimization_level_); - res.push_back(opt); - } - - const auto irArgCount = args.linker_ir2isa_args_count_; - if (irArgCount > 0) { - res.reserve(irArgCount); - const auto irArg = args.linker_ir2isa_args_; - for (size_t i = 0; i < irArgCount; i++) { - res.emplace_back(std::string(irArg[i])); - } - } - return res; -} - -// RTC Program Member Functions -RTCProgram::RTCProgram(std::string name) : name_(name) { - constexpr bool kComgrVersioned = true; - std::call_once(amd::Comgr::initialized, amd::Comgr::LoadLib, kComgrVersioned); - if (amd::Comgr::create_data_set(&exec_input_) != AMD_COMGR_STATUS_SUCCESS) { - crashWithMessage("Failed to allocate internal hiprtc structure"); - } -} - -bool RTCProgram::findIsa() { - -#ifdef BUILD_SHARED_LIBS - const char* libName; -#ifdef _WIN32 - std::string dll_name = std::string("amdhip64_" + std::to_string(HIP_VERSION_MAJOR) + ".dll"); - libName = dll_name.c_str(); -#else - libName = "libamdhip64.so"; -#endif - - void* handle = amd::Os::loadLibrary(libName); - - if (!handle) { - LogInfo("hip runtime failed to load using dlopen"); - build_log_ += - "hip runtime failed to load.\n" - "Error: Please provide architecture for which code is to be " - "generated.\n"; - return false; - } - - void* sym_hipGetDevice = amd::Os::getSymbol(handle, "hipGetDevice"); - void* sym_hipGetDeviceProperties = - amd::Os::getSymbol(handle, "hipGetDevicePropertiesR0600"); // Try to find the new symbol - if (sym_hipGetDeviceProperties == nullptr) { - sym_hipGetDeviceProperties = - amd::Os::getSymbol(handle, "hipGetDeviceProperties"); // Fall back to old one - } - - if (sym_hipGetDevice == nullptr || sym_hipGetDeviceProperties == nullptr) { - LogInfo("ISA cannot be found to dlsym failure"); - build_log_ += - "ISA cannot be found from hip runtime.\n" - "Error: Please provide architecture for which code is to be " - "generated.\n"; - return false; - } - - hipError_t (*dyn_hipGetDevice)(int*) = reinterpret_cast(sym_hipGetDevice); - - hipError_t (*dyn_hipGetDeviceProperties)(hipDeviceProp_t*, int) = - reinterpret_cast(sym_hipGetDeviceProperties); - - int device; - hipError_t status = dyn_hipGetDevice(&device); - if (status != hipSuccess) { - return false; - } - hipDeviceProp_t props; - status = dyn_hipGetDeviceProperties(&props, device); - if (status != hipSuccess) { - return false; - } - isa_ = "amdgcn-amd-amdhsa--"; - isa_.append(props.gcnArchName); - - amd::Os::unloadLibrary(handle); - return true; - -#else - int device; - hipError_t status = hipGetDevice(&device); - if (status != hipSuccess) { - return false; - } - hipDeviceProp_t props; - status = hipGetDeviceProperties(&props, device); - if (status != hipSuccess) { - return false; - } - isa_ = "amdgcn-amd-amdhsa--"; - isa_.append(props.gcnArchName); - - return true; -#endif -} // RTC Compile Program Member Functions -void RTCProgram::AppendOptions(const std::string app_env_var, std::vector* options) { - if (options == nullptr) { - LogError("Append options passed is nullptr."); - return; - } - - std::stringstream ss(app_env_var); - std::istream_iterator begin{ss}, end; - options->insert(options->end(), begin, end); -} - -// RTC Compile Program Member Functions -RTCCompileProgram::RTCCompileProgram(std::string name_) : RTCProgram(name_), fgpu_rdc_(false) { +RTCCompileProgram::RTCCompileProgram(std::string name_) : hip::RTCProgram(name_), fgpu_rdc_(false) { if ((amd::Comgr::create_data_set(&compile_input_) != AMD_COMGR_STATUS_SUCCESS) || (amd::Comgr::create_data_set(&link_input_) != AMD_COMGR_STATUS_SUCCESS)) { crashWithMessage("Failed to allocate internal hiprtc structure"); @@ -209,7 +90,7 @@ bool RTCCompileProgram::addSource(const std::string& source, const std::string& // objects bool RTCCompileProgram::addSource_impl() { std::vector vsource(source_code_.begin(), source_code_.end()); - if (!addCodeObjData(compile_input_, vsource, source_name_, AMD_COMGR_DATA_KIND_SOURCE)) { + if (!hip::helpers::addCodeObjData(compile_input_, vsource, source_name_, AMD_COMGR_DATA_KIND_SOURCE)) { return false; } return true; @@ -221,7 +102,7 @@ bool RTCCompileProgram::addHeader(const std::string& source, const std::string& return false; } std::vector vsource(source.begin(), source.end()); - if (!addCodeObjData(compile_input_, vsource, name, AMD_COMGR_DATA_KIND_INCLUDE)) { + if (!hip::helpers::addCodeObjData(compile_input_, vsource, name, AMD_COMGR_DATA_KIND_INCLUDE)) { return false; } return true; @@ -230,7 +111,7 @@ bool RTCCompileProgram::addHeader(const std::string& source, const std::string& bool RTCCompileProgram::addBuiltinHeader() { std::vector source(__hipRTC_header, __hipRTC_header + __hipRTC_header_size); std::string name{"hiprtc_runtime.h"}; - if (!addCodeObjData(compile_input_, source, name, AMD_COMGR_DATA_KIND_INCLUDE)) { + if (!hip::helpers::addCodeObjData(compile_input_, source, name, AMD_COMGR_DATA_KIND_INCLUDE)) { return false; } return true; @@ -301,9 +182,6 @@ bool RTCCompileProgram::transformOptions(std::vector& compile_optio return findIsa(); } -// HIPRTC Program lock -amd::Monitor RTCProgram::lock_(true); - bool RTCCompileProgram::compile(const std::vector& options, bool fgpu_rdc) { if (!addSource_impl()) { LogError("Error in hiprtc: unable to add source code"); @@ -323,13 +201,13 @@ bool RTCCompileProgram::compile(const std::vector& options, bool fg } if (fgpu_rdc_) { - if (!compileToBitCode(compile_input_, isa_, compileOpts, build_log_, LLVMBitcode_)) { + if (!hip::helpers::compileToBitCode(compile_input_, isa_, compileOpts, build_log_, LLVMBitcode_)) { LogError("Error in hiprtc: unable to compile source to bitcode"); return false; } } else { LogInfo("Using the new path of comgr"); - if (!compileToExecutable(compile_input_, isa_, compileOpts, link_options_, build_log_, + if (!hip::helpers::compileToExecutable(compile_input_, isa_, compileOpts, link_options_, build_log_, executable_)) { LogError("Failing to compile to realloc"); return false; @@ -338,7 +216,7 @@ bool RTCCompileProgram::compile(const std::vector& options, bool fg if (!mangled_names_.empty()) { auto& compile_step_output = fgpu_rdc_ ? LLVMBitcode_ : executable_; - if (!fillMangledNames(compile_step_output, mangled_names_, fgpu_rdc_)) { + if (!hip::helpers::fillMangledNames(compile_step_output, mangled_names_, fgpu_rdc_)) { LogError("Error in hiprtc: unable to fill mangled names"); return false; } @@ -410,273 +288,4 @@ bool RTCCompileProgram::GetBitcodeSize(size_t* bitcode_size) { *bitcode_size = LLVMBitcode_.size(); return true; } - -// RTC Link Program Member Functions -RTCLinkProgram::RTCLinkProgram(std::string name) : RTCProgram(name) { - if (amd::Comgr::create_data_set(&link_input_) != AMD_COMGR_STATUS_SUCCESS) { - crashWithMessage("Failed to allocate internal hiprtc structure"); - } - amd::ScopedLock lock(lock_); - linker_set_.insert(this); -} - -bool RTCLinkProgram::isLinkerValid(RTCLinkProgram* link_program) { - amd::ScopedLock lock(lock_); - if (linker_set_.find(link_program) == linker_set_.end()) { - return false; - } - return true; -} - -bool RTCLinkProgram::AddLinkerOptions(unsigned int num_options, hiprtcJIT_option* options_ptr, - void** options_vals_ptr) { - for (size_t opt_idx = 0; opt_idx < num_options; ++opt_idx) { - switch (options_ptr[opt_idx]) { - case HIPRTC_JIT_MAX_REGISTERS: - link_args_.max_registers_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_THREADS_PER_BLOCK: - link_args_.threads_per_block_ = - *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_WALL_TIME: - link_args_.wall_time_ = *(reinterpret_cast(options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_INFO_LOG_BUFFER: { - if (options_vals_ptr[opt_idx] == nullptr) { - LogError("Options value can not be nullptr"); - return false; - } - link_args_.info_log_ = (reinterpret_cast(options_vals_ptr[opt_idx])); - break; - } - case HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES: - link_args_.info_log_size_ = (reinterpret_cast(options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_ERROR_LOG_BUFFER: { - if (options_vals_ptr[opt_idx] == nullptr) { - LogError("Options value can not be nullptr"); - return false; - } - link_args_.error_log_ = reinterpret_cast(options_vals_ptr[opt_idx]); - break; - } - case HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES: - link_args_.error_log_size_ = (reinterpret_cast(options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_OPTIMIZATION_LEVEL: - link_args_.optimization_level_ = - *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_TARGET_FROM_HIPCONTEXT: - link_args_.target_from_hip_context_ = - *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_TARGET: - link_args_.jit_target_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_FALLBACK_STRATEGY: - link_args_.fallback_strategy_ = - *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_GENERATE_DEBUG_INFO: - link_args_.generate_debug_info_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_LOG_VERBOSE: - link_args_.log_verbose_ = reinterpret_cast(options_vals_ptr[opt_idx]); - break; - case HIPRTC_JIT_GENERATE_LINE_INFO: - link_args_.generate_line_info_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_CACHE_MODE: - link_args_.cache_mode_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_NEW_SM3X_OPT: - link_args_.sm3x_opt_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_FAST_COMPILE: - link_args_.fast_compile_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_GLOBAL_SYMBOL_NAMES: { - if (options_vals_ptr[opt_idx] == nullptr) { - LogError("Options value can not be nullptr"); - return false; - } - link_args_.global_symbol_names_ = reinterpret_cast(options_vals_ptr[opt_idx]); - break; - } - case HIPRTC_JIT_GLOBAL_SYMBOL_ADDRESS: { - if (options_vals_ptr[opt_idx] == nullptr) { - LogError("Options value can not be nullptr"); - return false; - } - link_args_.global_symbol_addresses_ = reinterpret_cast(options_vals_ptr[opt_idx]); - break; - } - case HIPRTC_JIT_GLOBAL_SYMBOL_COUNT: - link_args_.global_symbol_count_ = - *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_LTO: - link_args_.lto_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_FTZ: - link_args_.ftz_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_PREC_DIV: - link_args_.prec_div_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_PREC_SQRT: - link_args_.prec_sqrt_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_FMA: - link_args_.fma_ = *(reinterpret_cast(&options_vals_ptr[opt_idx])); - break; - case HIPRTC_JIT_IR_TO_ISA_OPT_EXT: { - if (options_vals_ptr[opt_idx] == nullptr) { - LogError("Options value can not be nullptr"); - return false; - } - link_args_.linker_ir2isa_args_ = reinterpret_cast(options_vals_ptr[opt_idx]); - break; - } - case HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT: - link_args_.linker_ir2isa_args_count_ = reinterpret_cast(options_vals_ptr[opt_idx]); - break; - default: - break; - } - } - - return true; -} - -amd_comgr_data_kind_t RTCLinkProgram::GetCOMGRDataKind(hiprtcJITInputType input_type) { - amd_comgr_data_kind_t data_kind = AMD_COMGR_DATA_KIND_UNDEF; - - // Map the hiprtc input type to comgr data kind - switch (input_type) { - case HIPRTC_JIT_INPUT_LLVM_BITCODE: - data_kind = AMD_COMGR_DATA_KIND_BC; - break; - case HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE: - data_kind = - HIPRTC_USE_RUNTIME_UNBUNDLER ? AMD_COMGR_DATA_KIND_BC : AMD_COMGR_DATA_KIND_BC_BUNDLE; - break; - case HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE: - data_kind = AMD_COMGR_DATA_KIND_AR_BUNDLE; - break; - default: - LogError("Cannot find the corresponding comgr data kind"); - break; - } - - return data_kind; -} - -bool RTCLinkProgram::AddLinkerDataImpl(std::vector& link_data, hiprtcJITInputType input_type, - std::string& link_file_name) { - std::vector llvm_bitcode; - // If this is bundled bitcode then unbundle this. - if (HIPRTC_USE_RUNTIME_UNBUNDLER && input_type == HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE) { - if (!findIsa()) { - return false; - } - - size_t co_offset = 0; - size_t co_size = 0; - if (!UnbundleBitCode(link_data, isa_, co_offset, co_size)) { - LogError("Error in hiprtc: unable to unbundle the llvm bitcode"); - return false; - } - - llvm_bitcode.assign(link_data.begin() + co_offset, link_data.begin() + co_offset + co_size); - } else { - llvm_bitcode.assign(link_data.begin(), link_data.end()); - } - - amd_comgr_data_kind_t data_kind; - if ((data_kind = GetCOMGRDataKind(input_type)) == AMD_COMGR_DATA_KIND_UNDEF) { - LogError("Cannot find the correct COMGR data kind"); - return false; - } - - if (!addCodeObjData(link_input_, llvm_bitcode, link_file_name, data_kind)) { - LogError("Error in hiprtc: unable to add linked code object"); - return false; - } - - return true; -} - -bool RTCLinkProgram::AddLinkerFile(std::string file_path, hiprtcJITInputType input_type) { - std::ifstream file_stream{file_path, std::ios_base::in | std::ios_base::binary}; - if (!file_stream.good()) { - return false; - } - - file_stream.seekg(0, std::ios::end); - std::streampos file_size = file_stream.tellg(); - file_stream.seekg(0, std::ios::beg); - - // Read the file contents - std::vector link_file_info(file_size); - file_stream.read(link_file_info.data(), file_size); - - file_stream.close(); - - std::string link_file_name("LinkerProgram"); - - return AddLinkerDataImpl(link_file_info, input_type, link_file_name); -} - -bool RTCLinkProgram::AddLinkerData(void* image_ptr, size_t image_size, std::string link_file_name, - hiprtcJITInputType input_type) { - char* image_char_buf = reinterpret_cast(image_ptr); - std::vector bundled_llvm_bitcode(image_char_buf, image_char_buf + image_size); - - return AddLinkerDataImpl(bundled_llvm_bitcode, input_type, link_file_name); -} - -bool RTCLinkProgram::LinkComplete(void** bin_out, size_t* size_out) { - if (!findIsa()) { - return false; - } - - AppendLinkerOptions(); - - std::vector linked_llvm_bitcode; - if (!linkLLVMBitcode(link_input_, isa_, link_options_, build_log_, linked_llvm_bitcode)) { - LogError("Error in hiprtc: unable to add device libs to linked bitcode"); - return false; - } - - std::string linkedFileName = "LLVMBitcode.bc"; - if (!addCodeObjData(exec_input_, linked_llvm_bitcode, linkedFileName, AMD_COMGR_DATA_KIND_BC)) { - LogError("Error in hiprtc: unable to add linked bitcode"); - return false; - } - - std::vector exe_options = getLinkOptions(link_args_); - LogPrintfInfo("Exe options forwarded to compiler: %s", - [&]() { - std::string ret; - for (const auto& i : exe_options) { - ret += i; - ret += " "; - } - return ret; - }() - .c_str()); - if (!createExecutable(exec_input_, isa_, exe_options, build_log_, executable_)) { - LogPrintfInfo("Error in hiprtc: unable to create exectuable: %s", build_log_.c_str()); - return false; - } - - *size_out = executable_.size(); - *bin_out = executable_.data(); - - return true; -} - } // namespace hiprtc diff --git a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp index ed754d85d2..1ef807a1e5 100644 --- a/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp +++ b/projects/clr/hipamd/src/hiprtc/hiprtcInternal.hpp @@ -19,6 +19,9 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ + +#pragma once + #include #include #include @@ -36,6 +39,10 @@ THE SOFTWARE. #include "rocclr/utils/debug.hpp" #include "rocclr/utils/flags.hpp" #include "rocclr/utils/macros.hpp" +#include "vdi_common.hpp" +#include "device/comgrctx.hpp" +#include "../hip_comgr_helper.hpp" + #ifdef __HIP_ENABLE_RTC extern "C" { @@ -44,8 +51,6 @@ extern unsigned __hipRTC_header_size; } #endif -#include "hiprtcComgrHelper.hpp" - namespace hiprtc { namespace internal { template inline std::string ToString(T v) { @@ -112,29 +117,7 @@ struct Settings { bool offloadArchProvided{false}; }; -class RTCProgram { - protected: - // Lock and control variables - static amd::Monitor lock_; - static std::once_flag initialized_; - - RTCProgram(std::string name); - ~RTCProgram() { amd::Comgr::destroy_data_set(exec_input_); } - - // Member Functions - bool findIsa(); - static void AppendOptions(std::string app_env_var, std::vector* options); - - // Data Members - std::string name_; - std::string isa_; - std::string build_log_; - std::vector executable_; - - amd_comgr_data_set_t exec_input_; -}; - -class RTCCompileProgram : public RTCProgram { +class RTCCompileProgram : public hip::RTCProgram { // Private Data Members Settings settings_; @@ -195,103 +178,6 @@ class RTCCompileProgram : public RTCProgram { size_t getLogSize() const { return build_log_.size(); } }; -// Linker Arguments passed via hipLinkCreate -struct LinkArguments { - unsigned int max_registers_; - unsigned int threads_per_block_; - float wall_time_; - size_t info_log_size_; - char* info_log_; - size_t error_log_size_; - char* error_log_; - unsigned int optimization_level_; - unsigned int target_from_hip_context_; - unsigned int jit_target_; - unsigned int fallback_strategy_; - int generate_debug_info_; - long log_verbose_; - int generate_line_info_; - unsigned int cache_mode_; - bool sm3x_opt_; - bool fast_compile_; - const char** global_symbol_names_; - void** global_symbol_addresses_; - unsigned int global_symbol_count_; - int lto_; - int ftz_; - int prec_div_; - int prec_sqrt_; - int fma_; - const char** linker_ir2isa_args_; - size_t linker_ir2isa_args_count_; - - LinkArguments() - : max_registers_{0}, - threads_per_block_{0}, - wall_time_{0.0f}, - info_log_size_{0}, - info_log_{nullptr}, - error_log_size_{0}, - error_log_{nullptr}, - optimization_level_{3}, - target_from_hip_context_{0}, - jit_target_{0}, - fallback_strategy_{0}, - generate_debug_info_{0}, - log_verbose_{0}, - generate_line_info_{0}, - cache_mode_{0}, - sm3x_opt_{false}, - fast_compile_{false}, - global_symbol_names_{nullptr}, - global_symbol_addresses_{nullptr}, - global_symbol_count_{0}, - lto_{0}, - ftz_{0}, - prec_div_{1}, - prec_sqrt_{1}, - fma_{1}, - linker_ir2isa_args_{nullptr}, - linker_ir2isa_args_count_{0} {} -}; - -class RTCLinkProgram : public RTCProgram { - // Private Member Functions (forbid these function calls) - RTCLinkProgram() = delete; - RTCLinkProgram(RTCLinkProgram&) = delete; - RTCLinkProgram& operator=(RTCLinkProgram&) = delete; - - amd_comgr_data_kind_t GetCOMGRDataKind(hiprtcJITInputType input_type); - - // Linker Argumenets at hipLinkCreate - LinkArguments link_args_; - - // Private Data Members - amd_comgr_data_set_t link_input_; - std::vector link_options_; - static std::unordered_set linker_set_; - - bool AddLinkerDataImpl(std::vector& link_data, hiprtcJITInputType input_type, - std::string& link_file_name); - - public: - RTCLinkProgram(std::string name); - ~RTCLinkProgram() { - amd::ScopedLock lock(lock_); - linker_set_.erase(this); - amd::Comgr::destroy_data_set(link_input_); - } - // Public Member Functions - bool AddLinkerOptions(unsigned int num_options, hiprtcJIT_option* options_ptr, - void** options_vals_ptr); - bool AddLinkerFile(std::string file_path, hiprtcJITInputType input_type); - bool AddLinkerData(void* image_ptr, size_t image_size, std::string link_file_name, - hiprtcJITInputType input_type); - bool LinkComplete(void** bin_out, size_t* size_out); - void AppendLinkerOptions() { AppendOptions(HIPRTC_LINK_OPTIONS_APPEND, &link_options_); } - static bool isLinkerValid(RTCLinkProgram* link_program); -}; - // Thread Local Storage Variables Aggregator Class class TlsAggregator { public: