From 27ec19116d00b32dae338947f555ea5eeafca18e Mon Sep 17 00:00:00 2001 From: Rahul Manocha <153310294+manocharahul@users.noreply.github.com> Date: Tue, 7 Oct 2025 08:28:56 -0700 Subject: [PATCH] SWDEV-557828 - fix hip-tests on cuda (#1152) Co-authored-by: Rahul Manocha --- .../include/hip/amd_detail/hip_api_trace.hpp | 8 +- .../include/hip/amd_detail/hip_prof_str.h | 531 +++++++++--------- projects/clr/hipamd/src/hip_api_trace.cpp | 8 +- projects/clr/hipamd/src/hip_library.cpp | 10 +- .../clr/hipamd/src/hip_table_interface.cpp | 8 +- .../catch/include/hip_test_common.hh | 4 + .../catch/unit/library/CMakeLists.txt | 15 +- .../catch/unit/library/library_negative.cc | 2 +- .../catch/unit/library/loadlib_rtc.cc | 6 +- .../catch/unit/memory/hipHostRegister.cc | 3 +- projects/hip/include/hip/hip_runtime_api.h | 8 +- .../nvidia_detail/nvidia_hip_runtime_api.h | 28 +- .../include/rocprofiler-sdk/hip/api_args.h | 32 +- 13 files changed, 331 insertions(+), 332 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp index 84c7f3c3a2..051ee774cc 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp @@ -1088,15 +1088,15 @@ typedef hipError_t (*t_hipGetDriverEntryPoint_spt)(const char* symbol, void** fu unsigned long long flags, hipDriverEntryPointQueryResult* status); typedef hipError_t (*t_hipLibraryLoadData)(hipLibrary_t* library, const void* code, - hipJitOption** jitOptions, void** jitOptionsValues, + hipJitOption* jitOptions, void** jitOptionsValues, unsigned int numJitOptions, - hipLibraryOption** libraryOptions, + hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions); typedef hipError_t (*t_hipLibraryLoadFromFile)(hipLibrary_t* library, const char* fileName, - hipJitOption** jitOptions, void** jitOptionsValues, + hipJitOption* jitOptions, void** jitOptionsValues, unsigned int numJitOptions, - hipLibraryOption** libraryOptions, + hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions); typedef hipError_t (*t_hipLibraryUnload)(hipLibrary_t library); diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index 8b93feedf0..d9b3941439 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -100,7 +100,7 @@ enum hip_api_id_t { HIP_API_ID_hipGetDeviceFlags = 80, HIP_API_ID_hipGetDevicePropertiesR0000 = 81, HIP_API_ID_RESERVED_82 = 82, - HIP_API_ID_hipGetErrorString = 83, + HIP_API_ID_RESERVED_83 = 83, HIP_API_ID_hipGetLastError = 84, HIP_API_ID_hipGetMipmappedArrayLevel = 85, HIP_API_ID_hipGetSymbolAddress = 86, @@ -461,7 +461,8 @@ enum hip_api_id_t { HIP_API_ID_hipLibraryUnload = 441, HIP_API_ID_hipLibraryGetKernel = 442, HIP_API_ID_hipLibraryGetKernelCount = 443, - HIP_API_ID_LAST = 443, + HIP_API_ID_hipMemGetHandleForAddressRange = 444, + HIP_API_ID_LAST = 444, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -474,13 +475,11 @@ enum hip_api_id_t { HIP_API_ID_hipDestroyTextureObject = HIP_API_ID_NONE, HIP_API_ID_hipDeviceGetCount = HIP_API_ID_NONE, HIP_API_ID_hipDeviceGetTexture1DLinearMaxWidth = HIP_API_ID_NONE, - HIP_API_ID_hipGetDriverEntryPoint_spt = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureAlignmentOffset = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureObjectResourceDesc = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureObjectResourceViewDesc = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureObjectTextureDesc = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureReference = HIP_API_ID_NONE, - HIP_API_ID_hipMemGetHandleForAddressRange = HIP_API_ID_NONE, HIP_API_ID_hipTexObjectCreate = HIP_API_ID_NONE, HIP_API_ID_hipTexObjectDestroy = HIP_API_ID_NONE, HIP_API_ID_hipTexObjectGetResourceDesc = HIP_API_ID_NONE, @@ -613,7 +612,6 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGetDevicePropertiesR0000: return "hipGetDevicePropertiesR0000"; case HIP_API_ID_hipGetDevicePropertiesR0600: return "hipGetDevicePropertiesR0600"; case HIP_API_ID_hipGetDriverEntryPoint: return "hipGetDriverEntryPoint"; - case HIP_API_ID_hipGetErrorString: return "hipGetErrorString"; case HIP_API_ID_hipGetFuncBySymbol: return "hipGetFuncBySymbol"; case HIP_API_ID_hipGetLastError: return "hipGetLastError"; case HIP_API_ID_hipGetMipmappedArrayLevel: return "hipGetMipmappedArrayLevel"; @@ -734,6 +732,11 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipLaunchHostFunc: return "hipLaunchHostFunc"; case HIP_API_ID_hipLaunchKernel: return "hipLaunchKernel"; case HIP_API_ID_hipLaunchKernelExC: return "hipLaunchKernelExC"; + case HIP_API_ID_hipLibraryGetKernel: return "hipLibraryGetKernel"; + case HIP_API_ID_hipLibraryGetKernelCount: return "hipLibraryGetKernelCount"; + case HIP_API_ID_hipLibraryLoadData: return "hipLibraryLoadData"; + case HIP_API_ID_hipLibraryLoadFromFile: return "hipLibraryLoadFromFile"; + case HIP_API_ID_hipLibraryUnload: return "hipLibraryUnload"; case HIP_API_ID_hipLinkAddData: return "hipLinkAddData"; case HIP_API_ID_hipLinkAddFile: return "hipLinkAddFile"; case HIP_API_ID_hipLinkComplete: return "hipLinkComplete"; @@ -761,6 +764,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipMemGetAddressRange: return "hipMemGetAddressRange"; case HIP_API_ID_hipMemGetAllocationGranularity: return "hipMemGetAllocationGranularity"; case HIP_API_ID_hipMemGetAllocationPropertiesFromHandle: return "hipMemGetAllocationPropertiesFromHandle"; + case HIP_API_ID_hipMemGetHandleForAddressRange: return "hipMemGetHandleForAddressRange"; case HIP_API_ID_hipMemGetInfo: return "hipMemGetInfo"; case HIP_API_ID_hipMemImportFromShareableHandle: return "hipMemImportFromShareableHandle"; case HIP_API_ID_hipMemMap: return "hipMemMap"; @@ -846,15 +850,16 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipMipmappedArrayDestroy: return "hipMipmappedArrayDestroy"; case HIP_API_ID_hipMipmappedArrayGetLevel: return "hipMipmappedArrayGetLevel"; case HIP_API_ID_hipModuleGetFunction: return "hipModuleGetFunction"; + case HIP_API_ID_hipModuleGetFunctionCount: return "hipModuleGetFunctionCount"; case HIP_API_ID_hipModuleGetGlobal: return "hipModuleGetGlobal"; case HIP_API_ID_hipModuleGetTexRef: return "hipModuleGetTexRef"; case HIP_API_ID_hipModuleLaunchCooperativeKernel: return "hipModuleLaunchCooperativeKernel"; case HIP_API_ID_hipModuleLaunchCooperativeKernelMultiDevice: return "hipModuleLaunchCooperativeKernelMultiDevice"; case HIP_API_ID_hipModuleLaunchKernel: return "hipModuleLaunchKernel"; - case HIP_API_ID_hipModuleLoadFatBinary: return "hipModuleLoadFatBinary"; case HIP_API_ID_hipModuleLoad: return "hipModuleLoad"; case HIP_API_ID_hipModuleLoadData: return "hipModuleLoadData"; case HIP_API_ID_hipModuleLoadDataEx: return "hipModuleLoadDataEx"; + case HIP_API_ID_hipModuleLoadFatBinary: return "hipModuleLoadFatBinary"; case HIP_API_ID_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor: return "hipModuleOccupancyMaxActiveBlocksPerMultiprocessor"; case HIP_API_ID_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags: return "hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags"; case HIP_API_ID_hipModuleOccupancyMaxPotentialBlockSize: return "hipModuleOccupancyMaxPotentialBlockSize"; @@ -926,12 +931,6 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipUserObjectRelease: return "hipUserObjectRelease"; case HIP_API_ID_hipUserObjectRetain: return "hipUserObjectRetain"; case HIP_API_ID_hipWaitExternalSemaphoresAsync: return "hipWaitExternalSemaphoresAsync"; - case HIP_API_ID_hipModuleGetFunctionCount: return "hipModuleGetFunctionCount"; - case HIP_API_ID_hipLibraryLoadData: return "hipLibraryLoadData"; - case HIP_API_ID_hipLibraryLoadFromFile: return "hipLibraryLoadFromFile"; - case HIP_API_ID_hipLibraryUnload: return "hipLibraryUnload"; - case HIP_API_ID_hipLibraryGetKernel: return "hipLibraryGetKernel"; - case HIP_API_ID_hipLibraryGetKernelCount: return "hipLibraryGetKernelCount"; }; return "unknown"; }; @@ -1051,7 +1050,6 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipGetDevicePropertiesR0000", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0000; if (strcmp("hipGetDevicePropertiesR0600", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0600; if (strcmp("hipGetDriverEntryPoint", name) == 0) return HIP_API_ID_hipGetDriverEntryPoint; - if (strcmp("hipGetErrorString", name) == 0) return HIP_API_ID_hipGetErrorString; if (strcmp("hipGetFuncBySymbol", name) == 0) return HIP_API_ID_hipGetFuncBySymbol; if (strcmp("hipGetLastError", name) == 0) return HIP_API_ID_hipGetLastError; if (strcmp("hipGetMipmappedArrayLevel", name) == 0) return HIP_API_ID_hipGetMipmappedArrayLevel; @@ -1172,6 +1170,11 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipLaunchHostFunc", name) == 0) return HIP_API_ID_hipLaunchHostFunc; if (strcmp("hipLaunchKernel", name) == 0) return HIP_API_ID_hipLaunchKernel; if (strcmp("hipLaunchKernelExC", name) == 0) return HIP_API_ID_hipLaunchKernelExC; + if (strcmp("hipLibraryGetKernel", name) == 0) return HIP_API_ID_hipLibraryGetKernel; + if (strcmp("hipLibraryGetKernelCount", name) == 0) return HIP_API_ID_hipLibraryGetKernelCount; + if (strcmp("hipLibraryLoadData", name) == 0) return HIP_API_ID_hipLibraryLoadData; + if (strcmp("hipLibraryLoadFromFile", name) == 0) return HIP_API_ID_hipLibraryLoadFromFile; + if (strcmp("hipLibraryUnload", name) == 0) return HIP_API_ID_hipLibraryUnload; 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; @@ -1199,6 +1202,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipMemGetAddressRange", name) == 0) return HIP_API_ID_hipMemGetAddressRange; if (strcmp("hipMemGetAllocationGranularity", name) == 0) return HIP_API_ID_hipMemGetAllocationGranularity; if (strcmp("hipMemGetAllocationPropertiesFromHandle", name) == 0) return HIP_API_ID_hipMemGetAllocationPropertiesFromHandle; + if (strcmp("hipMemGetHandleForAddressRange", name) == 0) return HIP_API_ID_hipMemGetHandleForAddressRange; if (strcmp("hipMemGetInfo", name) == 0) return HIP_API_ID_hipMemGetInfo; if (strcmp("hipMemImportFromShareableHandle", name) == 0) return HIP_API_ID_hipMemImportFromShareableHandle; if (strcmp("hipMemMap", name) == 0) return HIP_API_ID_hipMemMap; @@ -1284,15 +1288,16 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipMipmappedArrayDestroy", name) == 0) return HIP_API_ID_hipMipmappedArrayDestroy; if (strcmp("hipMipmappedArrayGetLevel", name) == 0) return HIP_API_ID_hipMipmappedArrayGetLevel; if (strcmp("hipModuleGetFunction", name) == 0) return HIP_API_ID_hipModuleGetFunction; + if (strcmp("hipModuleGetFunctionCount", name) == 0) return HIP_API_ID_hipModuleGetFunctionCount; if (strcmp("hipModuleGetGlobal", name) == 0) return HIP_API_ID_hipModuleGetGlobal; if (strcmp("hipModuleGetTexRef", name) == 0) return HIP_API_ID_hipModuleGetTexRef; if (strcmp("hipModuleLaunchCooperativeKernel", name) == 0) return HIP_API_ID_hipModuleLaunchCooperativeKernel; if (strcmp("hipModuleLaunchCooperativeKernelMultiDevice", name) == 0) return HIP_API_ID_hipModuleLaunchCooperativeKernelMultiDevice; if (strcmp("hipModuleLaunchKernel", name) == 0) return HIP_API_ID_hipModuleLaunchKernel; - if (strcmp("hipModuleLoadFatBinary", name) == 0) return HIP_API_ID_hipModuleLoadFatBinary; if (strcmp("hipModuleLoad", name) == 0) return HIP_API_ID_hipModuleLoad; if (strcmp("hipModuleLoadData", name) == 0) return HIP_API_ID_hipModuleLoadData; if (strcmp("hipModuleLoadDataEx", name) == 0) return HIP_API_ID_hipModuleLoadDataEx; + if (strcmp("hipModuleLoadFatBinary", name) == 0) return HIP_API_ID_hipModuleLoadFatBinary; if (strcmp("hipModuleOccupancyMaxActiveBlocksPerMultiprocessor", name) == 0) return HIP_API_ID_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor; if (strcmp("hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags", name) == 0) return HIP_API_ID_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags; if (strcmp("hipModuleOccupancyMaxPotentialBlockSize", name) == 0) return HIP_API_ID_hipModuleOccupancyMaxPotentialBlockSize; @@ -1364,12 +1369,6 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipUserObjectRelease", name) == 0) return HIP_API_ID_hipUserObjectRelease; if (strcmp("hipUserObjectRetain", name) == 0) return HIP_API_ID_hipUserObjectRetain; if (strcmp("hipWaitExternalSemaphoresAsync", name) == 0) return HIP_API_ID_hipWaitExternalSemaphoresAsync; - if (strcmp("hipModuleGetFunctionCount", name) == 0) return HIP_API_ID_hipModuleGetFunctionCount; - if (strcmp("hipLibraryLoadData", name) == 0) return HIP_API_ID_hipLibraryLoadData; - if (strcmp("hipLibraryLoadFromFile", name) == 0) return HIP_API_ID_hipLibraryLoadFromFile; - if (strcmp("hipLibraryUnload", name) == 0) return HIP_API_ID_hipLibraryUnload; - if (strcmp("hipLibraryGetKernel", name) == 0) return HIP_API_ID_hipLibraryGetKernel; - if (strcmp("hipLibraryGetKernelCount", name) == 0) return HIP_API_ID_hipLibraryGetKernelCount; return HIP_API_ID_NONE; } @@ -2709,6 +2708,52 @@ typedef struct hip_api_data_s { void** args; void* args__val; } hipLaunchKernelExC; + struct { + hipKernel_t* pKernel; + hipKernel_t pKernel__val; + hipLibrary_t library; + const char* name; + char name__val; + } hipLibraryGetKernel; + struct { + unsigned int* count; + unsigned int count__val; + hipLibrary_t library; + } hipLibraryGetKernelCount; + struct { + hipLibrary_t* library; + hipLibrary_t library__val; + const void* code; + hipJitOption* jitOptions; + hipJitOption jitOptions__val; + void** jitOptionsValues; + void* jitOptionsValues__val; + unsigned int numJitOptions; + hipLibraryOption* libraryOptions; + hipLibraryOption libraryOptions__val; + void** libraryOptionValues; + void* libraryOptionValues__val; + unsigned int numLibraryOptions; + } hipLibraryLoadData; + struct { + hipLibrary_t* library; + hipLibrary_t library__val; + const char* fileName; + char fileName__val; + hipJitOption* jitOptions; + hipJitOption jitOptions__val; + void** jitOptionsValues; + void* jitOptionsValues__val; + unsigned int numJitOptions; + hipLibraryOption* libraryOptions; + hipLibraryOption libraryOptions__val; + void** libraryOptionValues; + void* libraryOptionValues__val; + unsigned int numLibraryOptions; + } hipLibraryLoadFromFile; + struct { + hipLibrary_t library; + } hipLibraryUnload; struct { hipLinkState_t state; hipJitInputType type; @@ -2898,6 +2943,13 @@ typedef struct hip_api_data_s { hipMemAllocationProp prop__val; hipMemGenericAllocationHandle_t handle; } hipMemGetAllocationPropertiesFromHandle; + struct { + void* handle; + hipDeviceptr_t dptr; + size_t size; + hipMemRangeHandleType handleType; + unsigned long long flags; + } hipMemGetHandleForAddressRange; struct { size_t* free; size_t free__val; @@ -3516,11 +3568,6 @@ typedef struct hip_api_data_s { void** extra; void* extra__val; } hipModuleLaunchKernel; - struct { - hipModule_t* module; - hipModule_t module__val; - const void* fatbin; - } hipModuleLoadFatBinary; struct { hipModule_t* module; hipModule_t module__val; @@ -3542,6 +3589,11 @@ typedef struct hip_api_data_s { void** optionsValues; void* optionsValues__val; } hipModuleLoadDataEx; + struct { + hipModule_t* module; + hipModule_t module__val; + const void* fatbin; + } hipModuleLoadFatBinary; struct { int* numBlocks; int numBlocks__val; @@ -3705,7 +3757,7 @@ typedef struct hip_api_data_s { struct { hipStream_t stream; hipLaunchAttributeID attr; - const hipLaunchAttributeValue* value_out; + hipLaunchAttributeValue* value_out; hipLaunchAttributeValue value_out__val; } hipStreamGetAttribute; struct { @@ -3951,44 +4003,6 @@ typedef struct hip_api_data_s { unsigned int numExtSems; hipStream_t stream; } hipWaitExternalSemaphoresAsync; - struct { - hipLibrary_t* library; - hipLibrary_t library__val; - const void* image; - hipJitOption** jitOptions; - void** jitOptionsValues; - unsigned int numJitOptions; - hipLibraryOption** libraryOptions; - void** libraryOptionValues; - unsigned int numLibraryOptions; - } hipLibraryLoadData; - struct { - hipLibrary_t* library; - hipLibrary_t library__val; - const char* fname; - char fname__val; - hipJitOption** jitOptions; - void** jitOptionsValues; - unsigned int numJitOptions; - hipLibraryOption** libraryOptions; - void** libraryOptionValues; - unsigned int numLibraryOptions; - } hipLibraryLoadFromFile; - struct { - hipLibrary_t library; - } hipLibraryUnload; - struct { - hipKernel_t* kernel; - hipKernel_t kernel__val; - hipLibrary_t library; - const char* kname; - char kname__val; - } hipLibraryGetKernel; - struct { - unsigned int *count; - unsigned int count__val; - hipLibrary_t library; - } hipLibraryGetKernelCount; } args; uint64_t *phase_data; } hip_api_data_t; @@ -4581,9 +4595,6 @@ typedef struct hip_api_data_s { cb_data.args.hipGetDriverEntryPoint.flags = (unsigned long long)flags; \ cb_data.args.hipGetDriverEntryPoint.driverStatus = (hipDriverEntryPointQueryResult*)status; \ }; -// hipGetErrorString[] -#define INIT_hipGetErrorString_CB_ARGS_DATA(cb_data) { \ -}; // hipGetFuncBySymbol[('hipFunction_t*', 'functionPtr'), ('const void*', 'symbolPtr')] #define INIT_hipGetFuncBySymbol_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipGetFuncBySymbol.functionPtr = (hipFunction_t*)functionPtr; \ @@ -5329,6 +5340,43 @@ typedef struct hip_api_data_s { cb_data.args.hipLaunchKernelExC.fPtr = (const void*)fPtr; \ cb_data.args.hipLaunchKernelExC.args = (void**)args; \ }; +// hipLibraryGetKernel[('hipKernel_t*', 'pKernel'), ('hipLibrary_t', 'library'), ('const char*', 'name')] +#define INIT_hipLibraryGetKernel_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLibraryGetKernel.pKernel = (hipKernel_t*)kernel; \ + cb_data.args.hipLibraryGetKernel.library = (hipLibrary_t)library; \ + cb_data.args.hipLibraryGetKernel.name = (kname) ? strdup(kname) : NULL; \ +}; +// hipLibraryGetKernelCount[('unsigned int*', 'count'), ('hipLibrary_t', 'library')] +#define INIT_hipLibraryGetKernelCount_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLibraryGetKernelCount.count = (unsigned int*)count; \ + cb_data.args.hipLibraryGetKernelCount.library = (hipLibrary_t)library; \ +}; +// hipLibraryLoadData[('hipLibrary_t*', 'library'), ('const void*', 'code'), ('hipJitOption*', 'jitOptions'), ('void**', 'jitOptionsValues'), ('unsigned int', 'numJitOptions'), ('hipLibraryOption*', 'libraryOptions'), ('void**', 'libraryOptionValues'), ('unsigned int', 'numLibraryOptions')] +#define INIT_hipLibraryLoadData_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLibraryLoadData.library = (hipLibrary_t*)library; \ + cb_data.args.hipLibraryLoadData.code = (const void*)image; \ + cb_data.args.hipLibraryLoadData.jitOptions = (hipJitOption*)jitOptions; \ + cb_data.args.hipLibraryLoadData.jitOptionsValues = (void**)jitOptionsValues; \ + cb_data.args.hipLibraryLoadData.numJitOptions = (unsigned int)numJitOptions; \ + cb_data.args.hipLibraryLoadData.libraryOptions = (hipLibraryOption*)libraryOptions; \ + cb_data.args.hipLibraryLoadData.libraryOptionValues = (void**)libraryOptionValues; \ + cb_data.args.hipLibraryLoadData.numLibraryOptions = (unsigned int)numLibraryOptions; \ +}; +// hipLibraryLoadFromFile[('hipLibrary_t*', 'library'), ('const char*', 'fileName'), ('hipJitOption*', 'jitOptions'), ('void**', 'jitOptionsValues'), ('unsigned int', 'numJitOptions'), ('hipLibraryOption*', 'libraryOptions'), ('void**', 'libraryOptionValues'), ('unsigned int', 'numLibraryOptions')] +#define INIT_hipLibraryLoadFromFile_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLibraryLoadFromFile.library = (hipLibrary_t*)library; \ + cb_data.args.hipLibraryLoadFromFile.fileName = (fname) ? strdup(fname) : NULL; \ + cb_data.args.hipLibraryLoadFromFile.jitOptions = (hipJitOption*)jitOptions; \ + cb_data.args.hipLibraryLoadFromFile.jitOptionsValues = (void**)jitOptionsValues; \ + cb_data.args.hipLibraryLoadFromFile.numJitOptions = (unsigned int)numJitOptions; \ + cb_data.args.hipLibraryLoadFromFile.libraryOptions = (hipLibraryOption*)libraryOptions; \ + cb_data.args.hipLibraryLoadFromFile.libraryOptionValues = (void**)libraryOptionValues; \ + cb_data.args.hipLibraryLoadFromFile.numLibraryOptions = (unsigned int)numLibraryOptions; \ +}; +// hipLibraryUnload[('hipLibrary_t', 'library')] +#define INIT_hipLibraryUnload_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLibraryUnload.library = (hipLibrary_t)library; \ +}; // 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; \ @@ -5507,6 +5555,14 @@ typedef struct hip_api_data_s { cb_data.args.hipMemGetAllocationPropertiesFromHandle.prop = (hipMemAllocationProp*)prop; \ cb_data.args.hipMemGetAllocationPropertiesFromHandle.handle = (hipMemGenericAllocationHandle_t)handle; \ }; +// hipMemGetHandleForAddressRange[('void*', 'handle'), ('hipDeviceptr_t', 'dptr'), ('size_t', 'size'), ('hipMemRangeHandleType', 'handleType'), ('unsigned long long', 'flags')] +#define INIT_hipMemGetHandleForAddressRange_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemGetHandleForAddressRange.handle = (void*)handle; \ + cb_data.args.hipMemGetHandleForAddressRange.dptr = (hipDeviceptr_t)dptr; \ + cb_data.args.hipMemGetHandleForAddressRange.size = (size_t)size; \ + cb_data.args.hipMemGetHandleForAddressRange.handleType = (hipMemRangeHandleType)handleType; \ + cb_data.args.hipMemGetHandleForAddressRange.flags = (unsigned long long)flags; \ +}; // hipMemGetInfo[('size_t*', 'free'), ('size_t*', 'total')] #define INIT_hipMemGetInfo_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMemGetInfo.free = (size_t*)free; \ @@ -6117,6 +6173,11 @@ typedef struct hip_api_data_s { cb_data.args.hipModuleGetFunction.module = (hipModule_t)hmod; \ cb_data.args.hipModuleGetFunction.kname = (name) ? strdup(name) : NULL; \ }; +// hipModuleGetFunctionCount[('unsigned int*', 'count'), ('hipModule_t', 'mod')] +#define INIT_hipModuleGetFunctionCount_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipModuleGetFunctionCount.count = (unsigned int*)count; \ + cb_data.args.hipModuleGetFunctionCount.mod = (hipModule_t)mod; \ +}; // hipModuleGetGlobal[('hipDeviceptr_t*', 'dptr'), ('size_t*', 'bytes'), ('hipModule_t', 'hmod'), ('const char*', 'name')] #define INIT_hipModuleGetGlobal_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipModuleGetGlobal.dptr = (hipDeviceptr_t*)dptr; \ @@ -6163,11 +6224,6 @@ typedef struct hip_api_data_s { cb_data.args.hipModuleLaunchKernel.kernelParams = (void**)kernelParams; \ cb_data.args.hipModuleLaunchKernel.extra = (void**)extra; \ }; -// hipModuleLoadFatBinary[('hipModule_t*', 'module'), ('const void*', 'fatbin')] -#define INIT_hipModuleLoadFatBinary_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipModuleLoadFatBinary.module = (hipModule_t*)module; \ - cb_data.args.hipModuleLoadFatBinary.fatbin = (const void*)fatbin; \ -}; // hipModuleLoad[('hipModule_t*', 'module'), ('const char*', 'fname')] #define INIT_hipModuleLoad_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipModuleLoad.module = (hipModule_t*)module; \ @@ -6186,6 +6242,11 @@ typedef struct hip_api_data_s { cb_data.args.hipModuleLoadDataEx.options = (hipJitOption*)options; \ cb_data.args.hipModuleLoadDataEx.optionsValues = (void**)optionsValues; \ }; +// hipModuleLoadFatBinary[('hipModule_t*', 'module'), ('const void*', 'fatbin')] +#define INIT_hipModuleLoadFatBinary_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipModuleLoadFatBinary.module = (hipModule_t*)module; \ + cb_data.args.hipModuleLoadFatBinary.fatbin = (const void*)fatbin; \ +}; // hipModuleOccupancyMaxActiveBlocksPerMultiprocessor[('int*', 'numBlocks'), ('hipFunction_t', 'f'), ('int', 'blockSize'), ('size_t', 'dynSharedMemPerBlk')] #define INIT_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.numBlocks = (int*)numBlocks; \ @@ -6360,7 +6421,7 @@ typedef struct hip_api_data_s { cb_data.args.hipStreamEndCapture.stream = (hipStream_t)stream; \ cb_data.args.hipStreamEndCapture.pGraph = (hipGraph_t*)pGraph; \ }; -// hipStreamGetAttribute[('hipStream_t', 'stream'), ('hipLaunchAttributeID', 'attr'), ('const hipLaunchAttributeValue*', 'value_out')] +// hipStreamGetAttribute[('hipStream_t', 'stream'), ('hipLaunchAttributeID', 'attr'), ('hipLaunchAttributeValue*', 'value_out')] #define INIT_hipStreamGetAttribute_CB_ARGS_DATA(cb_data) { \ }; // hipStreamGetCaptureInfo[('hipStream_t', 'stream'), ('hipStreamCaptureStatus*', 'pCaptureStatus'), ('unsigned long long*', 'pId')] @@ -6591,12 +6652,6 @@ typedef struct hip_api_data_s { cb_data.args.hipWaitExternalSemaphoresAsync.numExtSems = (unsigned int)numExtSems; \ cb_data.args.hipWaitExternalSemaphoresAsync.stream = (hipStream_t)stream; \ }; -// hipModuleGetFunctionCount[('unsigned int*', 'count'), ('hipModule_t', 'mod')] -#define INIT_hipModuleGetFunctionCount_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipModuleGetFunctionCount.count = (unsigned int*)count; \ - cb_data.args.hipModuleGetFunctionCount.mod = (hipModule_t)mod; \ -}; - #define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data) // Macros for non-public API primitives @@ -6616,8 +6671,6 @@ typedef struct hip_api_data_s { #define INIT_hipDeviceGetCount_CB_ARGS_DATA(cb_data) {}; // hipDeviceGetTexture1DLinearMaxWidth() #define INIT_hipDeviceGetTexture1DLinearMaxWidth_CB_ARGS_DATA(cb_data) {}; -// hipGetDriverEntryPoint_spt() -#define INIT_hipGetDriverEntryPoint_spt_CB_ARGS_DATA(cb_data) {}; // hipGetTextureAlignmentOffset() #define INIT_hipGetTextureAlignmentOffset_CB_ARGS_DATA(cb_data) {}; // hipGetTextureObjectResourceDesc() @@ -6628,8 +6681,6 @@ typedef struct hip_api_data_s { #define INIT_hipGetTextureObjectTextureDesc_CB_ARGS_DATA(cb_data) {}; // hipGetTextureReference() #define INIT_hipGetTextureReference_CB_ARGS_DATA(cb_data) {}; -// hipMemGetHandleForAddressRange() -#define INIT_hipMemGetHandleForAddressRange_CB_ARGS_DATA(cb_data) {}; // hipTexObjectCreate() #define INIT_hipTexObjectCreate_CB_ARGS_DATA(cb_data) {}; // hipTexObjectDestroy() @@ -6654,46 +6705,6 @@ typedef struct hip_api_data_s { #define INIT_hipTexRefSetMipmapFilterMode_CB_ARGS_DATA(cb_data) {}; // hipUnbindTexture() #define INIT_hipUnbindTexture_CB_ARGS_DATA(cb_data) {}; -// hipLibraryLoadData() -#define INIT_hipLibraryLoadData_CB_ARGS_DATA(cb_data) \ - { \ - cb_data.args.hipLibraryLoadData.library = (hipLibrary_t*)library; \ - cb_data.args.hipLibraryLoadData.image = (const void*)image; \ - cb_data.args.hipLibraryLoadData.jitOptions = (hipJitOption**)jitOptions; \ - cb_data.args.hipLibraryLoadData.jitOptionsValues = (void**)jitOptionsValues; \ - cb_data.args.hipLibraryLoadData.numJitOptions = (unsigned int)numJitOptions; \ - cb_data.args.hipLibraryLoadData.libraryOptions = (hipLibraryOption**)libraryOptions; \ - cb_data.args.hipLibraryLoadData.libraryOptionValues = (void**)libraryOptionValues; \ - cb_data.args.hipLibraryLoadData.numLibraryOptions = (unsigned int)numLibraryOptions; \ - }; -// hipLibraryLoadFromFile() -#define INIT_hipLibraryLoadFromFile_CB_ARGS_DATA(cb_data) \ - { \ - cb_data.args.hipLibraryLoadFromFile.library = (hipLibrary_t*)library; \ - cb_data.args.hipLibraryLoadFromFile.fname = (const char*)fname; \ - cb_data.args.hipLibraryLoadFromFile.jitOptions = (hipJitOption**)jitOptions; \ - cb_data.args.hipLibraryLoadFromFile.jitOptionsValues = (void**)jitOptionsValues; \ - cb_data.args.hipLibraryLoadFromFile.numJitOptions = (unsigned int)numJitOptions; \ - cb_data.args.hipLibraryLoadFromFile.libraryOptions = (hipLibraryOption**)libraryOptions; \ - cb_data.args.hipLibraryLoadFromFile.libraryOptionValues = (void**)libraryOptionValues; \ - cb_data.args.hipLibraryLoadFromFile.numLibraryOptions = (unsigned int)numLibraryOptions; \ - }; -// hipLibraryUnload() -#define INIT_hipLibraryUnload_CB_ARGS_DATA(cb_data) \ - { cb_data.args.hipLibraryUnload.library = (hipLibrary_t)library; }; -// hipLibraryGetKernel() -#define INIT_hipLibraryGetKernel_CB_ARGS_DATA(cb_data) \ - { \ - cb_data.args.hipLibraryGetKernel.kernel = (hipKernel_t *)kernel; \ - cb_data.args.hipLibraryGetKernel.library = (hipLibrary_t)library; \ - cb_data.args.hipLibraryGetKernel.kname = (const char *)kname; \ - }; -// hipLibraryGetKernelCount() -#define INIT_hipLibraryGetKernelCount_CB_ARGS_DATA(cb_data) \ - { \ - cb_data.args.hipLibraryGetKernelCount.count = (unsigned int *)count; \ - cb_data.args.hipLibraryGetKernelCount.library = (hipLibrary_t)library; \ - }; #define INIT_NONE_CB_ARGS_DATA(cb_data) {}; @@ -7137,9 +7148,6 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipGetDriverEntryPoint.funcPtr) data->args.hipGetDriverEntryPoint.funcPtr__val = *(data->args.hipGetDriverEntryPoint.funcPtr); if (data->args.hipGetDriverEntryPoint.driverStatus) data->args.hipGetDriverEntryPoint.driverStatus__val = *(data->args.hipGetDriverEntryPoint.driverStatus); break; -// hipGetErrorString[] - case HIP_API_ID_hipGetErrorString: - break; // hipGetFuncBySymbol[('hipFunction_t*', 'functionPtr'), ('const void*', 'symbolPtr')] case HIP_API_ID_hipGetFuncBySymbol: if (data->args.hipGetFuncBySymbol.functionPtr) data->args.hipGetFuncBySymbol.functionPtr__val = *(data->args.hipGetFuncBySymbol.functionPtr); @@ -7635,6 +7643,35 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipLaunchKernelExC.config) data->args.hipLaunchKernelExC.config__val = *(data->args.hipLaunchKernelExC.config); if (data->args.hipLaunchKernelExC.args) data->args.hipLaunchKernelExC.args__val = *(data->args.hipLaunchKernelExC.args); break; +// hipLibraryGetKernel[('hipKernel_t*', 'pKernel'), ('hipLibrary_t', 'library'), ('const char*', 'name')] + case HIP_API_ID_hipLibraryGetKernel: + if (data->args.hipLibraryGetKernel.pKernel) data->args.hipLibraryGetKernel.pKernel__val = *(data->args.hipLibraryGetKernel.pKernel); + if (data->args.hipLibraryGetKernel.name) data->args.hipLibraryGetKernel.name__val = *(data->args.hipLibraryGetKernel.name); + break; +// hipLibraryGetKernelCount[('unsigned int*', 'count'), ('hipLibrary_t', 'library')] + case HIP_API_ID_hipLibraryGetKernelCount: + if (data->args.hipLibraryGetKernelCount.count) data->args.hipLibraryGetKernelCount.count__val = *(data->args.hipLibraryGetKernelCount.count); + break; +// hipLibraryLoadData[('hipLibrary_t*', 'library'), ('const void*', 'code'), ('hipJitOption*', 'jitOptions'), ('void**', 'jitOptionsValues'), ('unsigned int', 'numJitOptions'), ('hipLibraryOption*', 'libraryOptions'), ('void**', 'libraryOptionValues'), ('unsigned int', 'numLibraryOptions')] + case HIP_API_ID_hipLibraryLoadData: + if (data->args.hipLibraryLoadData.library) data->args.hipLibraryLoadData.library__val = *(data->args.hipLibraryLoadData.library); + if (data->args.hipLibraryLoadData.jitOptions) data->args.hipLibraryLoadData.jitOptions__val = *(data->args.hipLibraryLoadData.jitOptions); + if (data->args.hipLibraryLoadData.jitOptionsValues) data->args.hipLibraryLoadData.jitOptionsValues__val = *(data->args.hipLibraryLoadData.jitOptionsValues); + if (data->args.hipLibraryLoadData.libraryOptions) data->args.hipLibraryLoadData.libraryOptions__val = *(data->args.hipLibraryLoadData.libraryOptions); + if (data->args.hipLibraryLoadData.libraryOptionValues) data->args.hipLibraryLoadData.libraryOptionValues__val = *(data->args.hipLibraryLoadData.libraryOptionValues); + break; +// hipLibraryLoadFromFile[('hipLibrary_t*', 'library'), ('const char*', 'fileName'), ('hipJitOption*', 'jitOptions'), ('void**', 'jitOptionsValues'), ('unsigned int', 'numJitOptions'), ('hipLibraryOption*', 'libraryOptions'), ('void**', 'libraryOptionValues'), ('unsigned int', 'numLibraryOptions')] + case HIP_API_ID_hipLibraryLoadFromFile: + if (data->args.hipLibraryLoadFromFile.library) data->args.hipLibraryLoadFromFile.library__val = *(data->args.hipLibraryLoadFromFile.library); + if (data->args.hipLibraryLoadFromFile.fileName) data->args.hipLibraryLoadFromFile.fileName__val = *(data->args.hipLibraryLoadFromFile.fileName); + if (data->args.hipLibraryLoadFromFile.jitOptions) data->args.hipLibraryLoadFromFile.jitOptions__val = *(data->args.hipLibraryLoadFromFile.jitOptions); + if (data->args.hipLibraryLoadFromFile.jitOptionsValues) data->args.hipLibraryLoadFromFile.jitOptionsValues__val = *(data->args.hipLibraryLoadFromFile.jitOptionsValues); + if (data->args.hipLibraryLoadFromFile.libraryOptions) data->args.hipLibraryLoadFromFile.libraryOptions__val = *(data->args.hipLibraryLoadFromFile.libraryOptions); + if (data->args.hipLibraryLoadFromFile.libraryOptionValues) data->args.hipLibraryLoadFromFile.libraryOptionValues__val = *(data->args.hipLibraryLoadFromFile.libraryOptionValues); + break; +// hipLibraryUnload[('hipLibrary_t', 'library')] + case HIP_API_ID_hipLibraryUnload: + 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); @@ -7754,6 +7791,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipMemGetAllocationPropertiesFromHandle: if (data->args.hipMemGetAllocationPropertiesFromHandle.prop) data->args.hipMemGetAllocationPropertiesFromHandle.prop__val = *(data->args.hipMemGetAllocationPropertiesFromHandle.prop); break; +// hipMemGetHandleForAddressRange[('void*', 'handle'), ('hipDeviceptr_t', 'dptr'), ('size_t', 'size'), ('hipMemRangeHandleType', 'handleType'), ('unsigned long long', 'flags')] + case HIP_API_ID_hipMemGetHandleForAddressRange: + break; // hipMemGetInfo[('size_t*', 'free'), ('size_t*', 'total')] case HIP_API_ID_hipMemGetInfo: if (data->args.hipMemGetInfo.free) data->args.hipMemGetInfo.free__val = *(data->args.hipMemGetInfo.free); @@ -8047,6 +8087,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipModuleGetFunction.function) data->args.hipModuleGetFunction.function__val = *(data->args.hipModuleGetFunction.function); if (data->args.hipModuleGetFunction.kname) data->args.hipModuleGetFunction.kname__val = *(data->args.hipModuleGetFunction.kname); break; +// hipModuleGetFunctionCount[('unsigned int*', 'count'), ('hipModule_t', 'mod')] + case HIP_API_ID_hipModuleGetFunctionCount: + if (data->args.hipModuleGetFunctionCount.count) data->args.hipModuleGetFunctionCount.count__val = *(data->args.hipModuleGetFunctionCount.count); + break; // hipModuleGetGlobal[('hipDeviceptr_t*', 'dptr'), ('size_t*', 'bytes'), ('hipModule_t', 'hmod'), ('const char*', 'name')] case HIP_API_ID_hipModuleGetGlobal: if (data->args.hipModuleGetGlobal.dptr) data->args.hipModuleGetGlobal.dptr__val = *(data->args.hipModuleGetGlobal.dptr); @@ -8071,10 +8115,6 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipModuleLaunchKernel.kernelParams) data->args.hipModuleLaunchKernel.kernelParams__val = *(data->args.hipModuleLaunchKernel.kernelParams); if (data->args.hipModuleLaunchKernel.extra) data->args.hipModuleLaunchKernel.extra__val = *(data->args.hipModuleLaunchKernel.extra); break; -// hipModuleLoadFatBinary[('hipModule_t*', 'module'), ('const void*', 'fatbin')] - case HIP_API_ID_hipModuleLoadFatBinary: - if (data->args.hipModuleLoadFatBinary.module) data->args.hipModuleLoadFatBinary.module__val = *(data->args.hipModuleLoadFatBinary.module); - break; // hipModuleLoad[('hipModule_t*', 'module'), ('const char*', 'fname')] case HIP_API_ID_hipModuleLoad: if (data->args.hipModuleLoad.module) data->args.hipModuleLoad.module__val = *(data->args.hipModuleLoad.module); @@ -8090,6 +8130,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipModuleLoadDataEx.options) data->args.hipModuleLoadDataEx.options__val = *(data->args.hipModuleLoadDataEx.options); if (data->args.hipModuleLoadDataEx.optionsValues) data->args.hipModuleLoadDataEx.optionsValues__val = *(data->args.hipModuleLoadDataEx.optionsValues); break; +// hipModuleLoadFatBinary[('hipModule_t*', 'module'), ('const void*', 'fatbin')] + case HIP_API_ID_hipModuleLoadFatBinary: + if (data->args.hipModuleLoadFatBinary.module) data->args.hipModuleLoadFatBinary.module__val = *(data->args.hipModuleLoadFatBinary.module); + break; // hipModuleOccupancyMaxActiveBlocksPerMultiprocessor[('int*', 'numBlocks'), ('hipFunction_t', 'f'), ('int', 'blockSize'), ('size_t', 'dynSharedMemPerBlk')] case HIP_API_ID_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor: if (data->args.hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.numBlocks) data->args.hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.numBlocks__val = *(data->args.hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.numBlocks); @@ -8202,7 +8246,7 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { case HIP_API_ID_hipStreamEndCapture: if (data->args.hipStreamEndCapture.pGraph) data->args.hipStreamEndCapture.pGraph__val = *(data->args.hipStreamEndCapture.pGraph); break; -// hipStreamGetAttribute[('hipStream_t', 'stream'), ('hipLaunchAttributeID', 'attr'), ('const hipLaunchAttributeValue*', 'value_out')] +// hipStreamGetAttribute[('hipStream_t', 'stream'), ('hipLaunchAttributeID', 'attr'), ('hipLaunchAttributeValue*', 'value_out')] case HIP_API_ID_hipStreamGetAttribute: if (data->args.hipStreamGetAttribute.value_out) data->args.hipStreamGetAttribute.value_out__val = *(data->args.hipStreamGetAttribute.value_out); break; @@ -8378,29 +8422,7 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipWaitExternalSemaphoresAsync.extSemArray) data->args.hipWaitExternalSemaphoresAsync.extSemArray__val = *(data->args.hipWaitExternalSemaphoresAsync.extSemArray); if (data->args.hipWaitExternalSemaphoresAsync.paramsArray) data->args.hipWaitExternalSemaphoresAsync.paramsArray__val = *(data->args.hipWaitExternalSemaphoresAsync.paramsArray); break; -// hipModuleGetFunctionCount[('unsigned int*', 'count'), ('hipModule_t', 'mod')] - case HIP_API_ID_hipModuleGetFunctionCount: - if (data->args.hipModuleGetFunctionCount.count) - data->args.hipModuleGetFunctionCount.count__val = - *(data->args.hipModuleGetFunctionCount.count); - break; - case HIP_API_ID_hipLibraryLoadData: - if (data->args.hipLibraryLoadData.library) - data->args.hipLibraryLoadData.library__val = *(data->args.hipLibraryLoadData.library); - break; - case HIP_API_ID_hipLibraryLoadFromFile: - if (data->args.hipLibraryLoadFromFile.library) - data->args.hipLibraryLoadFromFile.library__val = - *(data->args.hipLibraryLoadFromFile.library); - if (data->args.hipLibraryLoadFromFile.fname) - data->args.hipLibraryLoadFromFile.fname__val = *(data->args.hipLibraryLoadFromFile.fname); - break; - case HIP_API_ID_hipLibraryGetKernel: - if (data->args.hipLibraryGetKernel.kernel) - data->args.hipLibraryGetKernel.kernel__val = *(data->args.hipLibraryGetKernel.kernel); - break; - default: - break; + default: break; }; } @@ -9209,10 +9231,6 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << ", driverStatus="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDriverEntryPoint.driverStatus__val); } oss << ")"; break; - case HIP_API_ID_hipGetErrorString: - oss << "hipGetErrorString("; - oss << ")"; - break; case HIP_API_ID_hipGetFuncBySymbol: oss << "hipGetFuncBySymbol("; if (data->args.hipGetFuncBySymbol.functionPtr == NULL) oss << "functionPtr=NULL"; @@ -10219,6 +10237,62 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << ", args="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLaunchKernelExC.args__val); } oss << ")"; break; + case HIP_API_ID_hipLibraryGetKernel: + oss << "hipLibraryGetKernel("; + if (data->args.hipLibraryGetKernel.pKernel == NULL) oss << "pKernel=NULL"; + else { oss << "pKernel="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryGetKernel.pKernel__val); } + oss << ", library="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryGetKernel.library); + if (data->args.hipLibraryGetKernel.name == NULL) oss << ", name=NULL"; + else { oss << ", name="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryGetKernel.name__val); } + oss << ")"; + break; + case HIP_API_ID_hipLibraryGetKernelCount: + oss << "hipLibraryGetKernelCount("; + if (data->args.hipLibraryGetKernelCount.count == NULL) oss << "count=NULL"; + else { oss << "count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryGetKernelCount.count__val); } + oss << ", library="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryGetKernelCount.library); + oss << ")"; + break; + case HIP_API_ID_hipLibraryLoadData: + oss << "hipLibraryLoadData("; + if (data->args.hipLibraryLoadData.library == NULL) oss << "library=NULL"; + else { oss << "library="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.library__val); } + oss << ", code="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.code); + if (data->args.hipLibraryLoadData.jitOptions == NULL) oss << ", jitOptions=NULL"; + else { oss << ", jitOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.jitOptions__val); } + if (data->args.hipLibraryLoadData.jitOptionsValues == NULL) oss << ", jitOptionsValues=NULL"; + else { oss << ", jitOptionsValues="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.jitOptionsValues__val); } + oss << ", numJitOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.numJitOptions); + if (data->args.hipLibraryLoadData.libraryOptions == NULL) oss << ", libraryOptions=NULL"; + else { oss << ", libraryOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.libraryOptions__val); } + if (data->args.hipLibraryLoadData.libraryOptionValues == NULL) oss << ", libraryOptionValues=NULL"; + else { oss << ", libraryOptionValues="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.libraryOptionValues__val); } + oss << ", numLibraryOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.numLibraryOptions); + oss << ")"; + break; + case HIP_API_ID_hipLibraryLoadFromFile: + oss << "hipLibraryLoadFromFile("; + if (data->args.hipLibraryLoadFromFile.library == NULL) oss << "library=NULL"; + else { oss << "library="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.library__val); } + if (data->args.hipLibraryLoadFromFile.fileName == NULL) oss << ", fileName=NULL"; + else { oss << ", fileName="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.fileName__val); } + if (data->args.hipLibraryLoadFromFile.jitOptions == NULL) oss << ", jitOptions=NULL"; + else { oss << ", jitOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.jitOptions__val); } + if (data->args.hipLibraryLoadFromFile.jitOptionsValues == NULL) oss << ", jitOptionsValues=NULL"; + else { oss << ", jitOptionsValues="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.jitOptionsValues__val); } + oss << ", numJitOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.numJitOptions); + if (data->args.hipLibraryLoadFromFile.libraryOptions == NULL) oss << ", libraryOptions=NULL"; + else { oss << ", libraryOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.libraryOptions__val); } + if (data->args.hipLibraryLoadFromFile.libraryOptionValues == NULL) oss << ", libraryOptionValues=NULL"; + else { oss << ", libraryOptionValues="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.libraryOptionValues__val); } + oss << ", numLibraryOptions="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.numLibraryOptions); + oss << ")"; + break; + case HIP_API_ID_hipLibraryUnload: + oss << "hipLibraryUnload("; + oss << "library="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryUnload.library); + oss << ")"; + break; case HIP_API_ID_hipLinkAddData: oss << "hipLinkAddData("; oss << "state="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLinkAddData.state); @@ -10462,6 +10536,15 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", handle="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemGetAllocationPropertiesFromHandle.handle); oss << ")"; break; + case HIP_API_ID_hipMemGetHandleForAddressRange: + oss << "hipMemGetHandleForAddressRange("; + oss << "handle="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemGetHandleForAddressRange.handle); + oss << ", dptr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemGetHandleForAddressRange.dptr); + oss << ", size="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemGetHandleForAddressRange.size); + oss << ", handleType="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemGetHandleForAddressRange.handleType); + oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemGetHandleForAddressRange.flags); + oss << ")"; + break; case HIP_API_ID_hipMemGetInfo: oss << "hipMemGetInfo("; if (data->args.hipMemGetInfo.free == NULL) oss << "free=NULL"; @@ -11195,6 +11278,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << ", kname="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleGetFunction.kname__val); } oss << ")"; break; + case HIP_API_ID_hipModuleGetFunctionCount: + oss << "hipModuleGetFunctionCount("; + if (data->args.hipModuleGetFunctionCount.count == NULL) oss << "count=NULL"; + else { oss << "count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleGetFunctionCount.count__val); } + oss << ", mod="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleGetFunctionCount.mod); + oss << ")"; + break; case HIP_API_ID_hipModuleGetGlobal: oss << "hipModuleGetGlobal("; if (data->args.hipModuleGetGlobal.dptr == NULL) oss << "dptr=NULL"; @@ -11255,13 +11345,6 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << ", extra="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchKernel.extra__val); } oss << ")"; break; - case HIP_API_ID_hipModuleLoadFatBinary: - oss << "hipModuleLoadFatBinary("; - if (data->args.hipModuleLoadFatBinary.module == NULL) oss << "module=NULL"; - else { oss << "module="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLoadFatBinary.module__val); } - oss << ", fatbin="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLoadFatBinary.fatbin); - oss << ")"; - break; case HIP_API_ID_hipModuleLoad: oss << "hipModuleLoad("; if (data->args.hipModuleLoad.module == NULL) oss << "module=NULL"; @@ -11289,6 +11372,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da else { oss << ", optionsValues="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLoadDataEx.optionsValues__val); } oss << ")"; break; + case HIP_API_ID_hipModuleLoadFatBinary: + oss << "hipModuleLoadFatBinary("; + if (data->args.hipModuleLoadFatBinary.module == NULL) oss << "module=NULL"; + else { oss << "module="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLoadFatBinary.module__val); } + oss << ", fatbin="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLoadFatBinary.fatbin); + oss << ")"; + break; case HIP_API_ID_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor: oss << "hipModuleOccupancyMaxActiveBlocksPerMultiprocessor("; if (data->args.hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.numBlocks == NULL) oss << "numBlocks=NULL"; @@ -11846,116 +11936,7 @@ 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.hipWaitExternalSemaphoresAsync.stream); oss << ")"; break; - case HIP_API_ID_hipModuleGetFunctionCount: - oss << "hipModuleGetFunctionCount("; - if (data->args.hipModuleGetFunctionCount.count == NULL) oss << "count=NULL"; - else { oss << "count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleGetFunctionCount.count__val); } - oss << ", mod="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleGetFunctionCount.mod); - oss << ")"; - break; - case HIP_API_ID_hipLibraryLoadData: - oss << "hipLibraryLoadData("; - if (data->args.hipLibraryLoadData.library == NULL) - oss << "library=NULL"; - else { - oss << "library="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.library__val); - } - oss << ", image="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.image); - oss << ", jitOptions="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.jitOptions); - oss << ", jitOptionsValues="; - roctracer::hip_support::detail::operator<<(oss, - data->args.hipLibraryLoadData.jitOptionsValues); - oss << ", numJitOptions="; - roctracer::hip_support::detail::operator<<(oss, - data->args.hipLibraryLoadData.numJitOptions); - oss << ", libraryOptions="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadData.libraryOptions); - oss << ", libraryOptionsValues="; - roctracer::hip_support::detail::operator<<( - oss, data->args.hipLibraryLoadData.libraryOptionValues); - oss << ", numLibraryOptions="; - roctracer::hip_support::detail::operator<<(oss, - data->args.hipLibraryLoadData.numLibraryOptions); - oss << ")"; - break; - case HIP_API_ID_hipLibraryLoadFromFile: - oss << "hipLibraryLoadFromFile("; - if (data->args.hipLibraryLoadFromFile.library == NULL) - oss << "library=NULL"; - else { - oss << "library="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.library__val); - } - if (data->args.hipLibraryLoadFromFile.fname == NULL) - oss << "fname=NULL"; - else { - oss << "fname="; - roctracer::hip_support::detail::operator<<(oss, - data->args.hipLibraryLoadFromFile.fname__val); - } - oss << ", jitOptions="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.jitOptions); - oss << ")"; - oss << ", jitOptionsValues="; - roctracer::hip_support::detail::operator<<( - oss, data->args.hipLibraryLoadFromFile.jitOptionsValues); - oss << ")"; - oss << ", numJitOptions="; - roctracer::hip_support::detail::operator<<(oss, - data->args.hipLibraryLoadFromFile.numJitOptions); - oss << ")"; - oss << ", libraryOptions="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryLoadFromFile.libraryOptions); - oss << ")"; - oss << ", libraryOptionsValues="; - roctracer::hip_support::detail::operator<<( - oss, data->args.hipLibraryLoadFromFile.libraryOptionValues); - oss << ")"; - oss << ", numLibraryOptions="; - roctracer::hip_support::detail::operator<<(oss, - data->args.hipLibraryLoadFromFile.numLibraryOptions); - oss << ")"; - break; - case HIP_API_ID_hipLibraryUnload: - oss << "hipLibraryUnload("; - oss << ", library="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryUnload.library); - break; - case HIP_API_ID_hipLibraryGetKernel: - oss << "hipLibraryGetKernel("; - if (data->args.hipLibraryGetKernel.kernel == NULL) - oss << "kernel=NULL"; - else { - oss << "kernel="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryGetKernel.kernel__val); - } - oss << ", library="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryGetKernel.library); - if (data->args.hipLibraryGetKernel.kname == NULL) - oss << "kname=NULL"; - else { - oss << "kname="; - roctracer::hip_support::detail::operator<<(oss, data->args.hipLibraryGetKernel.kname__val); - } - break; - case HIP_API_ID_hipLibraryGetKernelCount: - oss << "hipLibraryGetKernelCount("; - if (data->args.hipLibraryGetKernelCount.count == NULL) - oss << "count=NULL"; - else { - oss << "count="; - roctracer::hip_support::detail::operator<<( - oss, data->args.hipLibraryGetKernelCount.count__val); - } - oss << ", library="; - roctracer::hip_support::detail::operator<<( - oss, data->args.hipLibraryGetKernelCount.library); - break; - default: - oss << "unknown"; + default: oss << "unknown"; }; return strdup(oss.str().c_str()); } diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index 16df4d59bd..c2db14f15f 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -863,13 +863,13 @@ hipError_t hipMemcpy3DBatchAsync(size_t numOps, struct hipMemcpy3DBatchOp* opLis unsigned long long flags, hipStream_t stream); hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms* p); hipError_t hipMemcpy3DPeerAsync(hipMemcpy3DPeerParms* p, hipStream_t stream); -hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* code, hipJitOption** jitOptions, +hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* code, hipJitOption* jitOptions, void** jitOptionsValues, unsigned int numJitOptions, - hipLibraryOption** libraryOptions, void** libraryOptionValues, + hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions); hipError_t hipLibraryLoadFromFile(hipLibrary_t* library, const char* fileName, - hipJitOption** jitOptions, void** jitOptionsValues, - unsigned int numJitOptions, hipLibraryOption** libraryOptions, + hipJitOption* jitOptions, void** jitOptionsValues, + unsigned int numJitOptions, hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions); hipError_t hipLibraryUnload(hipLibrary_t library); hipError_t hipLibraryGetKernel(hipKernel_t* pKernel, hipLibrary_t library, const char* name); diff --git a/projects/clr/hipamd/src/hip_library.cpp b/projects/clr/hipamd/src/hip_library.cpp index 42b3c203e7..2d69c7d918 100644 --- a/projects/clr/hipamd/src/hip_library.cpp +++ b/projects/clr/hipamd/src/hip_library.cpp @@ -106,9 +106,9 @@ hipError_t LibraryContainer::BuildIt() { return hipSuccess; } -hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* image, hipJitOption** jitOptions, +hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* image, hipJitOption* jitOptions, void** jitOptionsValues, unsigned int numJitOptions, - hipLibraryOption** libraryOptions, void** libraryOptionValues, + hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions) { HIP_INIT_API(hipLibraryLoadData, library, image, jitOptions, jitOptionsValues, numJitOptions, libraryOptions, libraryOptionValues, numLibraryOptions); @@ -127,8 +127,8 @@ hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* image, hipJitOp } hipError_t hipLibraryLoadFromFile(hipLibrary_t* library, const char* fname, - hipJitOption** jitOptions, void** jitOptionsValues, - unsigned int numJitOptions, hipLibraryOption** libraryOptions, + hipJitOption* jitOptions, void** jitOptionsValues, + unsigned int numJitOptions, hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions) { HIP_INIT_API(hipLibraryLoadFromFile, library, fname, jitOptions, jitOptionsValues, numJitOptions, libraryOptions, libraryOptionValues, numLibraryOptions); @@ -143,7 +143,7 @@ hipError_t hipLibraryLoadFromFile(hipLibrary_t* library, const char* fname, hipError_t hipLibraryUnload(hipLibrary_t library) { HIP_INIT_API(hipLibraryUnload, library); if (library == nullptr) { - HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidResourceHandle); } auto l = reinterpret_cast(library); delete l; diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index 1160a72745..90d27455f6 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -2011,17 +2011,17 @@ hipError_t hipGraphExecExternalSemaphoresWaitNodeSetParams( return hip::GetHipDispatchTable()->hipGraphExecExternalSemaphoresWaitNodeSetParams_fn( hGraphExec, hNode, nodeParams); } -hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* code, hipJitOption** jitOptions, +hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* code, hipJitOption* jitOptions, void** jitOptionsValues, unsigned int numJitOptions, - hipLibraryOption** libraryOptions, void** libraryOptionValues, + hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions) { return hip::GetHipDispatchTable()->hipLibraryLoadData_fn( library, code, jitOptions, jitOptionsValues, numJitOptions, libraryOptions, libraryOptionValues, numLibraryOptions); } hipError_t hipLibraryLoadFromFile(hipLibrary_t* library, const char* fileName, - hipJitOption** jitOptions, void** jitOptionsValues, - unsigned int numJitOptions, hipLibraryOption** libraryOptions, + hipJitOption* jitOptions, void** jitOptionsValues, + unsigned int numJitOptions, hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions) { return hip::GetHipDispatchTable()->hipLibraryLoadFromFile_fn( library, fileName, jitOptions, jitOptionsValues, numJitOptions, libraryOptions, diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index 7072943081..989f09e425 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -44,11 +44,15 @@ THE SOFTWARE. #define __HIP_ATOMIC_BACKWARD_COMPAT 1 #endif +#if HT_AMD #if defined(__has_extension) && __has_extension(clang_atomic_attributes) && __HIP_ATOMIC_BACKWARD_COMPAT #define HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY [[clang::atomic(fine_grained_memory, remote_memory)]] #else #define HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY #endif +#elif HT_NVIDIA +#define HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY +#endif #ifdef TEST_CLOCK_CYCLE #define clock_function() clock64() diff --git a/projects/hip-tests/catch/unit/library/CMakeLists.txt b/projects/hip-tests/catch/unit/library/CMakeLists.txt index 1e48f40179..88be4131bf 100644 --- a/projects/hip-tests/catch/unit/library/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/library/CMakeLists.txt @@ -12,9 +12,16 @@ add_custom_target(library_code_load.code set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${CMAKE_CURRENT_BINARY_DIR}/library_code_load.code) -hip_add_exe_to_target(NAME LibraryTests - TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests - LINKER_LIBS hiprtc) +if(HIP_PLATFORM MATCHES "amd") + hip_add_exe_to_target(NAME LibraryTests + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS hiprtc) +else() + hip_add_exe_to_target(NAME LibraryTests + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS nvrtc) +endif() add_dependencies(LibraryTests library_code_load.code) diff --git a/projects/hip-tests/catch/unit/library/library_negative.cc b/projects/hip-tests/catch/unit/library/library_negative.cc index d6b3b8dd4f..5624adbf44 100644 --- a/projects/hip-tests/catch/unit/library/library_negative.cc +++ b/projects/hip-tests/catch/unit/library/library_negative.cc @@ -29,7 +29,7 @@ TEST_CASE("Unit_library_negative") { HIP_CHECK_ERROR( hipLibraryLoadFromFile(nullptr, nullptr, nullptr, nullptr, 0, nullptr, nullptr, 0), hipErrorInvalidValue); - HIP_CHECK_ERROR(hipLibraryUnload(nullptr), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipLibraryUnload(nullptr), hipErrorInvalidResourceHandle); HIP_CHECK_ERROR(hipLibraryGetKernel(nullptr, nullptr, nullptr), hipErrorInvalidValue); HIP_CHECK_ERROR(hipLibraryGetKernelCount(nullptr, nullptr), hipErrorInvalidValue); } diff --git a/projects/hip-tests/catch/unit/library/loadlib_rtc.cc b/projects/hip-tests/catch/unit/library/loadlib_rtc.cc index 6a14bcabd8..18c463428a 100644 --- a/projects/hip-tests/catch/unit/library/loadlib_rtc.cc +++ b/projects/hip-tests/catch/unit/library/loadlib_rtc.cc @@ -29,7 +29,11 @@ THE SOFTWARE. static std::vector compile_using_hiprtc(const std::string& code, std::string gpu_arch) { hiprtcProgram prog; HIPRTC_CHECK(hiprtcCreateProgram(&prog, code.c_str(), "code.cu", 0, NULL, NULL)); - std::string offload_arch = "--offload-arch=" + gpu_arch; + #ifdef __HIP_PLATFORM_AMD__ + std::string offload_arch = "--offload-arch=" + gpu_arch; + #else + std::string offload_arch = "--fmad=false"; + #endif const char* opts[] = {offload_arch.c_str()}; HIPRTC_CHECK(hiprtcCompileProgram(prog, 1, opts)); size_t size; diff --git a/projects/hip-tests/catch/unit/memory/hipHostRegister.cc b/projects/hip-tests/catch/unit/memory/hipHostRegister.cc index 33ef4a24ca..667b44f77d 100644 --- a/projects/hip-tests/catch/unit/memory/hipHostRegister.cc +++ b/projects/hip-tests/catch/unit/memory/hipHostRegister.cc @@ -930,9 +930,11 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Flags", "", int, float, double) { #endif FlagType{0xF0, false}, FlagType{0xFFF2, false}, FlagType{0xFFFFFFFF, false}); +#if (HT_AMD == 1) && (HT_LINUX == 1) if (IsNavi4X() && (flags.value & hipExtHostRegisterUncached)) { return; } +#endif INFO("Testing hipHostRegister flag: " << flags.value); if (flags.valid) { HIP_CHECK(hipHostRegister(hostPtr, sizeBytes, flags.value)); @@ -942,7 +944,6 @@ TEMPLATE_TEST_CASE("Unit_hipHostRegister_Flags", "", int, float, double) { } free(hostPtr); } - /** * Test Description * ------------------------ diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index d91d9eeb5a..d354e0b8e4 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -6361,9 +6361,9 @@ hipError_t hipModuleGetFunctionCount(unsigned int* count, hipModule_t mod); * @param [in] numLibraryOptions Number of library options * @return #hipSuccess, #hipErrorInvalidValue, */ -hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* code, hipJitOption** jitOptions, +hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* code, hipJitOption* jitOptions, void** jitOptionsValues, unsigned int numJitOptions, - hipLibraryOption** libraryOptions, void** libraryOptionValues, + hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions); /** @@ -6380,8 +6380,8 @@ hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* code, hipJitOpt * @return #hipSuccess, #hipErrorInvalidValue */ hipError_t hipLibraryLoadFromFile(hipLibrary_t* library, const char* fileName, - hipJitOption** jitOptions, void** jitOptionsValues, - unsigned int numJitOptions, hipLibraryOption** libraryOptions, + hipJitOption* jitOptions, void** jitOptionsValues, + unsigned int numJitOptions, hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions); /** diff --git a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index 5172a494fd..9cc9975a62 100644 --- a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -482,7 +482,7 @@ typedef enum cudaSharedMemConfig hipSharedMemConfig; typedef CUfunc_cache hipFuncCache; typedef CUjitInputType hipJitInputType; typedef CUjit_option hipJitOption; -typedef enum cudaLibraryOption hipLibraryOption; +typedef CUlibraryOption hipLibraryOption; typedef CUdevice hipDevice_t; typedef enum cudaDeviceP2PAttr hipDeviceP2PAttr; #define hipDevP2PAttrPerformanceRank cudaDevP2PAttrPerformanceRank @@ -492,12 +492,15 @@ typedef enum cudaDeviceP2PAttr hipDeviceP2PAttr; #define hipFuncAttributeMaxDynamicSharedMemorySize cudaFuncAttributeMaxDynamicSharedMemorySize #define hipFuncAttributePreferredSharedMemoryCarveout cudaFuncAttributePreferredSharedMemoryCarveout +#define hipLibraryHostUniversalFunctionAndDataTable CU_LIBRARY_HOST_UNIVERSAL_FUNCTION_AND_DATA_TABLE +#define hipLibraryBinaryIsPreserved CU_LIBRARY_BINARY_IS_PRESERVED + typedef CUlinkState hipLinkState_t; typedef CUmodule hipModule_t; typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; -typedef cudaLibrary_t hipLibrary_t; -typedef cudaKernel_t hipKernel_t; +typedef CUlibrary hipLibrary_t; +typedef CUkernel hipKernel_t; typedef struct cudaArray* hipArray_t; typedef struct cudaArray* hipArray_const_t; typedef struct cudaFuncAttributes hipFuncAttributes; @@ -1790,7 +1793,6 @@ typedef cudaGraphEdgeData hipGraphEdgeData; typedef cudaLaunchConfig_t hipLaunchConfig_t; typedef cudaLaunchAttribute hipLaunchAttribute; typedef CUlaunchAttribute hipDrvLaunchAttribute; -typedef cudaKernel_t hipKernel_t; typedef CUlaunchConfig HIP_LAUNCH_CONFIG; typedef CUlaunchAttributeID hipDrvLaunchAttributeID; typedef CUlaunchAttributeValue hipDrvLaunchAttributeValue; @@ -3628,38 +3630,38 @@ inline static hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* im } inline static hipError_t hipLibraryLoadData(hipLibrary_t* library, const void* code, - hipJitOption** jitOptions, void** jitOptionsValues, + hipJitOption* jitOptions, void** jitOptionsValues, unsigned int numJitOptions, - hipLibraryOption** libraryOptions, + hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions) { - return hipCUResultTohipError(cudaLibraryLoadData(library, code, jitOptions, jitOptionsValues, + return hipCUResultTohipError(cuLibraryLoadData(library, code, jitOptions, jitOptionsValues, numJitOptions, libraryOptions, libraryOptionValues, numLibraryOptions)); } inline static hipError_t hipLibraryLoadFromFile(hipLibrary_t* library, const char* fileName, - hipJitOption** jitOptions, void** jitOptionsValues, + hipJitOption* jitOptions, void** jitOptionsValues, unsigned int numJitOptions, - hipLibraryOption** libraryOptions, + hipLibraryOption* libraryOptions, void** libraryOptionValues, unsigned int numLibraryOptions) { return hipCUResultTohipError( - cudaLibraryLoadFromFile(library, fileName, jitOptions, jitOptionsValues, numJitOptions, + cuLibraryLoadFromFile(library, fileName, jitOptions, jitOptionsValues, numJitOptions, libraryOptions, libraryOptionValues, numLibraryOptions)); } inline static hipError_t hipLibraryUnload(hipLibrary_t library) { - return hipCUResultTohipError(cudaLibraryUnload(library)); + return hipCUResultTohipError(cuLibraryUnload(library)); } inline static hipError_t hipLibraryGetKernel(hipKernel_t* pKernel, hipLibrary_t library, const char* name) { - return hipCUResultTohipError(cudaLibraryGetKernel(pKernel, library, name)); + return hipCUResultTohipError(cuLibraryGetKernel(pKernel, library, name)); } inline static hipError_t hipLibraryGetKernelCount(unsigned int* count, hipLibrary_t library) { - return hipCUResultTohipError(cudaLibraryGetKernelCount(count, library)); + return hipCUResultTohipError(cuLibraryGetKernelCount(count, library)); } inline static hipError_t hipLaunchKernel(const void* function_address, dim3 numBlocks, diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h index 105774ceb7..1ecca43ad6 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/hip/api_args.h @@ -3288,25 +3288,25 @@ typedef union rocprofiler_hip_api_args_t #if HIP_RUNTIME_API_TABLE_STEP_VERSION >= 15 struct { - hipLibrary_t* library; - const void* code; - hipJitOption** jitOptions; - void** jitOptionsValues; - unsigned int numJitOptions; - hipLibraryOption** libraryOptions; - void** libraryOptionValues; - unsigned int numLibraryOptions; + hipLibrary_t* library; + const void* code; + hipJitOption* jitOptions; + void** jitOptionsValues; + unsigned int numJitOptions; + hipLibraryOption* libraryOptions; + void** libraryOptionValues; + unsigned int numLibraryOptions; } hipLibraryLoadData; struct { - hipLibrary_t* library; - const char* fileName; - hipJitOption** jitOptions; - void** jitOptionsValues; - unsigned int numJitOptions; - hipLibraryOption** libraryOptions; - void** libraryOptionValues; - unsigned int numLibraryOptions; + hipLibrary_t* library; + const char* fileName; + hipJitOption* jitOptions; + void** jitOptionsValues; + unsigned int numJitOptions; + hipLibraryOption* libraryOptions; + void** libraryOptionValues; + unsigned int numLibraryOptions; } hipLibraryLoadFromFile; struct {