SWDEV-489106 - Linker API addition to runtime
1) Add Linker APIs to runtime to support SPIRV linking 2) Migrate Internal implementations to runtime and share with rtc 3) Add Support to bundled and unbundled SPIRV Code object linking. Change-Id: Ic1fd4431f842a208a2468e8aec54a65b5fa6b0e3
Este cometimento está contido em:
Rahul Manocha
cometido por
Rahul Manocha
Rahul Manocha
ascendente
9faaf20aae
cometimento
5930f047bb
@@ -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
|
||||
|
||||
// ******************************************************************************************* //
|
||||
//
|
||||
|
||||
@@ -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";
|
||||
|
||||
@@ -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()
|
||||
|
||||
#############################
|
||||
|
||||
@@ -485,3 +485,8 @@ hipGraphBatchMemOpNodeGetParams
|
||||
hipGraphBatchMemOpNodeSetParams
|
||||
hipGraphExecBatchMemOpNodeSetParams
|
||||
hipEventRecordWithFlags
|
||||
hipLinkAddData
|
||||
hipLinkAddFile
|
||||
hipLinkComplete
|
||||
hipLinkCreate
|
||||
hipLinkDestroy
|
||||
@@ -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(<table>, <functor>, 8)
|
||||
//
|
||||
// HIP_ENFORCE_ABI_VERSIONING(<table>, 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 "
|
||||
|
||||
@@ -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 <io.h>
|
||||
#endif
|
||||
|
||||
#include "../amd_hsa_elf.hpp"
|
||||
#include "../src/amd_hsa_elf.hpp"
|
||||
|
||||
namespace hiprtc {
|
||||
namespace hip {
|
||||
|
||||
std::unordered_set<LinkProgram*> 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<char>& 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<char>& source, const std::string& isa,
|
||||
std::vector<std::string>& linkOptions, std::string& buildLog,
|
||||
std::vector<char>& 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<std::string>& linkOptions, std::string& buildLog,
|
||||
std::vector<char>& 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<std::string>& linkOptions, std::string& buildLog,
|
||||
std::vector<char>& 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<std::string>& exeOptions, std::string& buildLog,
|
||||
std::vector<char>& executable) {
|
||||
std::vector<char>& 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<char>& dataVec, std::map<std::string, std::str
|
||||
}
|
||||
|
||||
} // namespace helpers
|
||||
} // namespace hiprtc
|
||||
|
||||
std::vector<std::string> getLinkOptions(const LinkArguments& args) {
|
||||
std::vector<std::string> 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<hipError_t (*)(int*)>(sym_hipGetDevice);
|
||||
|
||||
hipError_t (*dyn_hipGetDeviceProperties)(hipDeviceProp_t*, int) =
|
||||
reinterpret_cast<hipError_t (*)(hipDeviceProp_t*, int)>(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<std::string>* options) {
|
||||
if (options == nullptr) {
|
||||
LogError("Append options passed is nullptr.");
|
||||
return;
|
||||
}
|
||||
|
||||
std::stringstream ss(app_env_var);
|
||||
std::istream_iterator<std::string> 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<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionThreadsPerBlock:
|
||||
link_args_.threads_per_block_ =
|
||||
*(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionWallTime:
|
||||
link_args_.wall_time_ = *(reinterpret_cast<long*>(options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionInfoLogBuffer: {
|
||||
link_args_.info_log_ = (reinterpret_cast<char*>(options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
}
|
||||
case hipJitOptionInfoLogBufferSizeBytes:
|
||||
link_args_.info_log_size_ = (reinterpret_cast<size_t>(options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionErrorLogBuffer: {
|
||||
link_args_.error_log_ = reinterpret_cast<char*>(options_vals_ptr[opt_idx]);
|
||||
break;
|
||||
}
|
||||
case hipJitOptionErrorLogBufferSizeBytes:
|
||||
link_args_.error_log_size_ = (reinterpret_cast<size_t>(options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionOptimizationLevel:
|
||||
link_args_.optimization_level_ =
|
||||
*(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionTargetFromContext:
|
||||
link_args_.target_from_hip_context_ =
|
||||
*(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionTarget:
|
||||
link_args_.jit_target_ = *(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionFallbackStrategy:
|
||||
link_args_.fallback_strategy_ =
|
||||
*(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionGenerateDebugInfo:
|
||||
link_args_.generate_debug_info_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionLogVerbose:
|
||||
link_args_.log_verbose_ = reinterpret_cast<size_t>(options_vals_ptr[opt_idx]);
|
||||
break;
|
||||
case hipJitOptionGenerateLineInfo:
|
||||
link_args_.generate_line_info_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionCacheMode:
|
||||
link_args_.cache_mode_ = *(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionSm3xOpt:
|
||||
link_args_.sm3x_opt_ = *(reinterpret_cast<bool*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionFastCompile:
|
||||
link_args_.fast_compile_ = *(reinterpret_cast<bool*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionGlobalSymbolNames: {
|
||||
link_args_.global_symbol_names_ = reinterpret_cast<const char**>(options_vals_ptr[opt_idx]);
|
||||
break;
|
||||
}
|
||||
case hipJitOptionGlobalSymbolAddresses: {
|
||||
link_args_.global_symbol_addresses_ = reinterpret_cast<void**>(options_vals_ptr[opt_idx]);
|
||||
break;
|
||||
}
|
||||
case hipJitOptionGlobalSymbolCount:
|
||||
link_args_.global_symbol_count_ =
|
||||
*(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionLto:
|
||||
link_args_.lto_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionFtz:
|
||||
link_args_.ftz_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionPrecDiv:
|
||||
link_args_.prec_div_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionPrecSqrt:
|
||||
link_args_.prec_sqrt_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionFma:
|
||||
link_args_.fma_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case hipJitOptionPositionIndependentCode:
|
||||
link_args_.pic_ = *(reinterpret_cast<int *>(&options_vals_ptr[opt_idx]));
|
||||
case hipJitOptionMinCTAPerSM:
|
||||
link_args_.min_cta_per_sm_ = *(reinterpret_cast<int *>(&options_vals_ptr[opt_idx]));
|
||||
case hipJitOptionMaxThreadsPerBlock:
|
||||
link_args_.max_threads_per_block_ = *(reinterpret_cast<int *>(&options_vals_ptr[opt_idx]));
|
||||
case hipJitOptionOverrideDirectiveValues:
|
||||
link_args_.override_directive_values_ = *(reinterpret_cast<int *>(&options_vals_ptr[opt_idx]));
|
||||
case hipJitOptionIRtoISAOptExt: {
|
||||
link_args_.linker_ir2isa_args_ = reinterpret_cast<const char**>(options_vals_ptr[opt_idx]);
|
||||
break;
|
||||
}
|
||||
case hipJitOptionIRtoISAOptCountExt:
|
||||
link_args_.linker_ir2isa_args_count_ = reinterpret_cast<size_t>(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<char>& link_data, hipJitInputType input_type,
|
||||
std::string& link_file_name) {
|
||||
std::vector<char> 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<char> 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<char*>(image_ptr);
|
||||
std::vector<char> 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<char> 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<char> 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<std::string> 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
|
||||
@@ -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 <vector>
|
||||
#include <string>
|
||||
|
||||
|
||||
#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<char>& 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<char>& 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<char>& bin);
|
||||
bool createAction(amd_comgr_action_info_t& action, std::vector<std::string>& 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<std::string>& compileOptions,
|
||||
std::vector<std::string>& linkOptions, std::string& buildLog,
|
||||
std::vector<char>& exe);
|
||||
bool compileToBitCode(const amd_comgr_data_set_t compileInputs, const std::string& isa,
|
||||
std::vector<std::string>& compileOptions, std::string& buildLog,
|
||||
std::vector<char>& LLVMBitcode);
|
||||
bool linkLLVMBitcode(const amd_comgr_data_set_t linkInputs, const std::string& isa,
|
||||
std::vector<std::string>& linkOptions, std::string& buildLog,
|
||||
std::vector<char>& LinkedLLVMBitcode);
|
||||
bool createExecutable(const amd_comgr_data_set_t linkInputs, const std::string& isa,
|
||||
std::vector<std::string>& exeOptions, std::string& buildLog,
|
||||
std::vector<char>& executable, bool spirv_bc = false);
|
||||
bool convertSPIRVToLLVMBC(const amd_comgr_data_set_t linkInputs, const std::string& isa,
|
||||
std::vector<std::string>& linkOptions, std::string& buildLog,
|
||||
std::vector<char>& linkedSPIRVBitcode);
|
||||
bool dumpIsaFromBC(const amd_comgr_data_set_t isaInputs, const std::string& isa,
|
||||
std::vector<std::string>& 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<char>& executable,
|
||||
std::map<std::string, std::string>& mangledNames, bool isBitcode);
|
||||
void GenerateUniqueFileName(std::string& name);
|
||||
|
||||
bool CheckIfBundled(std::vector<char>& llvm_bitcode);
|
||||
|
||||
bool UnbundleUsingComgr(std::vector<char>& source, const std::string& isa,
|
||||
std::vector<std::string>& linkOptions, std::string& buildLog,
|
||||
std::vector<char>& 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<std::string>* options);
|
||||
|
||||
// Data Members
|
||||
std::string name_;
|
||||
std::string isa_;
|
||||
std::string build_log_;
|
||||
std::vector<char> 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<std::string> link_options_;
|
||||
static std::unordered_set<LinkProgram*> linker_set_;
|
||||
|
||||
bool AddLinkerDataImpl(std::vector<char>& 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
|
||||
@@ -588,6 +588,11 @@ global:
|
||||
hipGraphBatchMemOpNodeSetParams;
|
||||
hipGraphExecBatchMemOpNodeSetParams;
|
||||
hipEventRecordWithFlags;
|
||||
hipLinkAddData;
|
||||
hipLinkAddFile;
|
||||
hipLinkComplete;
|
||||
hipLinkCreate;
|
||||
hipLinkDestroy;
|
||||
local:
|
||||
*;
|
||||
} hip_6.2;
|
||||
} hip_6.2;
|
||||
@@ -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<LinkProgram*>(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<LinkProgram*>(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<hipLinkState_t>(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<LinkProgram*>(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<LinkProgram*>(hip_link_state);
|
||||
|
||||
if (!LinkProgram::isLinkerValid(hip_link_prog_ptr)) {
|
||||
HIP_RETURN(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
delete hip_link_prog_ptr;
|
||||
HIP_RETURN(hipSuccess);
|
||||
}
|
||||
|
||||
} // namespace hip
|
||||
|
||||
@@ -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(
|
||||
|
||||
@@ -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()
|
||||
|
||||
|
||||
@@ -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 <hip/hiprtc.h>
|
||||
#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<hiprtc::RTCLinkProgram*>(hip_link_state);
|
||||
hip::LinkProgram* rtc_link_prog_ptr =
|
||||
reinterpret_cast<hip::LinkProgram*>(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<hiprtc::RTCLinkProgram*>(hip_link_state);
|
||||
hip::LinkProgram* rtc_link_prog_ptr =
|
||||
reinterpret_cast<hip::LinkProgram*>(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<hiprtc::RTCLinkProgram*>(hip_link_state);
|
||||
hip::LinkProgram* rtc_link_prog_ptr =
|
||||
reinterpret_cast<hip::LinkProgram*>(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<hiprtc::RTCLinkProgram*>(hip_link_state);
|
||||
hip::LinkProgram* rtc_link_prog_ptr =
|
||||
reinterpret_cast<hip::LinkProgram*>(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);
|
||||
}
|
||||
|
||||
|
||||
@@ -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 <vector>
|
||||
#include <string>
|
||||
|
||||
#include "vdi_common.hpp"
|
||||
#include "rocclr/utils/debug.hpp"
|
||||
#include "device/comgrctx.hpp"
|
||||
|
||||
namespace hiprtc {
|
||||
namespace helpers {
|
||||
bool UnbundleBitCode(const std::vector<char>& 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<char>& 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<char>& bin);
|
||||
bool createAction(amd_comgr_action_info_t& action, std::vector<std::string>& 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<std::string>& compileOptions,
|
||||
std::vector<std::string>& linkOptions, std::string& buildLog,
|
||||
std::vector<char>& exe);
|
||||
bool compileToBitCode(const amd_comgr_data_set_t compileInputs, const std::string& isa,
|
||||
std::vector<std::string>& compileOptions, std::string& buildLog,
|
||||
std::vector<char>& LLVMBitcode);
|
||||
bool linkLLVMBitcode(const amd_comgr_data_set_t linkInputs, const std::string& isa,
|
||||
std::vector<std::string>& linkOptions, std::string& buildLog,
|
||||
std::vector<char>& LinkedLLVMBitcode);
|
||||
bool createExecutable(const amd_comgr_data_set_t linkInputs, const std::string& isa,
|
||||
std::vector<std::string>& exeOptions, std::string& buildLog,
|
||||
std::vector<char>& executable);
|
||||
bool dumpIsaFromBC(const amd_comgr_data_set_t isaInputs, const std::string& isa,
|
||||
std::vector<std::string>& 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<char>& executable,
|
||||
std::map<std::string, std::string>& mangledNames, bool isBitcode);
|
||||
void GenerateUniqueFileName(std::string& name);
|
||||
} // namespace helpers
|
||||
} // namespace hiprtc
|
||||
@@ -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_set<RTCLinkProgram*>RTCLinkProgram::linker_set_;
|
||||
|
||||
std::vector<std::string> getLinkOptions(const LinkArguments& args) {
|
||||
std::vector<std::string> 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<hipError_t (*)(int*)>(sym_hipGetDevice);
|
||||
|
||||
hipError_t (*dyn_hipGetDeviceProperties)(hipDeviceProp_t*, int) =
|
||||
reinterpret_cast<hipError_t (*)(hipDeviceProp_t*, int)>(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<std::string>* options) {
|
||||
if (options == nullptr) {
|
||||
LogError("Append options passed is nullptr.");
|
||||
return;
|
||||
}
|
||||
|
||||
std::stringstream ss(app_env_var);
|
||||
std::istream_iterator<std::string> 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<char> 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<char> 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<char> 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<std::string>& compile_optio
|
||||
return findIsa();
|
||||
}
|
||||
|
||||
// HIPRTC Program lock
|
||||
amd::Monitor RTCProgram::lock_(true);
|
||||
|
||||
bool RTCCompileProgram::compile(const std::vector<std::string>& 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<std::string>& 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<std::string>& 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<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_THREADS_PER_BLOCK:
|
||||
link_args_.threads_per_block_ =
|
||||
*(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_WALL_TIME:
|
||||
link_args_.wall_time_ = *(reinterpret_cast<long*>(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<char*>(options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
}
|
||||
case HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES:
|
||||
link_args_.info_log_size_ = (reinterpret_cast<size_t>(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<char*>(options_vals_ptr[opt_idx]);
|
||||
break;
|
||||
}
|
||||
case HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES:
|
||||
link_args_.error_log_size_ = (reinterpret_cast<size_t>(options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_OPTIMIZATION_LEVEL:
|
||||
link_args_.optimization_level_ =
|
||||
*(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_TARGET_FROM_HIPCONTEXT:
|
||||
link_args_.target_from_hip_context_ =
|
||||
*(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_TARGET:
|
||||
link_args_.jit_target_ = *(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_FALLBACK_STRATEGY:
|
||||
link_args_.fallback_strategy_ =
|
||||
*(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_GENERATE_DEBUG_INFO:
|
||||
link_args_.generate_debug_info_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_LOG_VERBOSE:
|
||||
link_args_.log_verbose_ = reinterpret_cast<size_t>(options_vals_ptr[opt_idx]);
|
||||
break;
|
||||
case HIPRTC_JIT_GENERATE_LINE_INFO:
|
||||
link_args_.generate_line_info_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_CACHE_MODE:
|
||||
link_args_.cache_mode_ = *(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_NEW_SM3X_OPT:
|
||||
link_args_.sm3x_opt_ = *(reinterpret_cast<bool*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_FAST_COMPILE:
|
||||
link_args_.fast_compile_ = *(reinterpret_cast<bool*>(&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<const char**>(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<void**>(options_vals_ptr[opt_idx]);
|
||||
break;
|
||||
}
|
||||
case HIPRTC_JIT_GLOBAL_SYMBOL_COUNT:
|
||||
link_args_.global_symbol_count_ =
|
||||
*(reinterpret_cast<unsigned int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_LTO:
|
||||
link_args_.lto_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_FTZ:
|
||||
link_args_.ftz_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_PREC_DIV:
|
||||
link_args_.prec_div_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_PREC_SQRT:
|
||||
link_args_.prec_sqrt_ = *(reinterpret_cast<int*>(&options_vals_ptr[opt_idx]));
|
||||
break;
|
||||
case HIPRTC_JIT_FMA:
|
||||
link_args_.fma_ = *(reinterpret_cast<int*>(&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<const char**>(options_vals_ptr[opt_idx]);
|
||||
break;
|
||||
}
|
||||
case HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT:
|
||||
link_args_.linker_ir2isa_args_count_ = reinterpret_cast<size_t>(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<char>& link_data, hiprtcJITInputType input_type,
|
||||
std::string& link_file_name) {
|
||||
std::vector<char> 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<char> 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<char*>(image_ptr);
|
||||
std::vector<char> 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<char> 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<std::string> 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
|
||||
|
||||
@@ -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 <hip/hip_runtime.h>
|
||||
#include <hip/hiprtc.h>
|
||||
#include <hip/hip_version.h>
|
||||
@@ -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 <typename T> 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<std::string>* options);
|
||||
|
||||
// Data Members
|
||||
std::string name_;
|
||||
std::string isa_;
|
||||
std::string build_log_;
|
||||
std::vector<char> 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<std::string> link_options_;
|
||||
static std::unordered_set<RTCLinkProgram*> linker_set_;
|
||||
|
||||
bool AddLinkerDataImpl(std::vector<char>& 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:
|
||||
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador