From 30185b4114ec7acbb99ba5f51561f07394d60147 Mon Sep 17 00:00:00 2001 From: Michael Liao Date: Fri, 8 Mar 2019 09:08:39 -0500 Subject: [PATCH 1/9] [hip] Re-implement hipLaunchKernelGGL as macros. --- include/hip/hcc_detail/hip_runtime.h | 19 ++++++++----------- 1 file changed, 8 insertions(+), 11 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 60d145c884..88cdf7e9ba 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -331,18 +331,15 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri typedef int hipLaunchParm; -template -inline void hipLaunchKernelGGL(F&& kernelName, const dim3& numblocks, const dim3& numthreads, - unsigned memperblock, hipStream_t streamId, Args... args) { - kernelName<<>>(args...); -} +#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ + do { \ + kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(hipLaunchParam{}, ##__VA_ARGS__); \ + } while (0) -template -inline void hipLaunchKernel(F&& kernel, const dim3& numBlocks, const dim3& dimBlocks, - std::uint32_t groupMemBytes, hipStream_t stream, Args... args) { - hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{}, - std::move(args)...); -} +#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ + do { \ + kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(__VA_ARGS__); \ + } while (0) #include From 950b6efe725c20578ac2b8302a75af05c2f5262d Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Fri, 15 Mar 2019 12:35:01 -0400 Subject: [PATCH 2/9] [Device Function] Support immediate argument. - `immarg`, immediate argument, is enabled on all AMDGPU intrinsics. Revise device functions using these intrinsics with immediate arguments. --- include/hip/hcc_detail/device_functions.h | 81 ++++++++++++----------- include/hip/hcc_detail/hip_runtime.h | 9 ++- 2 files changed, 48 insertions(+), 42 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 602c6be87a..3b18e69989 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -46,7 +46,7 @@ __device__ static inline unsigned int __popcll(unsigned long long int input) { } __device__ static inline int __clz(int input) { - return __ockl_clz_u32((uint)input); + return __ockl_clz_u32((uint)input); } __device__ static inline int __clzll(long long int input) { @@ -224,56 +224,59 @@ __device__ static inline unsigned int __lane_id() { return __mbcnt_hi(-1, __mbc HIP specific device functions */ -// utility union type -union __u { - int i; - unsigned int u; - float f; -}; - __device__ static inline unsigned __hip_ds_bpermute(int index, unsigned src) { - __u tmp; tmp.u = src; + union { int i; unsigned u; float f; } tmp; tmp.u = src; tmp.i = __llvm_amdgcn_ds_bpermute(index, tmp.i); return tmp.u; } __device__ static inline float __hip_ds_bpermutef(int index, float src) { - __u tmp; tmp.f = src; + union { int i; unsigned u; float f; } tmp; tmp.f = src; tmp.i = __llvm_amdgcn_ds_bpermute(index, tmp.i); return tmp.f; } __device__ static inline unsigned __hip_ds_permute(int index, unsigned src) { - __u tmp; tmp.u = src; - tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i); - return tmp.u; + union { int i; unsigned u; float f; } tmp; tmp.u = src; + tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i); + return tmp.u; } __device__ static inline float __hip_ds_permutef(int index, float src) { - __u tmp; tmp.u = src; - tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i); - return tmp.u; -} - -__device__ static inline unsigned __hip_ds_swizzle(unsigned int src, int pattern) { - __u tmp; tmp.u = src; - tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern); + union { int i; unsigned u; float f; } tmp; tmp.u = src; + tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i); return tmp.u; } -__device__ static inline float __hip_ds_swizzlef(float src, int pattern) { - __u tmp; tmp.f = src; - tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern); + +#define __hip_ds_swizzle(src, pattern) __hip_ds_swizzle_N<(pattern)>((src)) +#define __hip_ds_swizzlef(src, pattern) __hip_ds_swizzlef_N<(pattern)>((src)) + +template +__device__ static inline unsigned __hip_ds_swizzle_N(unsigned int src) { + union { int i; unsigned u; float f; } tmp; tmp.u = src; + tmp.i = __builtin_amdgcn_ds_swizzle(tmp.i, pattern); + return tmp.u; +} + +template +__device__ static inline float __hip_ds2_swizzlef_N(float src) { + union { int i; unsigned u; float f; } tmp; tmp.f = src; + tmp.i = __builtin_amdgcn_ds_swizzle(tmp.i, pattern); return tmp.f; } -__device__ static inline int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, - int bank_mask, bool bound_ctrl) { - return __llvm_amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl); +#define __hip_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl) \ + __hip_move_dpp_N<(dpp_ctrl), (row_mask), (bank_mask), (bound_ctrl)>((src)) + +template +__device__ static inline int __hip2_move_dpp_N(int src) { + return __llvm_amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, + bound_ctrl); } static constexpr int warpSize = 64; - __device__ +__device__ inline int __shfl(int var, int src_lane, int width = warpSize) { int self = __lane_id(); @@ -283,14 +286,14 @@ int __shfl(int var, int src_lane, int width = warpSize) { __device__ inline unsigned int __shfl(unsigned int var, int src_lane, int width = warpSize) { - __u tmp; tmp.u = var; + union { int i; unsigned u; float f; } tmp; tmp.u = var; tmp.i = __shfl(tmp.i, src_lane, width); return tmp.u; } __device__ inline float __shfl(float var, int src_lane, int width = warpSize) { - __u tmp; tmp.f = var; + union { int i; unsigned u; float f; } tmp; tmp.f = var; tmp.i = __shfl(tmp.i, src_lane, width); return tmp.f; } @@ -320,14 +323,14 @@ int __shfl_up(int var, unsigned int lane_delta, int width = warpSize) { __device__ inline unsigned int __shfl_up(unsigned int var, unsigned int lane_delta, int width = warpSize) { - __u tmp; tmp.u = var; + union { int i; unsigned u; float f; } tmp; tmp.u = var; tmp.i = __shfl_up(tmp.i, lane_delta, width); return tmp.u; } __device__ inline float __shfl_up(float var, unsigned int lane_delta, int width = warpSize) { - __u tmp; tmp.f = var; + union { int i; unsigned u; float f; } tmp; tmp.f = var; tmp.i = __shfl_up(tmp.i, lane_delta, width); return tmp.f; } @@ -357,14 +360,14 @@ int __shfl_down(int var, unsigned int lane_delta, int width = warpSize) { __device__ inline unsigned int __shfl_down(unsigned int var, unsigned int lane_delta, int width = warpSize) { - __u tmp; tmp.u = var; + union { int i; unsigned u; float f; } tmp; tmp.u = var; tmp.i = __shfl_down(tmp.i, lane_delta, width); return tmp.u; } __device__ inline float __shfl_down(float var, unsigned int lane_delta, int width = warpSize) { - __u tmp; tmp.f = var; + union { int i; unsigned u; float f; } tmp; tmp.f = var; tmp.i = __shfl_down(tmp.i, lane_delta, width); return tmp.f; } @@ -394,14 +397,14 @@ int __shfl_xor(int var, int lane_mask, int width = warpSize) { __device__ inline unsigned int __shfl_xor(unsigned int var, int lane_mask, int width = warpSize) { - __u tmp; tmp.u = var; + union { int i; unsigned u; float f; } tmp; tmp.u = var; tmp.i = __shfl_xor(tmp.i, lane_mask, width); return tmp.u; } __device__ inline float __shfl_xor(float var, int lane_mask, int width = warpSize) { - __u tmp; tmp.f = var; + union { int i; unsigned u; float f; } tmp; tmp.f = var; tmp.i = __shfl_xor(tmp.i, lane_mask, width); return tmp.f; } @@ -681,9 +684,9 @@ inline __attribute((always_inline)) long long int __clock64() { // ToDo: Unify HCC and HIP implementation. #if __HCC__ - return (long long int) __clock_u64(); + return (long long int) __clock_u64(); #else - return (long long int) __builtin_amdgcn_s_memrealtime(); + return (long long int) __builtin_amdgcn_s_memrealtime(); #endif } @@ -881,7 +884,7 @@ __device__ inline __attribute__((weak)) void abort() { - return __builtin_trap(); + return __builtin_trap(); } diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 60d145c884..df98b9cd63 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -190,10 +190,13 @@ __device__ float __hip_ds_bpermutef(int index, float src); __device__ unsigned __hip_ds_permute(int index, unsigned src); __device__ float __hip_ds_permutef(int index, float src); -__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern); -__device__ float __hip_ds_swizzlef(float src, int pattern); +template +__device__ unsigned __hip_ds_swizzle_N(unsigned int src); +template +__device__ float __hip_ds_swizzlef_N(float src); -__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl); +template +__device__ int __hip_move_dpp_N(int src); #endif //__HIP_ARCH_GFX803__ == 1 From 0bddfc09e192ea820a267b3e032e9a4267f8f6be Mon Sep 17 00:00:00 2001 From: Nicholas Malaya Date: Mon, 18 Mar 2019 14:51:18 -0500 Subject: [PATCH 3/9] Update hip_faq.md Making more clear what this list details. In particular, this list is intended to indicate what items for each CUDA release is supported, and which are not. --- docs/markdown/hip_faq.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/markdown/hip_faq.md b/docs/markdown/hip_faq.md index c311aae320..3cea1cb42c 100644 --- a/docs/markdown/hip_faq.md +++ b/docs/markdown/hip_faq.md @@ -77,7 +77,7 @@ HIP code provides the same performance as native CUDA code, plus the benefits of HIP APIs and features do not map to a specific CUDA version. HIP provides a strong subset of functionality provided in CUDA, and the hipify tools can scan code to identify any unsupported CUDA functions - this is useful for identifying the specific features required by a given application. -However, we can provide a rough summary of the features included in each CUDA SDK and the support level in HIP: +However, we can provide a rough summary of the features included in each CUDA SDK and the support level in HIP. Each bullet below lists the major new language features in each CUDA release and then indicate which are supported/not supported in HIP: - CUDA 4.0 and earlier : - HIP supports CUDA 4.0 except for the limitations described above. From e9bc4034d8ba2c17f2ca1bafcd277eced7780bc1 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Tue, 19 Mar 2019 11:11:49 -0400 Subject: [PATCH 4/9] Add declaration of symbol related API for VDI --- include/hip/hcc_detail/hip_runtime_api.h | 21 ++++++++++++++++++++- 1 file changed, 20 insertions(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index ea93d28bfe..7d227f2fac 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -1377,7 +1377,26 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream); -#if !__HIP_VDI__ +#if __HIP_VDI__ +hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, + hipModule_t hmod, const char* name); + +hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName); +hipError_t hipGetSymbolSize(size_t* size, const void* symbolName); +hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, + size_t sizeBytes, size_t offset __dparm(0), + hipMemcpyKind kind __dparm(hipMemcpyHostToDevice)); +hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, + size_t sizeBytes, size_t offset, + hipMemcpyKind kind, hipStream_t stream __dparm(0)); +hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, + size_t sizeBytes, size_t offset __dparm(0), + hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost)); +hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, + size_t sizeBytes, size_t offset, + hipMemcpyKind kind, + hipStream_t stream __dparm(0)); +#else __attribute__((visibility("hidden"))) hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*); From e51fa174c849a29a1e37ac97166505e0488c032e Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Tue, 19 Mar 2019 15:27:21 -0400 Subject: [PATCH 5/9] [Device Function] Fix typos. --- include/hip/hcc_detail/device_functions.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 3b18e69989..a24df22315 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -259,7 +259,7 @@ __device__ static inline unsigned __hip_ds_swizzle_N(unsigned int src) { } template -__device__ static inline float __hip_ds2_swizzlef_N(float src) { +__device__ static inline float __hip_ds_swizzlef_N(float src) { union { int i; unsigned u; float f; } tmp; tmp.f = src; tmp.i = __builtin_amdgcn_ds_swizzle(tmp.i, pattern); return tmp.f; @@ -269,7 +269,7 @@ __device__ static inline float __hip_ds2_swizzlef_N(float src) { __hip_move_dpp_N<(dpp_ctrl), (row_mask), (bank_mask), (bound_ctrl)>((src)) template -__device__ static inline int __hip2_move_dpp_N(int src) { +__device__ static inline int __hip_move_dpp_N(int src) { return __llvm_amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl); } From e2bf34cd5e6444ee04adae7c0c496bf52cff4f31 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Tue, 19 Mar 2019 16:25:07 -0400 Subject: [PATCH 6/9] Change HIP dll name to amdhip64.dll on Windows --- bin/hipcc | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/bin/hipcc b/bin/hipcc index d227eea267..22bf17f2f5 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -196,8 +196,10 @@ if ($HIP_PLATFORM eq "clang") { $HIPLDFLAGS .= " $HIP_DEVLIB_FLAGS -L$HIP_LIB_PATH"; if (not $isWindows) { $HIPLDFLAGS .= " -Wl,--rpath=$HIP_LIB_PATH"; + $HIPLDFLAGS .= " -lhip_hcc"; + } else { + $HIPLDFLAGS .= " -lamdhip64"; } - $HIPLDFLAGS .= " -lhip_hcc"; if ($HIP_CLANG_HCC_COMPAT_MODE) { ## Allow __fp16 as function parameter and return type. $HIPCXXFLAGS .= " -Xclang -fallow-half-arguments-and-returns -D__HIP_HCC_COMPAT_MODE__=1"; From 24d08beef8bb91c572c54d6bccf6d7144d21ee8a Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Wed, 20 Mar 2019 01:11:15 -0400 Subject: [PATCH 7/9] reimplement HIP_INIT as hip_impl::hip_init(), add hip_init() to some of the inlined API (#966) * reimplement HIP_INIT as a function, expose it as hip_impl::hip_init() so that it could be called from hipLaunchKernelGGL and other inlined HIP functions * Don't call hip_init from ihipPreLaunchKernel --- .../hip/hcc_detail/functional_grid_launch.hpp | 3 ++- include/hip/hcc_detail/hip_runtime_api.h | 8 ++++++-- src/hip_clang.cpp | 2 +- src/hip_fatbin.cpp | 2 +- src/hip_hcc.cpp | 18 ++++++++++-------- src/hip_hcc_internal.h | 11 ++--------- 6 files changed, 22 insertions(+), 22 deletions(-) diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index 29a389f642..0e541001bf 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -189,6 +189,7 @@ inline void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, std::uint32_t sharedMemBytes, hipStream_t stream, Args... args) { + hip_impl::hip_init(); auto kernarg = hip_impl::make_kernarg( kernel, std::tuple{std::move(args)...}); std::size_t kernarg_size = kernarg.size(); @@ -212,4 +213,4 @@ inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBloc std::uint32_t groupMemBytes, hipStream_t stream, Args... args) { hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{}, std::move(args)...); -} \ No newline at end of file +} diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index ea93d28bfe..3b4f8fac5c 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -78,6 +78,10 @@ THE SOFTWARE. #define __dparm(x) #endif +namespace hip_impl { +hipError_t hip_init(); +} // namespace hip_impl + // Structure definitions: #ifdef __cplusplus extern "C" { @@ -1396,7 +1400,7 @@ inline __attribute__((visibility("hidden"))) hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { //HIP_INIT_API(hipGetSymbolAddress, devPtr, symbolName); - + hip_impl::hip_init(); size_t size = 0; return hipModuleGetGlobal(devPtr, &size, 0, (const char*)symbolName); } @@ -1416,7 +1420,7 @@ inline __attribute__((visibility("hidden"))) hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) { // HIP_INIT_API(hipGetSymbolSize, size, symbolName); - + hip_impl::hip_init(); void* devPtr = nullptr; return hipModuleGetGlobal(&devPtr, size, 0, (const char*)symbolName); } diff --git a/src/hip_clang.cpp b/src/hip_clang.cpp index 0d210e2b8c..9af4f5a7a2 100644 --- a/src/hip_clang.cpp +++ b/src/hip_clang.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. extern "C" std::vector* __hipRegisterFatBinary(const void* data) { - HIP_INIT(); + hip_impl::hip_init(); tprintf(DB_FB, "Enter __hipRegisterFatBinary(%p)\n", data); const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast(data); diff --git a/src/hip_fatbin.cpp b/src/hip_fatbin.cpp index 8fe7740ed7..5aab0b7101 100644 --- a/src/hip_fatbin.cpp +++ b/src/hip_fatbin.cpp @@ -43,7 +43,7 @@ void __hipDumpCodeObject(const std::string& image) { const void* __hipExtractCodeObjectFromFatBinary(const void* data, const char* agent_name) { - HIP_INIT(); + hip_impl::hip_init(); tprintf(DB_FB, "Enter __hipExtractCodeObjectFromFatBinary(%p, \"%s\")\n", data, agent_name); diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 3f0c4ba6b8..8a291f3ea4 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -119,9 +119,6 @@ int HCC_OPT_FLUSH = 1; int HCC_OPT_FLUSH = 0; #endif - -std::once_flag hip_initialized; - // Array of pointers to devices. ihipDevice_t** g_deviceArray; @@ -1442,6 +1439,15 @@ void ihipInit() { g_numLogicalThreads); } +namespace hip_impl { +hipError_t hip_init() { + static std::once_flag hip_initialized; + std::call_once(hip_initialized, ihipInit); + ihipCtxStackUpdate(); + return hipSuccess; +} +} + hipError_t ihipStreamSynchronize(hipStream_t stream) { hipError_t e = hipSuccess; @@ -1561,7 +1567,6 @@ void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp, // Allows runtime to track some information about the stream. hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm* lp, const char* kernelNameStr) { - HIP_INIT(); stream = ihipSyncAndResolveStream(stream); lp->grid_dim.x = grid.x; lp->grid_dim.y = grid.y; @@ -1583,7 +1588,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm* lp, const char* kernelNameStr) { - HIP_INIT(); stream = ihipSyncAndResolveStream(stream); lp->grid_dim.x = grid; lp->grid_dim.y = 1; @@ -1604,7 +1608,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm* lp, const char* kernelNameStr) { - HIP_INIT(); stream = ihipSyncAndResolveStream(stream); lp->grid_dim.x = grid.x; lp->grid_dim.y = grid.y; @@ -1625,7 +1628,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm* lp, const char* kernelNameStr) { - HIP_INIT(); stream = ihipSyncAndResolveStream(stream); lp->grid_dim.x = grid; lp->grid_dim.y = 1; @@ -2485,4 +2487,4 @@ namespace hip_impl { std::terminate(); #endif } -} // Namespace hip_impl. \ No newline at end of file +} // Namespace hip_impl. diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index be794956f6..95a65b55c9 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -288,19 +288,13 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr); #define API_TRACE(IS_CMD, ...) tls_tidInfo.incApiSeqNum(); #endif - -// Just initialize the HIP runtime, but don't log any trace information. -#define HIP_INIT() \ - std::call_once(hip_initialized, ihipInit); \ - ihipCtxStackUpdate(); #define HIP_SET_DEVICE() ihipDeviceSetState(); - // This macro should be called at the beginning of every HIP API. // It initializes the hip runtime (exactly once), and // generates a trace string that can be output to stderr or to ATP file. #define HIP_INIT_API(cid, ...) \ - HIP_INIT() \ + hip_impl::hip_init(); \ API_TRACE(0, __VA_ARGS__); \ HIP_CB_SPAWNER_OBJECT(cid); @@ -309,7 +303,7 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr); // Replace HIP_INIT_API with this call inside HIP APIs that launch work on the GPU: // kernel launches, copy commands, memory sets, etc. #define HIP_INIT_SPECIAL_API(cid, tbit, ...) \ - HIP_INIT() \ + hip_impl::hip_init(); \ API_TRACE((HIP_TRACE_API & (1 << tbit)), __VA_ARGS__); \ HIP_CB_SPAWNER_OBJECT(cid); @@ -933,7 +927,6 @@ class ihipCtx_t { //================================================================================================= // Global variable definition: -extern std::once_flag hip_initialized; extern unsigned g_deviceCnt; extern hsa_agent_t g_cpu_agent; // the CPU agent. extern hsa_agent_t* g_allAgents; // CPU agents + all the visible GPU agents. From 673cdeea55c8ec4cb3e59c2e77346b1035b13288 Mon Sep 17 00:00:00 2001 From: eshcherb <33529668+eshcherb@users.noreply.github.com> Date: Wed, 20 Mar 2019 00:11:40 -0500 Subject: [PATCH 8/9] adding prof primitives generator (#967) * adding prof primitives generator * minor change, renaming * minor cosmetic changes, comments correcting and dead code removing * minor changes and renaming * minor chane, fixing comments --- .gitignore | 1 + CMakeLists.txt | 19 +- hip_prof_gen.py | 454 +++++ include/hip/hcc_detail/hip_prof_str.h | 2378 ------------------------- 4 files changed, 471 insertions(+), 2381 deletions(-) create mode 100755 hip_prof_gen.py delete mode 100644 include/hip/hcc_detail/hip_prof_str.h diff --git a/.gitignore b/.gitignore index 3dfb1afef5..3c94b6aaa0 100644 --- a/.gitignore +++ b/.gitignore @@ -13,5 +13,6 @@ bin/hipInfo bin/hipBusBandwidth bin/hipDispatchLatency bin/hipify-clang +include/hip/hcc_detail/hip_prof_str.h samples/1_Utils/hipInfo/hipInfo diff --git a/CMakeLists.txt b/CMakeLists.txt index d1a1d6c5a0..115f746cd9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -160,9 +160,22 @@ if(NOT DEFINED COMPILE_HIP_ATP_MARKER) endif() add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER) -################ -# Detect profiling API -################ +############################# +# Profiling API support +############################# +# Generate profiling API macros/structures header +set(PROF_API_STR "${CMAKE_CURRENT_SOURCE_DIR}/include/hip/hcc_detail/hip_prof_str.h") +set(PROF_API_HDR "${CMAKE_CURRENT_SOURCE_DIR}/include/hip/hcc_detail/hip_runtime_api.h") +set(PROF_API_SRC "${CMAKE_CURRENT_SOURCE_DIR}/src") +set(PROF_API_GEN "${CMAKE_CURRENT_SOURCE_DIR}/hip_prof_gen.py") +add_custom_command( + OUTPUT ${PROF_API_STR} + COMMAND ${PROF_API_GEN} ${OPT_PROF_API} ${PROF_API_HDR} ${PROF_API_SRC} ${PROF_API_STR} + DEPENDS ${PROF_API_GEN} ${PROF_API_HDR} +) +add_custom_target(prof ALL DEPENDS ${PROF_API_STR}) + +# Enable profiling API if(USE_PROF_API EQUAL 1) find_path(PROF_API_HEADER_DIR prof_protocol.h HINTS diff --git a/hip_prof_gen.py b/hip_prof_gen.py new file mode 100755 index 0000000000..81df5132dd --- /dev/null +++ b/hip_prof_gen.py @@ -0,0 +1,454 @@ +#!/usr/bin/python +import os, sys, re + +verbose = 0 +PROF_HEADER = "hip_prof_str.h" +OUTPUT = PROF_HEADER +REC_MAX_LEN = 1024 + +# Fatal error termination +inp_file = 'none' +line_num = -1 +def fatal(msg): + if line_num != -1: + print >>sys.stderr, "Error: " + msg + ", file '" + inp_file + "', line (" + str(line_num) + ")" + else: + print >>sys.stderr, "Error: " + msg + sys.exit(1) + +# Verbose message +def message(msg): + if verbose: print >>sys.stdout, msg + +############################################################# +# Normalizing API arguments +def filtr_api_args(args_str): + args_str = re.sub(r'^\s*', r'', args_str); + args_str = re.sub(r'\s*$', r'', args_str); + args_str = re.sub(r'\s*,\s*', r',', args_str); + args_str = re.sub(r'\s+', r' ', args_str); + args_str = re.sub(r'void \*', r'void* ', args_str); + args_str = re.sub(r'(enum|struct) ', '', args_str); + return args_str + +# Normalizing types +def norm_api_types(type_str): + type_str = re.sub(r'uint32_t', r'unsigned int', type_str) + type_str = re.sub(r'^unsigned$', r'unsigned int', type_str) + return type_str + +# Creating a list of arguments [(type, name), ...] +def list_api_args(args_str): + args_str = filtr_api_args(args_str) + args_list = [] + if args_str != '': + for arg_pair in args_str.split(','): + if arg_pair == 'void': continue + arg_pair = re.sub(r'\s*=\s*\S+$','', arg_pair); + m = re.match("^(.*)\s(\S+)$", arg_pair); + if m: + arg_type = norm_api_types(m.group(1)) + arg_name = m.group(2) + args_list.append((arg_type, arg_name)) + else: + fatal("bad args: args_str: '" + args_str + "' arg_pair: '" + arg_pair + "'") + return args_list; + +# Creating arguments string "type0, type1, ..." +def filtr_api_types(args_str): + args_list = list_api_args(args_str) + types_str = '' + for arg_tuple in args_list: + types_str += arg_tuple[0] + ', ' + return types_str + +# Creating options list [opt0, opt1, ...] +def filtr_api_opts(args_str): + args_list = list_api_args(args_str) + opts_list = [] + for arg_tuple in args_list: + opts_list.append(arg_tuple[1]) + return opts_list +############################################################# +# Parsing API header +# hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset); +def parse_api(inp_file_p, out): + global inp_file + global line_num + inp_file = inp_file_p + + beg_pattern = re.compile("^(hipError_t|const char\s*\*)\s+[^\(]+\("); + api_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)"); + end_pattern = re.compile("Texture"); + hidden_pattern = re.compile(r'__attribute__\(\(visibility\("hidden"\)\)\)') + nms_open_pattern = re.compile(r'namespace hip_impl {') + nms_close_pattern = re.compile(r'}') + + inp = open(inp_file, 'r') + + found = 0 + hidden = 0 + nms_level = 0; + record = "" + line_num = -1 + + for line in inp.readlines(): + record += re.sub(r'^\s+', r' ', line[:-1]) + line_num += 1 + + if len(record) > REC_MAX_LEN: + fatal("bad record \"" + record + "\"") + + if beg_pattern.match(record) and (hidden == 0) and (nms_level == 0): found = 1 + + if found != 0: + record = re.sub("\s__dparm\([^\)]*\)", '', record); + m = api_pattern.match(record) + if m: + found = 0 + if end_pattern.search(record): break + out[m.group(2)] = m.group(3) + else: continue + + hidden = 0 + if hidden_pattern.match(line): hidden = 1 + + if nms_open_pattern.match(line): nms_level += 1 + if (nms_level > 0) and nms_close_pattern.match(line): nms_level -= 1 + if nms_level < 0: + fatal("nms level < 0") + + record = "" + + inp.close() + line_num = -1 +############################################################# +# Parsing API implementation +# hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset) { +# HIP_INIT_CB(hipSetupArgument, arg, size, offset); +# inp_file - input implementation source file +# api_map - input public API map [] => +# out - output map [] => [opt0, opt1, ...] +def parse_content(inp_file_p, api_map, out): + global inp_file + global line_num + inp_file = inp_file_p + + # API definition begin pattern + beg_pattern = re.compile("^(hipError_t|const char\s*\*)\s+[^\(]+\("); + # API definition complete pattern + api_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*{"); + # API init macro pattern + init_pattern = re.compile("^\s*HIP_INIT[_\w]*_API\(([^,]+)(,|\))"); + target_pattern = re.compile("^(\s*HIP_INIT[^\(]*)(_API\()(.*)\);\s*$"); + + # Open input file + inp = open(inp_file, 'r') + + # API name + api_name = "" + # Valid public API found flag + api_valid = 0 + + # Input file patched content + content = '' + # Sub content for found API defiition + sub_content = '' + # Current record, accumulating several API definition related lines + record = '' + # Current input file line number + line_num = -1 + # API beginning found flag + found = 0 + + # Reading input file + for line in inp.readlines(): + # Accumulating record + record += re.sub(r'^\s+', r' ', line[:-1]) + line_num += 1 + + if len(record) > REC_MAX_LEN: + fatal("bad record \"" + record + "\"") + break; + + # Looking for API begin + if beg_pattern.match(record): found = 1 + + # Matching complete API definition + if found == 1: + record = re.sub("\s__dparm\([^\)]*\)", '', record); + m = api_pattern.match(record) + # Checking if complete API matched + if m: + found = 2 + api_name = m.group(2); + # Checking if API name is in the API map + if api_name in api_map: + # Getting API arguments + api_args = m.group(3) + # Getting etalon arguments from the API map + eta_args = api_map[api_name] + if eta_args == '': + eta_args = api_args + api_map[api_name] = eta_args + # Normalizing API arguments + api_types = filtr_api_types(api_args) + # Normalizing etalon arguments + eta_types = filtr_api_types(eta_args) + if api_types == eta_types: + # API is already found + if api_name in out: + fatal("API redefined \"" + api_name + "\", record \"" + record + "\"") + # Set valid public API found flag + api_valid = 1 + # Set output API map with API arguments list + out[api_name] = filtr_api_opts(api_args) + else: + # Warning about mismatched API, possible non public overloaded version + api_diff = '\t\t' + inp_file + " line(" + str(line_num) + ")\n\t\tapi: " + api_types + "\n\t\teta: " + eta_types + message("\t" + api_name + ':\n' + api_diff + '\n') + + # API found action + if found == 2: + # Looking for INIT macro + m = init_pattern.match(line) + if m: + found = 0 + if api_valid == 1: + api_valid = 0 + message("\t" + api_name) + else: + # Registering dummy API for non public API if the name in INIT is not NONE + init_name = m.group(1) + # Ignore if it is initialized as NONE + if init_name != 'NONE': + # Check if init name matching API name + if init_name != api_name: + fatal("init name mismatch: '" + init_name + "' <> '" + api_name + "'") + # If init name is not in public API map then it is private API + # else it was not identified and will be checked on finish + if not init_name in api_map: + if init_name in out: + fatal("API reinit \"" + api_name + "\", record \"" + record + "\"") + out[init_name] = [] + elif re.search('}', line): + found = 0 + # Expect INIT macro for valid public API + if api_valid == 1: + api_valid = 0 + if api_name in out: + del out[api_name] + del api_map[api_name] + out['.' + api_name] = 1 + else: + fatal("API is not in out \"" + api_name + "\", record \"" + record + "\"") + + if found != 1: record = "" + content += line + + inp.close() + line_num = -1 + + if len(out) != 0: + return content + else: + return '' + +# src path walk +def parse_src(api_map, src_path, src_patt, out): + pattern = re.compile(src_patt) + src_path = re.sub(r'\s', '', src_path) + for src_dir in src_path.split(':'): + message("Parsing " + src_dir + " for '" + src_patt + "'") + for root, dirs, files in os.walk(src_dir): + for fnm in files: + if pattern.search(fnm): + file = root + '/' + fnm + message(file) + content = parse_content(file, api_map, out); + if content != '': + f = open(file, 'w') + f.write(content) + f.close() +############################################################# +# Generating profiling primitives header +# api_map - public API map [] => [(type, name), ...] +# opts_map - opts map [] => [opt0, opt1, ...] +def generate_prof_header(f, api_map, opts_map): + # Private API list + priv_lst = [] + + f.write('// automatically generated sources\n') + f.write('#ifndef _HIP_PROF_STR_H\n'); + f.write('#define _HIP_PROF_STR_H\n'); + f.write('#include \n'); + f.write('#include \n'); + + # Generating dummy macro for non-public API + f.write('\n// Dummy API primitives\n') + f.write('#define INIT_NONE_CB_ARGS_DATA(cb_data) {};\n') + for name in opts_map: + if not name in api_map: + opts_lst = opts_map[name] + if len(opts_lst) != 0: + fatal("bad dummy API \"" + name + "\", args: " + str(opts_lst)) + f.write('#define INIT_'+ name + '_CB_ARGS_DATA(cb_data) {};\n') + priv_lst.append(name) + + for name in priv_lst: + message("Private: " + name) + + # Generating the callbacks ID enumaration + f.write('\n// HIP API callbacks ID enumaration\n') + f.write('enum hip_api_id_t {\n') + cb_id = 0 + for name in api_map.keys(): + f.write(' HIP_API_ID_' + name + ' = ' + str(cb_id) + ',\n') + cb_id += 1 + f.write(' HIP_API_ID_NUMBER = ' + str(cb_id) + ',\n') + f.write(' HIP_API_ID_ANY = ' + str(cb_id + 1) + ',\n') + f.write('\n') + f.write(' HIP_API_ID_NONE = HIP_API_ID_NUMBER,\n') + for name in priv_lst: + f.write(' HIP_API_ID_' + name + ' = HIP_API_ID_NUMBER,\n') + f.write('};\n') + + # Generating the callbacks ID enumaration + f.write('\n// Return HIP API string\n') + f.write('static const char* hip_api_name(const uint32_t& id) {\n') + f.write(' switch(id) {\n') + for name in api_map.keys(): + f.write(' case HIP_API_ID_' + name + ': return "' + name + '";\n') + f.write(' };\n') + f.write(' return "unknown";\n') + f.write('};\n') + + # Generating the callbacks data structure + f.write('\n// HIP API callbacks data structure\n') + f.write( + 'struct hip_api_data_t {\n' + + ' uint64_t correlation_id;\n' + + ' uint32_t phase;\n' + + ' union {\n' + ) + for name, args in api_map.items(): + if len(args) != 0: + f.write(' struct {\n') + for arg_tuple in args: + f.write(' ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') + f.write(' } ' + name + ';\n') + f.write( + ' } args;\n' + + '};\n' + ) + + # Generating the callbacks args data filling macros + f.write('\n// HIP API callbacks args data filling macros\n') + for name, args in api_map.items(): + f.write('// ' + name + str(args) + '\n') + f.write('#define INIT_' + name + '_CB_ARGS_DATA(cb_data) { \\\n') + if name in opts_map: + opts_list = opts_map[name] + if len(args) != len(opts_list): + fatal("\"" + name + "\" API args and opts mismatch, args: " + str(args) + ", opts: " + str(opts_list)) + # API args iterating: + # type is args[][0] + # name is args[][1] + for ind in range(0, len(args)): + arg_tuple = args[ind] + fld_name = arg_tuple[1] + arg_name = opts_list[ind] + f.write(' cb_data.args.' + name + '.' + fld_name + ' = ' + arg_name + '; \\\n') + f.write('};\n') + f.write('#define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)\n') + + # Generating the method for the API string, name and parameters + f.write('\n') + f.write('#if 0\n') + f.write('// HIP API string method, method name and parameters\n') + f.write('const char* hipApiString(hip_api_id_t id, const hip_api_data_t* data) {\n') + f.write(' std::ostringstream oss;\n') + f.write(' switch (id) {\n') + for name, args in api_map.items(): + f.write(' case HIP_API_ID_' + name + ':\n') + f.write(' oss << "' + name + '("') + for ind in range(0, len(args)): + arg_tuple = args[ind] + arg_name = arg_tuple[1] + if ind != 0: f.write(' << ","') + f.write('\n << " ' + arg_name + '=" << data->args.' + name + '.' + arg_name) + f.write('\n << ")";\n') + f.write(' break;\n') + f.write(' default: oss << "unknown";\n') + f.write(' };\n') + f.write(' return strdup(oss.str().c_str());\n') + f.write('};\n') + f.write('#endif\n') + + f.write('#endif // _HIP_PROF_STR_H\n'); + +############################################################# +# main +# Usage +if (len(sys.argv) > 1) and (sys.argv[1] == '-v'): + verbose = 1 + sys.argv.pop(1) + +if (len(sys.argv) < 3): + fatal ("Usage: " + sys.argv[0] + " [-v] \n" + + " -v - verbose messages\n" + + " example:\n" + + " $ hipap.py hip/include/hip/hcc_detail/hip_runtime_api.h hip/src") + +# API header file given as an argument +api_hfile = sys.argv[1] +if not os.path.isfile(api_hfile): + fatal("input file '" + api_hfile + "' not found") + +# Srcs directory given as an argument +src_pat = "\.cpp$" +src_dir = sys.argv[2] +if not os.path.isdir(src_dir): + fatal("src directory " + src_dir + "' not found") + +if len(sys.argv) > 3: OUTPUT = sys.argv[3] + +# API declaration map +api_map = { + 'hipHccModuleLaunchKernel': '' +} +# API options map +opts_map = {} + +# Parsing API header +parse_api(api_hfile, api_map) + +# Parsing sources +parse_src(api_map, src_dir, src_pat, opts_map) + +# Checking for non-conformant APIs +for name in opts_map.keys(): + m = re.match(r'\.(\S*)', name) + if m: + message("Init missing: " + m.group(1)) + del opts_map[name] + +# Converting api map to map of lists +# Checking for not found APIs +not_found = 0 +if len(opts_map) != 0: + for name in api_map.keys(): + args_str = api_map[name]; + api_map[name] = list_api_args(args_str) + if not name in opts_map: + fatal("not found: " + name) + not_found += 1 +if not_found != 0: + fatal(not_found + " API calls not found") + +# Generating output header file +with open(OUTPUT, 'w') as f: + generate_prof_header(f, api_map, opts_map) + +# Successfull exit +sys.exit(0) diff --git a/include/hip/hcc_detail/hip_prof_str.h b/include/hip/hcc_detail/hip_prof_str.h deleted file mode 100644 index 0eb7636de8..0000000000 --- a/include/hip/hcc_detail/hip_prof_str.h +++ /dev/null @@ -1,2378 +0,0 @@ -// automatically generated sources -#ifndef _HIP_PROF_STR_H -#define _HIP_PROF_STR_H -#include -#include - -// Dummy API primitives -#define INIT_NONE_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipMemcpyToSymbolAsync_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipMemcpyFromSymbolAsync_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipDestroySurfaceObject_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipTexRefSetAddress_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipGetTextureObjectTextureDesc_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipBindTexture2D_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipTexRefSetAddressMode_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipCreateTextureObject_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipBindTextureToMipmappedArray_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipBindTextureToArray_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipTexRefSetFormat_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipMemcpyFromSymbol_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipMemcpyHtoH_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipGetTextureReference_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipDestroyTextureObject_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipTexRefSetArray_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipTexRefSetAddress2D_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipGetTextureObjectResourceViewDesc_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipUnbindTexture_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipTexRefSetFilterMode_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipCreateSurfaceObject_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipGetChannelDesc_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipHccGetAcceleratorView_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipExtModuleLaunchKernel_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipGetTextureObjectResourceDesc_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipMemcpyToSymbol_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipGetTextureAlignmentOffset_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipTexRefSetFlags_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipBindTexture_CB_ARGS_DATA(cb_data) {}; -#define INIT_hipHccGetAccelerator_CB_ARGS_DATA(cb_data) {}; - -// HIP API callbacks ID enumaration -enum hip_api_id_t { - HIP_API_ID_hipStreamCreateWithPriority = 0, - HIP_API_ID_hipMallocPitch = 1, - HIP_API_ID_hipMalloc = 2, - HIP_API_ID_hipEventRecord = 3, - HIP_API_ID_hipCtxSynchronize = 4, - HIP_API_ID_hipSetDevice = 5, - HIP_API_ID_hipCtxGetApiVersion = 6, - HIP_API_ID_hipSetupArgument = 7, - HIP_API_ID_hipMemcpyDtoD = 8, - HIP_API_ID_hipHostFree = 9, - HIP_API_ID_hipMemcpy2DToArray = 10, - HIP_API_ID_hipCtxGetCacheConfig = 11, - HIP_API_ID_hipStreamWaitEvent = 12, - HIP_API_ID_hipDeviceGetStreamPriorityRange = 13, - HIP_API_ID_hipModuleLoad = 14, - HIP_API_ID_hipDevicePrimaryCtxSetFlags = 15, - HIP_API_ID_hipMemcpyAsync = 16, - HIP_API_ID_hipMalloc3DArray = 17, - HIP_API_ID_hipStreamCreate = 18, - HIP_API_ID_hipCtxGetCurrent = 19, - HIP_API_ID_hipDevicePrimaryCtxGetState = 20, - HIP_API_ID_hipEventQuery = 21, - HIP_API_ID_hipEventCreate = 22, - HIP_API_ID_hipMemGetAddressRange = 23, - HIP_API_ID_hipArrayCreate = 24, - HIP_API_ID_hipStreamGetFlags = 25, - HIP_API_ID_hipMallocArray = 26, - HIP_API_ID_hipCtxGetSharedMemConfig = 27, - HIP_API_ID_hipMemPtrGetInfo = 28, - HIP_API_ID_hipCtxGetFlags = 29, - HIP_API_ID_hipStreamDestroy = 30, - HIP_API_ID_hipMemset3DAsync = 31, - HIP_API_ID_hipMemcpy3D = 32, - HIP_API_ID_hipInit = 33, - HIP_API_ID_hipMemcpyAtoH = 34, - HIP_API_ID_hipStreamGetPriority = 35, - HIP_API_ID_hipMemset2D = 36, - HIP_API_ID_hipMemset2DAsync = 37, - HIP_API_ID_hipDeviceCanAccessPeer = 38, - HIP_API_ID_hipDeviceEnablePeerAccess = 39, - HIP_API_ID_hipModuleUnload = 40, - HIP_API_ID_hipHostUnregister = 41, - HIP_API_ID_hipProfilerStop = 42, - HIP_API_ID_hipLaunchByPtr = 43, - HIP_API_ID_hipStreamSynchronize = 44, - HIP_API_ID_hipDeviceSetCacheConfig = 45, - HIP_API_ID_hipGetErrorName = 46, - HIP_API_ID_hipMemcpyHtoD = 47, - HIP_API_ID_hipMemcpyHtoA = 48, - HIP_API_ID_hipCtxCreate = 49, - HIP_API_ID_hipMemcpy2D = 50, - HIP_API_ID_hipIpcCloseMemHandle = 51, - HIP_API_ID_hipChooseDevice = 52, - HIP_API_ID_hipDeviceSetSharedMemConfig = 53, - HIP_API_ID_hipDeviceComputeCapability = 54, - HIP_API_ID_hipDeviceGet = 55, - HIP_API_ID_hipProfilerStart = 56, - HIP_API_ID_hipCtxSetCacheConfig = 57, - HIP_API_ID_hipFuncSetCacheConfig = 58, - HIP_API_ID_hipMemcpyPeerAsync = 59, - HIP_API_ID_hipEventElapsedTime = 60, - HIP_API_ID_hipDevicePrimaryCtxReset = 61, - HIP_API_ID_hipEventDestroy = 62, - HIP_API_ID_hipCtxPopCurrent = 63, - HIP_API_ID_hipHostGetFlags = 64, - HIP_API_ID_hipHostMalloc = 65, - HIP_API_ID_hipDriverGetVersion = 66, - HIP_API_ID_hipMemGetInfo = 67, - HIP_API_ID_hipDeviceReset = 68, - HIP_API_ID_hipMemset = 69, - HIP_API_ID_hipMemsetD8 = 70, - HIP_API_ID_hipHostRegister = 71, - HIP_API_ID_hipCtxSetSharedMemConfig = 72, - HIP_API_ID_hipArray3DCreate = 73, - HIP_API_ID_hipIpcOpenMemHandle = 74, - HIP_API_ID_hipGetLastError = 75, - HIP_API_ID_hipCtxDestroy = 76, - HIP_API_ID_hipDeviceGetSharedMemConfig = 77, - HIP_API_ID_hipSetDeviceFlags = 78, - HIP_API_ID_hipHccModuleLaunchKernel = 79, - HIP_API_ID_hipFree = 80, - HIP_API_ID_hipDeviceGetAttribute = 81, - HIP_API_ID_hipMemcpyDtoH = 82, - HIP_API_ID_hipCtxDisablePeerAccess = 83, - HIP_API_ID_hipDeviceGetByPCIBusId = 84, - HIP_API_ID_hipIpcGetMemHandle = 85, - HIP_API_ID_hipMemcpyHtoDAsync = 86, - HIP_API_ID_hipCtxGetDevice = 87, - HIP_API_ID_hipMemset3D = 88, - HIP_API_ID_hipModuleLoadData = 89, - HIP_API_ID_hipDeviceTotalMem = 90, - HIP_API_ID_hipCtxSetCurrent = 91, - HIP_API_ID_hipGetErrorString = 92, - HIP_API_ID_hipDevicePrimaryCtxRetain = 93, - HIP_API_ID_hipDeviceDisablePeerAccess = 94, - HIP_API_ID_hipStreamCreateWithFlags = 95, - HIP_API_ID_hipMemcpyFromArray = 96, - HIP_API_ID_hipMemcpy2DAsync = 97, - HIP_API_ID_hipEventCreateWithFlags = 98, - HIP_API_ID_hipStreamQuery = 99, - HIP_API_ID_hipDeviceGetPCIBusId = 100, - HIP_API_ID_hipMemcpy = 101, - HIP_API_ID_hipPeekAtLastError = 102, - HIP_API_ID_hipStreamAddCallback = 103, - HIP_API_ID_hipMemcpyToArray = 104, - HIP_API_ID_hipMemsetD32 = 105, - HIP_API_ID_hipDeviceSynchronize = 106, - HIP_API_ID_hipDeviceGetCacheConfig = 107, - HIP_API_ID_hipMalloc3D = 108, - HIP_API_ID_hipPointerGetAttributes = 109, - HIP_API_ID_hipMemsetAsync = 110, - HIP_API_ID_hipDeviceGetName = 111, - HIP_API_ID_hipCtxPushCurrent = 112, - HIP_API_ID_hipMemcpyPeer = 113, - HIP_API_ID_hipEventSynchronize = 114, - HIP_API_ID_hipMemcpyDtoDAsync = 115, - HIP_API_ID_hipCtxEnablePeerAccess = 116, - HIP_API_ID_hipMemcpyDtoHAsync = 117, - HIP_API_ID_hipModuleLaunchKernel = 118, - HIP_API_ID_hipModuleGetTexRef = 119, - HIP_API_ID_hipDeviceGetLimit = 120, - HIP_API_ID_hipModuleLoadDataEx = 121, - HIP_API_ID_hipRuntimeGetVersion = 122, - HIP_API_ID_hipGetDeviceProperties = 123, - HIP_API_ID_hipFreeArray = 124, - HIP_API_ID_hipDevicePrimaryCtxRelease = 125, - HIP_API_ID_hipHostGetDevicePointer = 126, - HIP_API_ID_hipMemcpyParam2D = 127, - HIP_API_ID_hipModuleGetFunction = 128, - HIP_API_ID_hipMemsetD32Async = 129, - HIP_API_ID_hipGetDevice = 130, - HIP_API_ID_hipGetDeviceCount = 131, - HIP_API_ID_NUMBER = 132, - HIP_API_ID_ANY = 133, - - HIP_API_ID_NONE = HIP_API_ID_NUMBER, - HIP_API_ID_hipMemcpyToSymbolAsync = HIP_API_ID_NUMBER, - HIP_API_ID_hipMemcpyFromSymbolAsync = HIP_API_ID_NUMBER, - HIP_API_ID_hipDestroySurfaceObject = HIP_API_ID_NUMBER, - HIP_API_ID_hipTexRefSetAddress = HIP_API_ID_NUMBER, - HIP_API_ID_hipGetTextureObjectTextureDesc = HIP_API_ID_NUMBER, - HIP_API_ID_hipBindTexture2D = HIP_API_ID_NUMBER, - HIP_API_ID_hipTexRefSetAddressMode = HIP_API_ID_NUMBER, - HIP_API_ID_hipCreateTextureObject = HIP_API_ID_NUMBER, - HIP_API_ID_hipBindTextureToMipmappedArray = HIP_API_ID_NUMBER, - HIP_API_ID_hipBindTextureToArray = HIP_API_ID_NUMBER, - HIP_API_ID_hipTexRefSetFormat = HIP_API_ID_NUMBER, - HIP_API_ID_hipMemcpyFromSymbol = HIP_API_ID_NUMBER, - HIP_API_ID_hipMemcpyHtoH = HIP_API_ID_NUMBER, - HIP_API_ID_hipGetTextureReference = HIP_API_ID_NUMBER, - HIP_API_ID_hipDestroyTextureObject = HIP_API_ID_NUMBER, - HIP_API_ID_hipTexRefSetArray = HIP_API_ID_NUMBER, - HIP_API_ID_hipTexRefSetAddress2D = HIP_API_ID_NUMBER, - HIP_API_ID_hipGetTextureObjectResourceViewDesc = HIP_API_ID_NUMBER, - HIP_API_ID_hipUnbindTexture = HIP_API_ID_NUMBER, - HIP_API_ID_hipTexRefSetFilterMode = HIP_API_ID_NUMBER, - HIP_API_ID_hipCreateSurfaceObject = HIP_API_ID_NUMBER, - HIP_API_ID_hipGetChannelDesc = HIP_API_ID_NUMBER, - HIP_API_ID_hipHccGetAcceleratorView = HIP_API_ID_NUMBER, - HIP_API_ID_hipExtModuleLaunchKernel = HIP_API_ID_NUMBER, - HIP_API_ID_hipGetTextureObjectResourceDesc = HIP_API_ID_NUMBER, - HIP_API_ID_hipMemcpyToSymbol = HIP_API_ID_NUMBER, - HIP_API_ID_hipGetTextureAlignmentOffset = HIP_API_ID_NUMBER, - HIP_API_ID_hipTexRefSetFlags = HIP_API_ID_NUMBER, - HIP_API_ID_hipBindTexture = HIP_API_ID_NUMBER, - HIP_API_ID_hipHccGetAccelerator = HIP_API_ID_NUMBER, -}; - -// Return HIP API string -static const char* hip_api_name(const uint32_t& id) { - switch(id) { - case HIP_API_ID_hipStreamCreateWithPriority: return "hipStreamCreateWithPriority"; - case HIP_API_ID_hipMallocPitch: return "hipMallocPitch"; - case HIP_API_ID_hipMalloc: return "hipMalloc"; - case HIP_API_ID_hipEventRecord: return "hipEventRecord"; - case HIP_API_ID_hipCtxSynchronize: return "hipCtxSynchronize"; - case HIP_API_ID_hipSetDevice: return "hipSetDevice"; - case HIP_API_ID_hipCtxGetApiVersion: return "hipCtxGetApiVersion"; - case HIP_API_ID_hipSetupArgument: return "hipSetupArgument"; - case HIP_API_ID_hipMemcpyDtoD: return "hipMemcpyDtoD"; - case HIP_API_ID_hipHostFree: return "hipHostFree"; - case HIP_API_ID_hipMemcpy2DToArray: return "hipMemcpy2DToArray"; - case HIP_API_ID_hipCtxGetCacheConfig: return "hipCtxGetCacheConfig"; - case HIP_API_ID_hipStreamWaitEvent: return "hipStreamWaitEvent"; - case HIP_API_ID_hipDeviceGetStreamPriorityRange: return "hipDeviceGetStreamPriorityRange"; - case HIP_API_ID_hipModuleLoad: return "hipModuleLoad"; - case HIP_API_ID_hipDevicePrimaryCtxSetFlags: return "hipDevicePrimaryCtxSetFlags"; - case HIP_API_ID_hipMemcpyAsync: return "hipMemcpyAsync"; - case HIP_API_ID_hipMalloc3DArray: return "hipMalloc3DArray"; - case HIP_API_ID_hipStreamCreate: return "hipStreamCreate"; - case HIP_API_ID_hipCtxGetCurrent: return "hipCtxGetCurrent"; - case HIP_API_ID_hipDevicePrimaryCtxGetState: return "hipDevicePrimaryCtxGetState"; - case HIP_API_ID_hipEventQuery: return "hipEventQuery"; - case HIP_API_ID_hipEventCreate: return "hipEventCreate"; - case HIP_API_ID_hipMemGetAddressRange: return "hipMemGetAddressRange"; - case HIP_API_ID_hipArrayCreate: return "hipArrayCreate"; - case HIP_API_ID_hipStreamGetFlags: return "hipStreamGetFlags"; - case HIP_API_ID_hipMallocArray: return "hipMallocArray"; - case HIP_API_ID_hipCtxGetSharedMemConfig: return "hipCtxGetSharedMemConfig"; - case HIP_API_ID_hipMemPtrGetInfo: return "hipMemPtrGetInfo"; - case HIP_API_ID_hipCtxGetFlags: return "hipCtxGetFlags"; - case HIP_API_ID_hipStreamDestroy: return "hipStreamDestroy"; - case HIP_API_ID_hipMemset3DAsync: return "hipMemset3DAsync"; - case HIP_API_ID_hipMemcpy3D: return "hipMemcpy3D"; - case HIP_API_ID_hipInit: return "hipInit"; - case HIP_API_ID_hipMemcpyAtoH: return "hipMemcpyAtoH"; - case HIP_API_ID_hipStreamGetPriority: return "hipStreamGetPriority"; - case HIP_API_ID_hipMemset2D: return "hipMemset2D"; - case HIP_API_ID_hipMemset2DAsync: return "hipMemset2DAsync"; - case HIP_API_ID_hipDeviceCanAccessPeer: return "hipDeviceCanAccessPeer"; - case HIP_API_ID_hipDeviceEnablePeerAccess: return "hipDeviceEnablePeerAccess"; - case HIP_API_ID_hipModuleUnload: return "hipModuleUnload"; - case HIP_API_ID_hipHostUnregister: return "hipHostUnregister"; - case HIP_API_ID_hipProfilerStop: return "hipProfilerStop"; - case HIP_API_ID_hipLaunchByPtr: return "hipLaunchByPtr"; - case HIP_API_ID_hipStreamSynchronize: return "hipStreamSynchronize"; - case HIP_API_ID_hipDeviceSetCacheConfig: return "hipDeviceSetCacheConfig"; - case HIP_API_ID_hipGetErrorName: return "hipGetErrorName"; - case HIP_API_ID_hipMemcpyHtoD: return "hipMemcpyHtoD"; - case HIP_API_ID_hipMemcpyHtoA: return "hipMemcpyHtoA"; - case HIP_API_ID_hipCtxCreate: return "hipCtxCreate"; - case HIP_API_ID_hipMemcpy2D: return "hipMemcpy2D"; - case HIP_API_ID_hipIpcCloseMemHandle: return "hipIpcCloseMemHandle"; - case HIP_API_ID_hipChooseDevice: return "hipChooseDevice"; - case HIP_API_ID_hipDeviceSetSharedMemConfig: return "hipDeviceSetSharedMemConfig"; - case HIP_API_ID_hipDeviceComputeCapability: return "hipDeviceComputeCapability"; - case HIP_API_ID_hipDeviceGet: return "hipDeviceGet"; - case HIP_API_ID_hipProfilerStart: return "hipProfilerStart"; - case HIP_API_ID_hipCtxSetCacheConfig: return "hipCtxSetCacheConfig"; - case HIP_API_ID_hipFuncSetCacheConfig: return "hipFuncSetCacheConfig"; - case HIP_API_ID_hipMemcpyPeerAsync: return "hipMemcpyPeerAsync"; - case HIP_API_ID_hipEventElapsedTime: return "hipEventElapsedTime"; - case HIP_API_ID_hipDevicePrimaryCtxReset: return "hipDevicePrimaryCtxReset"; - case HIP_API_ID_hipEventDestroy: return "hipEventDestroy"; - case HIP_API_ID_hipCtxPopCurrent: return "hipCtxPopCurrent"; - case HIP_API_ID_hipHostGetFlags: return "hipHostGetFlags"; - case HIP_API_ID_hipHostMalloc: return "hipHostMalloc"; - case HIP_API_ID_hipDriverGetVersion: return "hipDriverGetVersion"; - case HIP_API_ID_hipMemGetInfo: return "hipMemGetInfo"; - case HIP_API_ID_hipDeviceReset: return "hipDeviceReset"; - case HIP_API_ID_hipMemset: return "hipMemset"; - case HIP_API_ID_hipMemsetD8: return "hipMemsetD8"; - case HIP_API_ID_hipHostRegister: return "hipHostRegister"; - case HIP_API_ID_hipCtxSetSharedMemConfig: return "hipCtxSetSharedMemConfig"; - case HIP_API_ID_hipArray3DCreate: return "hipArray3DCreate"; - case HIP_API_ID_hipIpcOpenMemHandle: return "hipIpcOpenMemHandle"; - case HIP_API_ID_hipGetLastError: return "hipGetLastError"; - case HIP_API_ID_hipCtxDestroy: return "hipCtxDestroy"; - case HIP_API_ID_hipDeviceGetSharedMemConfig: return "hipDeviceGetSharedMemConfig"; - case HIP_API_ID_hipSetDeviceFlags: return "hipSetDeviceFlags"; - case HIP_API_ID_hipHccModuleLaunchKernel: return "hipHccModuleLaunchKernel"; - case HIP_API_ID_hipFree: return "hipFree"; - case HIP_API_ID_hipDeviceGetAttribute: return "hipDeviceGetAttribute"; - case HIP_API_ID_hipMemcpyDtoH: return "hipMemcpyDtoH"; - case HIP_API_ID_hipCtxDisablePeerAccess: return "hipCtxDisablePeerAccess"; - case HIP_API_ID_hipDeviceGetByPCIBusId: return "hipDeviceGetByPCIBusId"; - case HIP_API_ID_hipIpcGetMemHandle: return "hipIpcGetMemHandle"; - case HIP_API_ID_hipMemcpyHtoDAsync: return "hipMemcpyHtoDAsync"; - case HIP_API_ID_hipCtxGetDevice: return "hipCtxGetDevice"; - case HIP_API_ID_hipMemset3D: return "hipMemset3D"; - case HIP_API_ID_hipModuleLoadData: return "hipModuleLoadData"; - case HIP_API_ID_hipDeviceTotalMem: return "hipDeviceTotalMem"; - case HIP_API_ID_hipCtxSetCurrent: return "hipCtxSetCurrent"; - case HIP_API_ID_hipGetErrorString: return "hipGetErrorString"; - case HIP_API_ID_hipDevicePrimaryCtxRetain: return "hipDevicePrimaryCtxRetain"; - case HIP_API_ID_hipDeviceDisablePeerAccess: return "hipDeviceDisablePeerAccess"; - case HIP_API_ID_hipStreamCreateWithFlags: return "hipStreamCreateWithFlags"; - case HIP_API_ID_hipMemcpyFromArray: return "hipMemcpyFromArray"; - case HIP_API_ID_hipMemcpy2DAsync: return "hipMemcpy2DAsync"; - case HIP_API_ID_hipEventCreateWithFlags: return "hipEventCreateWithFlags"; - case HIP_API_ID_hipStreamQuery: return "hipStreamQuery"; - case HIP_API_ID_hipDeviceGetPCIBusId: return "hipDeviceGetPCIBusId"; - case HIP_API_ID_hipMemcpy: return "hipMemcpy"; - case HIP_API_ID_hipPeekAtLastError: return "hipPeekAtLastError"; - case HIP_API_ID_hipStreamAddCallback: return "hipStreamAddCallback"; - case HIP_API_ID_hipMemcpyToArray: return "hipMemcpyToArray"; - case HIP_API_ID_hipMemsetD32: return "hipMemsetD32"; - case HIP_API_ID_hipDeviceSynchronize: return "hipDeviceSynchronize"; - case HIP_API_ID_hipDeviceGetCacheConfig: return "hipDeviceGetCacheConfig"; - case HIP_API_ID_hipMalloc3D: return "hipMalloc3D"; - case HIP_API_ID_hipPointerGetAttributes: return "hipPointerGetAttributes"; - case HIP_API_ID_hipMemsetAsync: return "hipMemsetAsync"; - case HIP_API_ID_hipDeviceGetName: return "hipDeviceGetName"; - case HIP_API_ID_hipCtxPushCurrent: return "hipCtxPushCurrent"; - case HIP_API_ID_hipMemcpyPeer: return "hipMemcpyPeer"; - case HIP_API_ID_hipEventSynchronize: return "hipEventSynchronize"; - case HIP_API_ID_hipMemcpyDtoDAsync: return "hipMemcpyDtoDAsync"; - case HIP_API_ID_hipCtxEnablePeerAccess: return "hipCtxEnablePeerAccess"; - case HIP_API_ID_hipMemcpyDtoHAsync: return "hipMemcpyDtoHAsync"; - case HIP_API_ID_hipModuleLaunchKernel: return "hipModuleLaunchKernel"; - case HIP_API_ID_hipModuleGetTexRef: return "hipModuleGetTexRef"; - case HIP_API_ID_hipDeviceGetLimit: return "hipDeviceGetLimit"; - case HIP_API_ID_hipModuleLoadDataEx: return "hipModuleLoadDataEx"; - case HIP_API_ID_hipRuntimeGetVersion: return "hipRuntimeGetVersion"; - case HIP_API_ID_hipGetDeviceProperties: return "hipGetDeviceProperties"; - case HIP_API_ID_hipFreeArray: return "hipFreeArray"; - case HIP_API_ID_hipDevicePrimaryCtxRelease: return "hipDevicePrimaryCtxRelease"; - case HIP_API_ID_hipHostGetDevicePointer: return "hipHostGetDevicePointer"; - case HIP_API_ID_hipMemcpyParam2D: return "hipMemcpyParam2D"; - case HIP_API_ID_hipModuleGetFunction: return "hipModuleGetFunction"; - case HIP_API_ID_hipMemsetD32Async: return "hipMemsetD32Async"; - case HIP_API_ID_hipGetDevice: return "hipGetDevice"; - case HIP_API_ID_hipGetDeviceCount: return "hipGetDeviceCount"; - }; - return "unknown"; -}; - -// HIP API callbacks data structure -struct hip_api_data_t { - uint64_t correlation_id; - uint32_t phase; - union { - struct { - hipStream_t* stream; - unsigned int flags; - int priority; - } hipStreamCreateWithPriority; - struct { - void** ptr; - size_t* pitch; - size_t width; - size_t height; - } hipMallocPitch; - struct { - void** ptr; - size_t size; - } hipMalloc; - struct { - hipEvent_t event; - hipStream_t stream; - } hipEventRecord; - struct { - int deviceId; - } hipSetDevice; - struct { - hipCtx_t ctx; - int* apiVersion; - } hipCtxGetApiVersion; - struct { - const void* arg; - size_t size; - size_t offset; - } hipSetupArgument; - struct { - hipDeviceptr_t dst; - hipDeviceptr_t src; - size_t sizeBytes; - } hipMemcpyDtoD; - struct { - void* ptr; - } hipHostFree; - struct { - hipArray* dst; - size_t wOffset; - size_t hOffset; - const void* src; - size_t spitch; - size_t width; - size_t height; - hipMemcpyKind kind; - } hipMemcpy2DToArray; - struct { - hipFuncCache_t* cacheConfig; - } hipCtxGetCacheConfig; - struct { - hipStream_t stream; - hipEvent_t event; - unsigned int flags; - } hipStreamWaitEvent; - struct { - int* leastPriority; - int* greatestPriority; - } hipDeviceGetStreamPriorityRange; - struct { - hipModule_t* module; - const char* fname; - } hipModuleLoad; - struct { - hipDevice_t dev; - unsigned int flags; - } hipDevicePrimaryCtxSetFlags; - struct { - void* dst; - const void* src; - size_t sizeBytes; - hipMemcpyKind kind; - hipStream_t stream; - } hipMemcpyAsync; - struct { - hipArray** array; - const hipChannelFormatDesc* desc; - hipExtent extent; - unsigned int flags; - } hipMalloc3DArray; - struct { - hipStream_t* stream; - } hipStreamCreate; - struct { - hipCtx_t* ctx; - } hipCtxGetCurrent; - struct { - hipDevice_t dev; - unsigned int* flags; - int* active; - } hipDevicePrimaryCtxGetState; - struct { - hipEvent_t event; - } hipEventQuery; - struct { - hipEvent_t* event; - } hipEventCreate; - struct { - hipDeviceptr_t* pbase; - size_t* psize; - hipDeviceptr_t dptr; - } hipMemGetAddressRange; - struct { - hipArray** pHandle; - const HIP_ARRAY_DESCRIPTOR* pAllocateArray; - } hipArrayCreate; - struct { - hipStream_t stream; - unsigned int* flags; - } hipStreamGetFlags; - struct { - hipArray** array; - const hipChannelFormatDesc* desc; - size_t width; - size_t height; - unsigned int flags; - } hipMallocArray; - struct { - hipSharedMemConfig* pConfig; - } hipCtxGetSharedMemConfig; - struct { - void* ptr; - size_t* size; - } hipMemPtrGetInfo; - struct { - unsigned int* flags; - } hipCtxGetFlags; - struct { - hipStream_t stream; - } hipStreamDestroy; - struct { - hipPitchedPtr pitchedDevPtr; - int value; - hipExtent extent; - hipStream_t stream; - } hipMemset3DAsync; - struct { - const hipMemcpy3DParms* p; - } hipMemcpy3D; - struct { - unsigned int flags; - } hipInit; - struct { - void* dst; - hipArray* srcArray; - size_t srcOffset; - size_t count; - } hipMemcpyAtoH; - struct { - hipStream_t stream; - int* priority; - } hipStreamGetPriority; - struct { - void* dst; - size_t pitch; - int value; - size_t width; - size_t height; - } hipMemset2D; - struct { - void* dst; - size_t pitch; - int value; - size_t width; - size_t height; - hipStream_t stream; - } hipMemset2DAsync; - struct { - int* canAccessPeer; - int deviceId; - int peerDeviceId; - } hipDeviceCanAccessPeer; - struct { - int peerDeviceId; - unsigned int flags; - } hipDeviceEnablePeerAccess; - struct { - hipModule_t module; - } hipModuleUnload; - struct { - void* hostPtr; - } hipHostUnregister; - struct { - const void* func; - } hipLaunchByPtr; - struct { - hipStream_t stream; - } hipStreamSynchronize; - struct { - hipFuncCache_t cacheConfig; - } hipDeviceSetCacheConfig; - struct { - hipError_t hip_error; - } hipGetErrorName; - struct { - hipDeviceptr_t dst; - void* src; - size_t sizeBytes; - } hipMemcpyHtoD; - struct { - hipArray* dstArray; - size_t dstOffset; - const void* srcHost; - size_t count; - } hipMemcpyHtoA; - struct { - hipCtx_t* ctx; - unsigned int flags; - hipDevice_t device; - } hipCtxCreate; - struct { - void* dst; - size_t dpitch; - const void* src; - size_t spitch; - size_t width; - size_t height; - hipMemcpyKind kind; - } hipMemcpy2D; - struct { - void* devPtr; - } hipIpcCloseMemHandle; - struct { - int* device; - const hipDeviceProp_t* prop; - } hipChooseDevice; - struct { - hipSharedMemConfig config; - } hipDeviceSetSharedMemConfig; - struct { - int* major; - int* minor; - hipDevice_t device; - } hipDeviceComputeCapability; - struct { - hipDevice_t* device; - int ordinal; - } hipDeviceGet; - struct { - hipFuncCache_t cacheConfig; - } hipCtxSetCacheConfig; - struct { - const void* func; - hipFuncCache_t config; - } hipFuncSetCacheConfig; - struct { - void* dst; - int dstDeviceId; - const void* src; - int srcDevice; - size_t sizeBytes; - hipStream_t stream; - } hipMemcpyPeerAsync; - struct { - float* ms; - hipEvent_t start; - hipEvent_t stop; - } hipEventElapsedTime; - struct { - hipDevice_t dev; - } hipDevicePrimaryCtxReset; - struct { - hipEvent_t event; - } hipEventDestroy; - struct { - hipCtx_t* ctx; - } hipCtxPopCurrent; - struct { - unsigned int* flagsPtr; - void* hostPtr; - } hipHostGetFlags; - struct { - void** ptr; - size_t size; - unsigned int flags; - } hipHostMalloc; - struct { - int* driverVersion; - } hipDriverGetVersion; - struct { - size_t* free; - size_t* total; - } hipMemGetInfo; - struct { - void* dst; - int value; - size_t sizeBytes; - } hipMemset; - struct { - hipDeviceptr_t dest; - unsigned char value; - size_t sizeBytes; - } hipMemsetD8; - struct { - void* hostPtr; - size_t sizeBytes; - unsigned int flags; - } hipHostRegister; - struct { - hipSharedMemConfig config; - } hipCtxSetSharedMemConfig; - struct { - hipArray** array; - const HIP_ARRAY_DESCRIPTOR* pAllocateArray; - } hipArray3DCreate; - struct { - void** devPtr; - hipIpcMemHandle_t handle; - unsigned int flags; - } hipIpcOpenMemHandle; - struct { - hipCtx_t ctx; - } hipCtxDestroy; - struct { - hipSharedMemConfig* pConfig; - } hipDeviceGetSharedMemConfig; - struct { - unsigned int flags; - } hipSetDeviceFlags; - struct { - hipFunction_t f; - unsigned int globalWorkSizeX; - unsigned int globalWorkSizeY; - unsigned int globalWorkSizeZ; - unsigned int localWorkSizeX; - unsigned int localWorkSizeY; - unsigned int localWorkSizeZ; - size_t sharedMemBytes; - hipStream_t hStream; - void** kernelParams; - void** extra; - hipEvent_t startEvent; - hipEvent_t stopEvent; - } hipHccModuleLaunchKernel; - struct { - void* ptr; - } hipFree; - struct { - int* pi; - hipDeviceAttribute_t attr; - int deviceId; - } hipDeviceGetAttribute; - struct { - void* dst; - hipDeviceptr_t src; - size_t sizeBytes; - } hipMemcpyDtoH; - struct { - hipCtx_t peerCtx; - } hipCtxDisablePeerAccess; - struct { - int* device; - const char* pciBusId; - } hipDeviceGetByPCIBusId; - struct { - hipIpcMemHandle_t* handle; - void* devPtr; - } hipIpcGetMemHandle; - struct { - hipDeviceptr_t dst; - void* src; - size_t sizeBytes; - hipStream_t stream; - } hipMemcpyHtoDAsync; - struct { - hipDevice_t* device; - } hipCtxGetDevice; - struct { - hipPitchedPtr pitchedDevPtr; - int value; - hipExtent extent; - } hipMemset3D; - struct { - hipModule_t* module; - const void* image; - } hipModuleLoadData; - struct { - size_t* bytes; - hipDevice_t device; - } hipDeviceTotalMem; - struct { - hipCtx_t ctx; - } hipCtxSetCurrent; - struct { - hipError_t hipError; - } hipGetErrorString; - struct { - hipCtx_t* pctx; - hipDevice_t dev; - } hipDevicePrimaryCtxRetain; - struct { - int peerDeviceId; - } hipDeviceDisablePeerAccess; - struct { - hipStream_t* stream; - unsigned int flags; - } hipStreamCreateWithFlags; - struct { - void* dst; - hipArray_const_t srcArray; - size_t wOffset; - size_t hOffset; - size_t count; - hipMemcpyKind kind; - } hipMemcpyFromArray; - struct { - void* dst; - size_t dpitch; - const void* src; - size_t spitch; - size_t width; - size_t height; - hipMemcpyKind kind; - hipStream_t stream; - } hipMemcpy2DAsync; - struct { - hipEvent_t* event; - unsigned int flags; - } hipEventCreateWithFlags; - struct { - hipStream_t stream; - } hipStreamQuery; - struct { - char* pciBusId; - int len; - int device; - } hipDeviceGetPCIBusId; - struct { - void* dst; - const void* src; - size_t sizeBytes; - hipMemcpyKind kind; - } hipMemcpy; - struct { - hipStream_t stream; - hipStreamCallback_t callback; - void* userData; - unsigned int flags; - } hipStreamAddCallback; - struct { - hipArray* dst; - size_t wOffset; - size_t hOffset; - const void* src; - size_t count; - hipMemcpyKind kind; - } hipMemcpyToArray; - struct { - hipDeviceptr_t dest; - int value; - size_t count; - } hipMemsetD32; - struct { - hipFuncCache_t* cacheConfig; - } hipDeviceGetCacheConfig; - struct { - hipPitchedPtr* pitchedDevPtr; - hipExtent extent; - } hipMalloc3D; - struct { - hipPointerAttribute_t* attributes; - const void* ptr; - } hipPointerGetAttributes; - struct { - void* dst; - int value; - size_t sizeBytes; - hipStream_t stream; - } hipMemsetAsync; - struct { - char* name; - int len; - hipDevice_t device; - } hipDeviceGetName; - struct { - hipCtx_t ctx; - } hipCtxPushCurrent; - struct { - void* dst; - int dstDeviceId; - const void* src; - int srcDeviceId; - size_t sizeBytes; - } hipMemcpyPeer; - struct { - hipEvent_t event; - } hipEventSynchronize; - struct { - hipDeviceptr_t dst; - hipDeviceptr_t src; - size_t sizeBytes; - hipStream_t stream; - } hipMemcpyDtoDAsync; - struct { - hipCtx_t peerCtx; - unsigned int flags; - } hipCtxEnablePeerAccess; - struct { - void* dst; - hipDeviceptr_t src; - size_t sizeBytes; - hipStream_t stream; - } hipMemcpyDtoHAsync; - struct { - hipFunction_t f; - unsigned int gridDimX; - unsigned int gridDimY; - unsigned int gridDimZ; - unsigned int blockDimX; - unsigned int blockDimY; - unsigned int blockDimZ; - unsigned int sharedMemBytes; - hipStream_t stream; - void** kernelParams; - void** extra; - } hipModuleLaunchKernel; - struct { - textureReference** texRef; - hipModule_t hmod; - const char* name; - } hipModuleGetTexRef; - struct { - size_t* pValue; - hipLimit_t limit; - } hipDeviceGetLimit; - struct { - hipModule_t* module; - const void* image; - unsigned int numOptions; - hipJitOption* options; - void** optionValues; - } hipModuleLoadDataEx; - struct { - int* runtimeVersion; - } hipRuntimeGetVersion; - struct { - hipDeviceProp_t* prop; - int deviceId; - } hipGetDeviceProperties; - struct { - hipArray* array; - } hipFreeArray; - struct { - hipDevice_t dev; - } hipDevicePrimaryCtxRelease; - struct { - void** devPtr; - void* hstPtr; - unsigned int flags; - } hipHostGetDevicePointer; - struct { - const hip_Memcpy2D* pCopy; - } hipMemcpyParam2D; - struct { - hipFunction_t* function; - hipModule_t module; - const char* kname; - } hipModuleGetFunction; - struct { - hipDeviceptr_t dst; - int value; - size_t count; - hipStream_t stream; - } hipMemsetD32Async; - struct { - int* deviceId; - } hipGetDevice; - struct { - int* count; - } hipGetDeviceCount; - } args; -}; - -// HIP API callbacks args data filling macros -#define INIT_hipStreamCreateWithPriority_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamCreateWithPriority.stream = stream; \ - cb_data.args.hipStreamCreateWithPriority.flags = flags; \ - cb_data.args.hipStreamCreateWithPriority.priority = priority; \ -}; -#define INIT_hipMallocPitch_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMallocPitch.ptr = ptr; \ - cb_data.args.hipMallocPitch.pitch = pitch; \ - cb_data.args.hipMallocPitch.width = width; \ - cb_data.args.hipMallocPitch.height = height; \ -}; -#define INIT_hipMalloc_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMalloc.ptr = ptr; \ - cb_data.args.hipMalloc.size = sizeBytes; \ -}; -#define INIT_hipEventRecord_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipEventRecord.event = event; \ - cb_data.args.hipEventRecord.stream = stream; \ -}; -#define INIT_hipCtxSynchronize_CB_ARGS_DATA(cb_data) { \ -}; -#define INIT_hipSetDevice_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipSetDevice.deviceId = deviceId; \ -}; -#define INIT_hipCtxGetApiVersion_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxGetApiVersion.ctx = ctx; \ - cb_data.args.hipCtxGetApiVersion.apiVersion = apiVersion; \ -}; -#define INIT_hipSetupArgument_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipSetupArgument.arg = arg; \ - cb_data.args.hipSetupArgument.size = size; \ - cb_data.args.hipSetupArgument.offset = offset; \ -}; -#define INIT_hipMemcpyDtoD_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyDtoD.dst = dst; \ - cb_data.args.hipMemcpyDtoD.src = src; \ - cb_data.args.hipMemcpyDtoD.sizeBytes = sizeBytes; \ -}; -#define INIT_hipHostFree_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipHostFree.ptr = ptr; \ -}; -#define INIT_hipMemcpy2DToArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpy2DToArray.dst = dst; \ - cb_data.args.hipMemcpy2DToArray.wOffset = wOffset; \ - cb_data.args.hipMemcpy2DToArray.hOffset = hOffset; \ - cb_data.args.hipMemcpy2DToArray.src = src; \ - cb_data.args.hipMemcpy2DToArray.spitch = spitch; \ - cb_data.args.hipMemcpy2DToArray.width = width; \ - cb_data.args.hipMemcpy2DToArray.height = height; \ - cb_data.args.hipMemcpy2DToArray.kind = kind; \ -}; -#define INIT_hipCtxGetCacheConfig_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxGetCacheConfig.cacheConfig = cacheConfig; \ -}; -#define INIT_hipStreamWaitEvent_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamWaitEvent.stream = stream; \ - cb_data.args.hipStreamWaitEvent.event = event; \ - cb_data.args.hipStreamWaitEvent.flags = flags; \ -}; -#define INIT_hipDeviceGetStreamPriorityRange_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceGetStreamPriorityRange.leastPriority = leastPriority; \ - cb_data.args.hipDeviceGetStreamPriorityRange.greatestPriority = greatestPriority; \ -}; -#define INIT_hipModuleLoad_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipModuleLoad.module = module; \ - cb_data.args.hipModuleLoad.fname = fname; \ -}; -#define INIT_hipDevicePrimaryCtxSetFlags_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDevicePrimaryCtxSetFlags.dev = dev; \ - cb_data.args.hipDevicePrimaryCtxSetFlags.flags = flags; \ -}; -#define INIT_hipMemcpyAsync_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyAsync.dst = dst; \ - cb_data.args.hipMemcpyAsync.src = src; \ - cb_data.args.hipMemcpyAsync.sizeBytes = sizeBytes; \ - cb_data.args.hipMemcpyAsync.kind = kind; \ - cb_data.args.hipMemcpyAsync.stream = stream; \ -}; -#define INIT_hipMalloc3DArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMalloc3DArray.array = array; \ - cb_data.args.hipMalloc3DArray.desc = desc; \ - cb_data.args.hipMalloc3DArray.extent = extent; \ - cb_data.args.hipMalloc3DArray.flags = flags; \ -}; -#define INIT_hipStreamCreate_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamCreate.stream = stream; \ -}; -#define INIT_hipCtxGetCurrent_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxGetCurrent.ctx = ctx; \ -}; -#define INIT_hipDevicePrimaryCtxGetState_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDevicePrimaryCtxGetState.dev = dev; \ - cb_data.args.hipDevicePrimaryCtxGetState.flags = flags; \ - cb_data.args.hipDevicePrimaryCtxGetState.active = active; \ -}; -#define INIT_hipEventQuery_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipEventQuery.event = event; \ -}; -#define INIT_hipEventCreate_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipEventCreate.event = event; \ -}; -#define INIT_hipMemGetAddressRange_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemGetAddressRange.pbase = pbase; \ - cb_data.args.hipMemGetAddressRange.psize = psize; \ - cb_data.args.hipMemGetAddressRange.dptr = dptr; \ -}; -#define INIT_hipArrayCreate_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipArrayCreate.pHandle = array; \ - cb_data.args.hipArrayCreate.pAllocateArray = pAllocateArray; \ -}; -#define INIT_hipStreamGetFlags_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamGetFlags.stream = stream; \ - cb_data.args.hipStreamGetFlags.flags = flags; \ -}; -#define INIT_hipMallocArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMallocArray.array = array; \ - cb_data.args.hipMallocArray.desc = desc; \ - cb_data.args.hipMallocArray.width = width; \ - cb_data.args.hipMallocArray.height = height; \ - cb_data.args.hipMallocArray.flags = flags; \ -}; -#define INIT_hipCtxGetSharedMemConfig_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxGetSharedMemConfig.pConfig = pConfig; \ -}; -#define INIT_hipMemPtrGetInfo_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemPtrGetInfo.ptr = ptr; \ - cb_data.args.hipMemPtrGetInfo.size = size; \ -}; -#define INIT_hipCtxGetFlags_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxGetFlags.flags = flags; \ -}; -#define INIT_hipStreamDestroy_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamDestroy.stream = stream; \ -}; -#define INIT_hipMemset3DAsync_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemset3DAsync.pitchedDevPtr = pitchedDevPtr; \ - cb_data.args.hipMemset3DAsync.value = value; \ - cb_data.args.hipMemset3DAsync.extent = extent; \ - cb_data.args.hipMemset3DAsync.stream = stream; \ -}; -#define INIT_hipMemcpy3D_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpy3D.p = p; \ -}; -#define INIT_hipInit_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipInit.flags = flags; \ -}; -#define INIT_hipMemcpyAtoH_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyAtoH.dst = dst; \ - cb_data.args.hipMemcpyAtoH.srcArray = srcArray; \ - cb_data.args.hipMemcpyAtoH.srcOffset = srcOffset; \ - cb_data.args.hipMemcpyAtoH.count = count; \ -}; -#define INIT_hipStreamGetPriority_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamGetPriority.stream = stream; \ - cb_data.args.hipStreamGetPriority.priority = priority; \ -}; -#define INIT_hipMemset2D_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemset2D.dst = dst; \ - cb_data.args.hipMemset2D.pitch = pitch; \ - cb_data.args.hipMemset2D.value = value; \ - cb_data.args.hipMemset2D.width = width; \ - cb_data.args.hipMemset2D.height = height; \ -}; -#define INIT_hipMemset2DAsync_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemset2DAsync.dst = dst; \ - cb_data.args.hipMemset2DAsync.pitch = pitch; \ - cb_data.args.hipMemset2DAsync.value = value; \ - cb_data.args.hipMemset2DAsync.width = width; \ - cb_data.args.hipMemset2DAsync.height = height; \ - cb_data.args.hipMemset2DAsync.stream = stream; \ -}; -#define INIT_hipDeviceCanAccessPeer_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceCanAccessPeer.canAccessPeer = canAccessPeer; \ - cb_data.args.hipDeviceCanAccessPeer.deviceId = deviceId; \ - cb_data.args.hipDeviceCanAccessPeer.peerDeviceId = peerDeviceId; \ -}; -#define INIT_hipDeviceEnablePeerAccess_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceEnablePeerAccess.peerDeviceId = peerDeviceId; \ - cb_data.args.hipDeviceEnablePeerAccess.flags = flags; \ -}; -#define INIT_hipModuleUnload_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipModuleUnload.module = hmod; \ -}; -#define INIT_hipHostUnregister_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipHostUnregister.hostPtr = hostPtr; \ -}; -#define INIT_hipProfilerStop_CB_ARGS_DATA(cb_data) { \ -}; -#define INIT_hipLaunchByPtr_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipLaunchByPtr.func = hostFunction; \ -}; -#define INIT_hipStreamSynchronize_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamSynchronize.stream = stream; \ -}; -#define INIT_hipDeviceSetCacheConfig_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceSetCacheConfig.cacheConfig = cacheConfig; \ -}; -#define INIT_hipGetErrorName_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGetErrorName.hip_error = hip_error; \ -}; -#define INIT_hipMemcpyHtoD_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyHtoD.dst = dst; \ - cb_data.args.hipMemcpyHtoD.src = src; \ - cb_data.args.hipMemcpyHtoD.sizeBytes = sizeBytes; \ -}; -#define INIT_hipMemcpyHtoA_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyHtoA.dstArray = dstArray; \ - cb_data.args.hipMemcpyHtoA.dstOffset = dstOffset; \ - cb_data.args.hipMemcpyHtoA.srcHost = srcHost; \ - cb_data.args.hipMemcpyHtoA.count = count; \ -}; -#define INIT_hipCtxCreate_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxCreate.ctx = ctx; \ - cb_data.args.hipCtxCreate.flags = flags; \ - cb_data.args.hipCtxCreate.device = device; \ -}; -#define INIT_hipMemcpy2D_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpy2D.dst = dst; \ - cb_data.args.hipMemcpy2D.dpitch = dpitch; \ - cb_data.args.hipMemcpy2D.src = src; \ - cb_data.args.hipMemcpy2D.spitch = spitch; \ - cb_data.args.hipMemcpy2D.width = width; \ - cb_data.args.hipMemcpy2D.height = height; \ - cb_data.args.hipMemcpy2D.kind = kind; \ -}; -#define INIT_hipIpcCloseMemHandle_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipIpcCloseMemHandle.devPtr = devPtr; \ -}; -#define INIT_hipChooseDevice_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipChooseDevice.device = device; \ - cb_data.args.hipChooseDevice.prop = prop; \ -}; -#define INIT_hipDeviceSetSharedMemConfig_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceSetSharedMemConfig.config = config; \ -}; -#define INIT_hipDeviceComputeCapability_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceComputeCapability.major = major; \ - cb_data.args.hipDeviceComputeCapability.minor = minor; \ - cb_data.args.hipDeviceComputeCapability.device = device; \ -}; -#define INIT_hipDeviceGet_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceGet.device = device; \ - cb_data.args.hipDeviceGet.ordinal = deviceId; \ -}; -#define INIT_hipProfilerStart_CB_ARGS_DATA(cb_data) { \ -}; -#define INIT_hipCtxSetCacheConfig_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxSetCacheConfig.cacheConfig = cacheConfig; \ -}; -#define INIT_hipFuncSetCacheConfig_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipFuncSetCacheConfig.func = func; \ - cb_data.args.hipFuncSetCacheConfig.config = cacheConfig; \ -}; -#define INIT_hipMemcpyPeerAsync_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyPeerAsync.dst = dst; \ - cb_data.args.hipMemcpyPeerAsync.dstDeviceId = dstDevice; \ - cb_data.args.hipMemcpyPeerAsync.src = src; \ - cb_data.args.hipMemcpyPeerAsync.srcDevice = srcDevice; \ - cb_data.args.hipMemcpyPeerAsync.sizeBytes = sizeBytes; \ - cb_data.args.hipMemcpyPeerAsync.stream = stream; \ -}; -#define INIT_hipEventElapsedTime_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipEventElapsedTime.ms = ms; \ - cb_data.args.hipEventElapsedTime.start = start; \ - cb_data.args.hipEventElapsedTime.stop = stop; \ -}; -#define INIT_hipDevicePrimaryCtxReset_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDevicePrimaryCtxReset.dev = dev; \ -}; -#define INIT_hipEventDestroy_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipEventDestroy.event = event; \ -}; -#define INIT_hipCtxPopCurrent_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxPopCurrent.ctx = ctx; \ -}; -#define INIT_hipHostGetFlags_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipHostGetFlags.flagsPtr = flagsPtr; \ - cb_data.args.hipHostGetFlags.hostPtr = hostPtr; \ -}; -#define INIT_hipHostMalloc_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipHostMalloc.ptr = ptr; \ - cb_data.args.hipHostMalloc.size = sizeBytes; \ - cb_data.args.hipHostMalloc.flags = flags; \ -}; -#define INIT_hipDriverGetVersion_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDriverGetVersion.driverVersion = driverVersion; \ -}; -#define INIT_hipMemGetInfo_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemGetInfo.free = free; \ - cb_data.args.hipMemGetInfo.total = total; \ -}; -#define INIT_hipDeviceReset_CB_ARGS_DATA(cb_data) { \ -}; -#define INIT_hipMemset_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemset.dst = dst; \ - cb_data.args.hipMemset.value = value; \ - cb_data.args.hipMemset.sizeBytes = sizeBytes; \ -}; -#define INIT_hipMemsetD8_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemsetD8.dest = dst; \ - cb_data.args.hipMemsetD8.value = value; \ - cb_data.args.hipMemsetD8.sizeBytes = sizeBytes; \ -}; -#define INIT_hipHostRegister_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipHostRegister.hostPtr = hostPtr; \ - cb_data.args.hipHostRegister.sizeBytes = sizeBytes; \ - cb_data.args.hipHostRegister.flags = flags; \ -}; -#define INIT_hipCtxSetSharedMemConfig_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxSetSharedMemConfig.config = config; \ -}; -#define INIT_hipArray3DCreate_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipArray3DCreate.array = array; \ - cb_data.args.hipArray3DCreate.pAllocateArray = pAllocateArray; \ -}; -#define INIT_hipIpcOpenMemHandle_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipIpcOpenMemHandle.devPtr = devPtr; \ - cb_data.args.hipIpcOpenMemHandle.handle = handle; \ - cb_data.args.hipIpcOpenMemHandle.flags = flags; \ -}; -#define INIT_hipGetLastError_CB_ARGS_DATA(cb_data) { \ -}; -#define INIT_hipCtxDestroy_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxDestroy.ctx = ctx; \ -}; -#define INIT_hipDeviceGetSharedMemConfig_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceGetSharedMemConfig.pConfig = pConfig; \ -}; -#define INIT_hipSetDeviceFlags_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipSetDeviceFlags.flags = flags; \ -}; -#define INIT_hipHccModuleLaunchKernel_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipHccModuleLaunchKernel.f = f; \ - cb_data.args.hipHccModuleLaunchKernel.globalWorkSizeX = globalWorkSizeX; \ - cb_data.args.hipHccModuleLaunchKernel.globalWorkSizeY = globalWorkSizeY; \ - cb_data.args.hipHccModuleLaunchKernel.globalWorkSizeZ = globalWorkSizeZ; \ - cb_data.args.hipHccModuleLaunchKernel.localWorkSizeX = localWorkSizeX; \ - cb_data.args.hipHccModuleLaunchKernel.localWorkSizeY = localWorkSizeY; \ - cb_data.args.hipHccModuleLaunchKernel.localWorkSizeZ = localWorkSizeZ; \ - cb_data.args.hipHccModuleLaunchKernel.sharedMemBytes = sharedMemBytes; \ - cb_data.args.hipHccModuleLaunchKernel.hStream = hStream; \ - cb_data.args.hipHccModuleLaunchKernel.kernelParams = kernelParams; \ - cb_data.args.hipHccModuleLaunchKernel.extra = extra; \ - cb_data.args.hipHccModuleLaunchKernel.startEvent = startEvent; \ - cb_data.args.hipHccModuleLaunchKernel.stopEvent = stopEvent; \ -}; -#define INIT_hipFree_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipFree.ptr = ptr; \ -}; -#define INIT_hipDeviceGetAttribute_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceGetAttribute.pi = pi; \ - cb_data.args.hipDeviceGetAttribute.attr = attr; \ - cb_data.args.hipDeviceGetAttribute.deviceId = device; \ -}; -#define INIT_hipMemcpyDtoH_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyDtoH.dst = dst; \ - cb_data.args.hipMemcpyDtoH.src = src; \ - cb_data.args.hipMemcpyDtoH.sizeBytes = sizeBytes; \ -}; -#define INIT_hipCtxDisablePeerAccess_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxDisablePeerAccess.peerCtx = peerCtx; \ -}; -#define INIT_hipDeviceGetByPCIBusId_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceGetByPCIBusId.device = device; \ - cb_data.args.hipDeviceGetByPCIBusId.pciBusId = pciBusId; \ -}; -#define INIT_hipIpcGetMemHandle_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipIpcGetMemHandle.handle = handle; \ - cb_data.args.hipIpcGetMemHandle.devPtr = devPtr; \ -}; -#define INIT_hipMemcpyHtoDAsync_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyHtoDAsync.dst = dst; \ - cb_data.args.hipMemcpyHtoDAsync.src = src; \ - cb_data.args.hipMemcpyHtoDAsync.sizeBytes = sizeBytes; \ - cb_data.args.hipMemcpyHtoDAsync.stream = stream; \ -}; -#define INIT_hipCtxGetDevice_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxGetDevice.device = device; \ -}; -#define INIT_hipMemset3D_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemset3D.pitchedDevPtr = pitchedDevPtr; \ - cb_data.args.hipMemset3D.value = value; \ - cb_data.args.hipMemset3D.extent = extent; \ -}; -#define INIT_hipModuleLoadData_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipModuleLoadData.module = module; \ - cb_data.args.hipModuleLoadData.image = image; \ -}; -#define INIT_hipDeviceTotalMem_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceTotalMem.bytes = bytes; \ - cb_data.args.hipDeviceTotalMem.device = device; \ -}; -#define INIT_hipCtxSetCurrent_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxSetCurrent.ctx = ctx; \ -}; -#define INIT_hipGetErrorString_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGetErrorString.hipError = hip_error; \ -}; -#define INIT_hipDevicePrimaryCtxRetain_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDevicePrimaryCtxRetain.pctx = pctx; \ - cb_data.args.hipDevicePrimaryCtxRetain.dev = dev; \ -}; -#define INIT_hipDeviceDisablePeerAccess_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceDisablePeerAccess.peerDeviceId = peerDeviceId; \ -}; -#define INIT_hipStreamCreateWithFlags_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamCreateWithFlags.stream = stream; \ - cb_data.args.hipStreamCreateWithFlags.flags = flags; \ -}; -#define INIT_hipMemcpyFromArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyFromArray.dst = dst; \ - cb_data.args.hipMemcpyFromArray.srcArray = srcArray; \ - cb_data.args.hipMemcpyFromArray.wOffset = wOffset; \ - cb_data.args.hipMemcpyFromArray.hOffset = hOffset; \ - cb_data.args.hipMemcpyFromArray.count = count; \ - cb_data.args.hipMemcpyFromArray.kind = kind; \ -}; -#define INIT_hipMemcpy2DAsync_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpy2DAsync.dst = dst; \ - cb_data.args.hipMemcpy2DAsync.dpitch = dpitch; \ - cb_data.args.hipMemcpy2DAsync.src = src; \ - cb_data.args.hipMemcpy2DAsync.spitch = spitch; \ - cb_data.args.hipMemcpy2DAsync.width = width; \ - cb_data.args.hipMemcpy2DAsync.height = height; \ - cb_data.args.hipMemcpy2DAsync.kind = kind; \ - cb_data.args.hipMemcpy2DAsync.stream = stream; \ -}; -#define INIT_hipEventCreateWithFlags_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipEventCreateWithFlags.event = event; \ - cb_data.args.hipEventCreateWithFlags.flags = flags; \ -}; -#define INIT_hipStreamQuery_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamQuery.stream = stream; \ -}; -#define INIT_hipDeviceGetPCIBusId_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceGetPCIBusId.pciBusId = pciBusId; \ - cb_data.args.hipDeviceGetPCIBusId.len = len; \ - cb_data.args.hipDeviceGetPCIBusId.device = device; \ -}; -#define INIT_hipMemcpy_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpy.dst = dst; \ - cb_data.args.hipMemcpy.src = src; \ - cb_data.args.hipMemcpy.sizeBytes = sizeBytes; \ - cb_data.args.hipMemcpy.kind = kind; \ -}; -#define INIT_hipPeekAtLastError_CB_ARGS_DATA(cb_data) { \ -}; -#define INIT_hipStreamAddCallback_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipStreamAddCallback.stream = stream; \ - cb_data.args.hipStreamAddCallback.callback = callback; \ - cb_data.args.hipStreamAddCallback.userData = userData; \ - cb_data.args.hipStreamAddCallback.flags = flags; \ -}; -#define INIT_hipMemcpyToArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyToArray.dst = dst; \ - cb_data.args.hipMemcpyToArray.wOffset = wOffset; \ - cb_data.args.hipMemcpyToArray.hOffset = hOffset; \ - cb_data.args.hipMemcpyToArray.src = src; \ - cb_data.args.hipMemcpyToArray.count = count; \ - cb_data.args.hipMemcpyToArray.kind = kind; \ -}; -#define INIT_hipMemsetD32_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemsetD32.dest = dst; \ - cb_data.args.hipMemsetD32.value = value; \ - cb_data.args.hipMemsetD32.count = count; \ -}; -#define INIT_hipDeviceSynchronize_CB_ARGS_DATA(cb_data) { \ -}; -#define INIT_hipDeviceGetCacheConfig_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceGetCacheConfig.cacheConfig = cacheConfig; \ -}; -#define INIT_hipMalloc3D_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMalloc3D.pitchedDevPtr = pitchedDevPtr; \ - cb_data.args.hipMalloc3D.extent = extent; \ -}; -#define INIT_hipPointerGetAttributes_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipPointerGetAttributes.attributes = attributes; \ - cb_data.args.hipPointerGetAttributes.ptr = ptr; \ -}; -#define INIT_hipMemsetAsync_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemsetAsync.dst = dst; \ - cb_data.args.hipMemsetAsync.value = value; \ - cb_data.args.hipMemsetAsync.sizeBytes = sizeBytes; \ - cb_data.args.hipMemsetAsync.stream = stream; \ -}; -#define INIT_hipDeviceGetName_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceGetName.name = name; \ - cb_data.args.hipDeviceGetName.len = len; \ - cb_data.args.hipDeviceGetName.device = device; \ -}; -#define INIT_hipCtxPushCurrent_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxPushCurrent.ctx = ctx; \ -}; -#define INIT_hipMemcpyPeer_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyPeer.dst = dst; \ - cb_data.args.hipMemcpyPeer.dstDeviceId = dstDevice; \ - cb_data.args.hipMemcpyPeer.src = src; \ - cb_data.args.hipMemcpyPeer.srcDeviceId = srcDevice; \ - cb_data.args.hipMemcpyPeer.sizeBytes = sizeBytes; \ -}; -#define INIT_hipEventSynchronize_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipEventSynchronize.event = event; \ -}; -#define INIT_hipMemcpyDtoDAsync_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyDtoDAsync.dst = dst; \ - cb_data.args.hipMemcpyDtoDAsync.src = src; \ - cb_data.args.hipMemcpyDtoDAsync.sizeBytes = sizeBytes; \ - cb_data.args.hipMemcpyDtoDAsync.stream = stream; \ -}; -#define INIT_hipCtxEnablePeerAccess_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipCtxEnablePeerAccess.peerCtx = peerCtx; \ - cb_data.args.hipCtxEnablePeerAccess.flags = flags; \ -}; -#define INIT_hipMemcpyDtoHAsync_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyDtoHAsync.dst = dst; \ - cb_data.args.hipMemcpyDtoHAsync.src = src; \ - cb_data.args.hipMemcpyDtoHAsync.sizeBytes = sizeBytes; \ - cb_data.args.hipMemcpyDtoHAsync.stream = stream; \ -}; -#define INIT_hipModuleLaunchKernel_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipModuleLaunchKernel.f = f; \ - cb_data.args.hipModuleLaunchKernel.gridDimX = gridDimX; \ - cb_data.args.hipModuleLaunchKernel.gridDimY = gridDimY; \ - cb_data.args.hipModuleLaunchKernel.gridDimZ = gridDimZ; \ - cb_data.args.hipModuleLaunchKernel.blockDimX = blockDimX; \ - cb_data.args.hipModuleLaunchKernel.blockDimY = blockDimY; \ - cb_data.args.hipModuleLaunchKernel.blockDimZ = blockDimZ; \ - cb_data.args.hipModuleLaunchKernel.sharedMemBytes = sharedMemBytes; \ - cb_data.args.hipModuleLaunchKernel.stream = hStream; \ - cb_data.args.hipModuleLaunchKernel.kernelParams = kernelParams; \ - cb_data.args.hipModuleLaunchKernel.extra = extra; \ -}; -#define INIT_hipModuleGetTexRef_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipModuleGetTexRef.texRef = texRef; \ - cb_data.args.hipModuleGetTexRef.hmod = hmod; \ - cb_data.args.hipModuleGetTexRef.name = name; \ -}; -#define INIT_hipDeviceGetLimit_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDeviceGetLimit.pValue = pValue; \ - cb_data.args.hipDeviceGetLimit.limit = limit; \ -}; -#define INIT_hipModuleLoadDataEx_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipModuleLoadDataEx.module = module; \ - cb_data.args.hipModuleLoadDataEx.image = image; \ - cb_data.args.hipModuleLoadDataEx.numOptions = numOptions; \ - cb_data.args.hipModuleLoadDataEx.options = options; \ - cb_data.args.hipModuleLoadDataEx.optionValues = optionValues; \ -}; -#define INIT_hipRuntimeGetVersion_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipRuntimeGetVersion.runtimeVersion = runtimeVersion; \ -}; -#define INIT_hipGetDeviceProperties_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGetDeviceProperties.prop = props; \ - cb_data.args.hipGetDeviceProperties.deviceId = device; \ -}; -#define INIT_hipFreeArray_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipFreeArray.array = array; \ -}; -#define INIT_hipDevicePrimaryCtxRelease_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipDevicePrimaryCtxRelease.dev = dev; \ -}; -#define INIT_hipHostGetDevicePointer_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipHostGetDevicePointer.devPtr = devicePointer; \ - cb_data.args.hipHostGetDevicePointer.hstPtr = hostPointer; \ - cb_data.args.hipHostGetDevicePointer.flags = flags; \ -}; -#define INIT_hipMemcpyParam2D_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemcpyParam2D.pCopy = pCopy; \ -}; -#define INIT_hipModuleGetFunction_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipModuleGetFunction.function = hfunc; \ - cb_data.args.hipModuleGetFunction.module = hmod; \ - cb_data.args.hipModuleGetFunction.kname = name; \ -}; -#define INIT_hipMemsetD32Async_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipMemsetD32Async.dst = dst; \ - cb_data.args.hipMemsetD32Async.value = value; \ - cb_data.args.hipMemsetD32Async.count = count; \ - cb_data.args.hipMemsetD32Async.stream = stream; \ -}; -#define INIT_hipGetDevice_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGetDevice.deviceId = deviceId; \ -}; -#define INIT_hipGetDeviceCount_CB_ARGS_DATA(cb_data) { \ - cb_data.args.hipGetDeviceCount.count = count; \ -}; -#define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data) - -#if 0 -// HIP API string method, method name and parameters -const char* hipApiString(hip_api_id_t id, const hip_api_data_t* data) { - std::ostringstream oss; - switch (id) { - case HIP_API_ID_hipStreamCreateWithPriority: - oss << "hipStreamCreateWithPriority(" - << " stream=" << data->args.hipStreamCreateWithPriority.stream << "," - << " flags=" << data->args.hipStreamCreateWithPriority.flags << "," - << " priority=" << data->args.hipStreamCreateWithPriority.priority - << ")"; - break; - case HIP_API_ID_hipMallocPitch: - oss << "hipMallocPitch(" - << " ptr=" << data->args.hipMallocPitch.ptr << "," - << " pitch=" << data->args.hipMallocPitch.pitch << "," - << " width=" << data->args.hipMallocPitch.width << "," - << " height=" << data->args.hipMallocPitch.height - << ")"; - break; - case HIP_API_ID_hipMalloc: - oss << "hipMalloc(" - << " ptr=" << data->args.hipMalloc.ptr << "," - << " size=" << data->args.hipMalloc.size - << ")"; - break; - case HIP_API_ID_hipEventRecord: - oss << "hipEventRecord(" - << " event=" << data->args.hipEventRecord.event << "," - << " stream=" << data->args.hipEventRecord.stream - << ")"; - break; - case HIP_API_ID_hipCtxSynchronize: - oss << "hipCtxSynchronize(" - << ")"; - break; - case HIP_API_ID_hipSetDevice: - oss << "hipSetDevice(" - << " deviceId=" << data->args.hipSetDevice.deviceId - << ")"; - break; - case HIP_API_ID_hipCtxGetApiVersion: - oss << "hipCtxGetApiVersion(" - << " ctx=" << data->args.hipCtxGetApiVersion.ctx << "," - << " apiVersion=" << data->args.hipCtxGetApiVersion.apiVersion - << ")"; - break; - case HIP_API_ID_hipSetupArgument: - oss << "hipSetupArgument(" - << " arg=" << data->args.hipSetupArgument.arg << "," - << " size=" << data->args.hipSetupArgument.size << "," - << " offset=" << data->args.hipSetupArgument.offset - << ")"; - break; - case HIP_API_ID_hipMemcpyDtoD: - oss << "hipMemcpyDtoD(" - << " dst=" << data->args.hipMemcpyDtoD.dst << "," - << " src=" << data->args.hipMemcpyDtoD.src << "," - << " sizeBytes=" << data->args.hipMemcpyDtoD.sizeBytes - << ")"; - break; - case HIP_API_ID_hipHostFree: - oss << "hipHostFree(" - << " ptr=" << data->args.hipHostFree.ptr - << ")"; - break; - case HIP_API_ID_hipMemcpy2DToArray: - oss << "hipMemcpy2DToArray(" - << " dst=" << data->args.hipMemcpy2DToArray.dst << "," - << " wOffset=" << data->args.hipMemcpy2DToArray.wOffset << "," - << " hOffset=" << data->args.hipMemcpy2DToArray.hOffset << "," - << " src=" << data->args.hipMemcpy2DToArray.src << "," - << " spitch=" << data->args.hipMemcpy2DToArray.spitch << "," - << " width=" << data->args.hipMemcpy2DToArray.width << "," - << " height=" << data->args.hipMemcpy2DToArray.height << "," - << " kind=" << data->args.hipMemcpy2DToArray.kind - << ")"; - break; - case HIP_API_ID_hipCtxGetCacheConfig: - oss << "hipCtxGetCacheConfig(" - << " cacheConfig=" << data->args.hipCtxGetCacheConfig.cacheConfig - << ")"; - break; - case HIP_API_ID_hipStreamWaitEvent: - oss << "hipStreamWaitEvent(" - << " stream=" << data->args.hipStreamWaitEvent.stream << "," - << " event=" << data->args.hipStreamWaitEvent.event << "," - << " flags=" << data->args.hipStreamWaitEvent.flags - << ")"; - break; - case HIP_API_ID_hipDeviceGetStreamPriorityRange: - oss << "hipDeviceGetStreamPriorityRange(" - << " leastPriority=" << data->args.hipDeviceGetStreamPriorityRange.leastPriority << "," - << " greatestPriority=" << data->args.hipDeviceGetStreamPriorityRange.greatestPriority - << ")"; - break; - case HIP_API_ID_hipModuleLoad: - oss << "hipModuleLoad(" - << " module=" << data->args.hipModuleLoad.module << "," - << " fname=" << data->args.hipModuleLoad.fname - << ")"; - break; - case HIP_API_ID_hipDevicePrimaryCtxSetFlags: - oss << "hipDevicePrimaryCtxSetFlags(" - << " dev=" << data->args.hipDevicePrimaryCtxSetFlags.dev << "," - << " flags=" << data->args.hipDevicePrimaryCtxSetFlags.flags - << ")"; - break; - case HIP_API_ID_hipMemcpyAsync: - oss << "hipMemcpyAsync(" - << " dst=" << data->args.hipMemcpyAsync.dst << "," - << " src=" << data->args.hipMemcpyAsync.src << "," - << " sizeBytes=" << data->args.hipMemcpyAsync.sizeBytes << "," - << " kind=" << data->args.hipMemcpyAsync.kind << "," - << " stream=" << data->args.hipMemcpyAsync.stream - << ")"; - break; - case HIP_API_ID_hipMalloc3DArray: - oss << "hipMalloc3DArray(" - << " array=" << data->args.hipMalloc3DArray.array << "," - << " desc=" << data->args.hipMalloc3DArray.desc << "," - << " extent=" << data->args.hipMalloc3DArray.extent << "," - << " flags=" << data->args.hipMalloc3DArray.flags - << ")"; - break; - case HIP_API_ID_hipStreamCreate: - oss << "hipStreamCreate(" - << " stream=" << data->args.hipStreamCreate.stream - << ")"; - break; - case HIP_API_ID_hipCtxGetCurrent: - oss << "hipCtxGetCurrent(" - << " ctx=" << data->args.hipCtxGetCurrent.ctx - << ")"; - break; - case HIP_API_ID_hipDevicePrimaryCtxGetState: - oss << "hipDevicePrimaryCtxGetState(" - << " dev=" << data->args.hipDevicePrimaryCtxGetState.dev << "," - << " flags=" << data->args.hipDevicePrimaryCtxGetState.flags << "," - << " active=" << data->args.hipDevicePrimaryCtxGetState.active - << ")"; - break; - case HIP_API_ID_hipEventQuery: - oss << "hipEventQuery(" - << " event=" << data->args.hipEventQuery.event - << ")"; - break; - case HIP_API_ID_hipEventCreate: - oss << "hipEventCreate(" - << " event=" << data->args.hipEventCreate.event - << ")"; - break; - case HIP_API_ID_hipMemGetAddressRange: - oss << "hipMemGetAddressRange(" - << " pbase=" << data->args.hipMemGetAddressRange.pbase << "," - << " psize=" << data->args.hipMemGetAddressRange.psize << "," - << " dptr=" << data->args.hipMemGetAddressRange.dptr - << ")"; - break; - case HIP_API_ID_hipArrayCreate: - oss << "hipArrayCreate(" - << " pHandle=" << data->args.hipArrayCreate.pHandle << "," - << " pAllocateArray=" << data->args.hipArrayCreate.pAllocateArray - << ")"; - break; - case HIP_API_ID_hipStreamGetFlags: - oss << "hipStreamGetFlags(" - << " stream=" << data->args.hipStreamGetFlags.stream << "," - << " flags=" << data->args.hipStreamGetFlags.flags - << ")"; - break; - case HIP_API_ID_hipMallocArray: - oss << "hipMallocArray(" - << " array=" << data->args.hipMallocArray.array << "," - << " desc=" << data->args.hipMallocArray.desc << "," - << " width=" << data->args.hipMallocArray.width << "," - << " height=" << data->args.hipMallocArray.height << "," - << " flags=" << data->args.hipMallocArray.flags - << ")"; - break; - case HIP_API_ID_hipCtxGetSharedMemConfig: - oss << "hipCtxGetSharedMemConfig(" - << " pConfig=" << data->args.hipCtxGetSharedMemConfig.pConfig - << ")"; - break; - case HIP_API_ID_hipMemPtrGetInfo: - oss << "hipMemPtrGetInfo(" - << " ptr=" << data->args.hipMemPtrGetInfo.ptr << "," - << " size=" << data->args.hipMemPtrGetInfo.size - << ")"; - break; - case HIP_API_ID_hipCtxGetFlags: - oss << "hipCtxGetFlags(" - << " flags=" << data->args.hipCtxGetFlags.flags - << ")"; - break; - case HIP_API_ID_hipStreamDestroy: - oss << "hipStreamDestroy(" - << " stream=" << data->args.hipStreamDestroy.stream - << ")"; - break; - case HIP_API_ID_hipMemset3DAsync: - oss << "hipMemset3DAsync(" - << " pitchedDevPtr=" << data->args.hipMemset3DAsync.pitchedDevPtr << "," - << " value=" << data->args.hipMemset3DAsync.value << "," - << " extent=" << data->args.hipMemset3DAsync.extent << "," - << " stream=" << data->args.hipMemset3DAsync.stream - << ")"; - break; - case HIP_API_ID_hipMemcpy3D: - oss << "hipMemcpy3D(" - << " p=" << data->args.hipMemcpy3D.p - << ")"; - break; - case HIP_API_ID_hipInit: - oss << "hipInit(" - << " flags=" << data->args.hipInit.flags - << ")"; - break; - case HIP_API_ID_hipMemcpyAtoH: - oss << "hipMemcpyAtoH(" - << " dst=" << data->args.hipMemcpyAtoH.dst << "," - << " srcArray=" << data->args.hipMemcpyAtoH.srcArray << "," - << " srcOffset=" << data->args.hipMemcpyAtoH.srcOffset << "," - << " count=" << data->args.hipMemcpyAtoH.count - << ")"; - break; - case HIP_API_ID_hipStreamGetPriority: - oss << "hipStreamGetPriority(" - << " stream=" << data->args.hipStreamGetPriority.stream << "," - << " priority=" << data->args.hipStreamGetPriority.priority - << ")"; - break; - case HIP_API_ID_hipMemset2D: - oss << "hipMemset2D(" - << " dst=" << data->args.hipMemset2D.dst << "," - << " pitch=" << data->args.hipMemset2D.pitch << "," - << " value=" << data->args.hipMemset2D.value << "," - << " width=" << data->args.hipMemset2D.width << "," - << " height=" << data->args.hipMemset2D.height - << ")"; - break; - case HIP_API_ID_hipMemset2DAsync: - oss << "hipMemset2DAsync(" - << " dst=" << data->args.hipMemset2DAsync.dst << "," - << " pitch=" << data->args.hipMemset2DAsync.pitch << "," - << " value=" << data->args.hipMemset2DAsync.value << "," - << " width=" << data->args.hipMemset2DAsync.width << "," - << " height=" << data->args.hipMemset2DAsync.height << "," - << " stream=" << data->args.hipMemset2DAsync.stream - << ")"; - break; - case HIP_API_ID_hipDeviceCanAccessPeer: - oss << "hipDeviceCanAccessPeer(" - << " canAccessPeer=" << data->args.hipDeviceCanAccessPeer.canAccessPeer << "," - << " deviceId=" << data->args.hipDeviceCanAccessPeer.deviceId << "," - << " peerDeviceId=" << data->args.hipDeviceCanAccessPeer.peerDeviceId - << ")"; - break; - case HIP_API_ID_hipDeviceEnablePeerAccess: - oss << "hipDeviceEnablePeerAccess(" - << " peerDeviceId=" << data->args.hipDeviceEnablePeerAccess.peerDeviceId << "," - << " flags=" << data->args.hipDeviceEnablePeerAccess.flags - << ")"; - break; - case HIP_API_ID_hipModuleUnload: - oss << "hipModuleUnload(" - << " module=" << data->args.hipModuleUnload.module - << ")"; - break; - case HIP_API_ID_hipHostUnregister: - oss << "hipHostUnregister(" - << " hostPtr=" << data->args.hipHostUnregister.hostPtr - << ")"; - break; - case HIP_API_ID_hipProfilerStop: - oss << "hipProfilerStop(" - << ")"; - break; - case HIP_API_ID_hipLaunchByPtr: - oss << "hipLaunchByPtr(" - << " func=" << data->args.hipLaunchByPtr.func - << ")"; - break; - case HIP_API_ID_hipStreamSynchronize: - oss << "hipStreamSynchronize(" - << " stream=" << data->args.hipStreamSynchronize.stream - << ")"; - break; - case HIP_API_ID_hipDeviceSetCacheConfig: - oss << "hipDeviceSetCacheConfig(" - << " cacheConfig=" << data->args.hipDeviceSetCacheConfig.cacheConfig - << ")"; - break; - case HIP_API_ID_hipGetErrorName: - oss << "hipGetErrorName(" - << " hip_error=" << data->args.hipGetErrorName.hip_error - << ")"; - break; - case HIP_API_ID_hipMemcpyHtoD: - oss << "hipMemcpyHtoD(" - << " dst=" << data->args.hipMemcpyHtoD.dst << "," - << " src=" << data->args.hipMemcpyHtoD.src << "," - << " sizeBytes=" << data->args.hipMemcpyHtoD.sizeBytes - << ")"; - break; - case HIP_API_ID_hipMemcpyHtoA: - oss << "hipMemcpyHtoA(" - << " dstArray=" << data->args.hipMemcpyHtoA.dstArray << "," - << " dstOffset=" << data->args.hipMemcpyHtoA.dstOffset << "," - << " srcHost=" << data->args.hipMemcpyHtoA.srcHost << "," - << " count=" << data->args.hipMemcpyHtoA.count - << ")"; - break; - case HIP_API_ID_hipCtxCreate: - oss << "hipCtxCreate(" - << " ctx=" << data->args.hipCtxCreate.ctx << "," - << " flags=" << data->args.hipCtxCreate.flags << "," - << " device=" << data->args.hipCtxCreate.device - << ")"; - break; - case HIP_API_ID_hipMemcpy2D: - oss << "hipMemcpy2D(" - << " dst=" << data->args.hipMemcpy2D.dst << "," - << " dpitch=" << data->args.hipMemcpy2D.dpitch << "," - << " src=" << data->args.hipMemcpy2D.src << "," - << " spitch=" << data->args.hipMemcpy2D.spitch << "," - << " width=" << data->args.hipMemcpy2D.width << "," - << " height=" << data->args.hipMemcpy2D.height << "," - << " kind=" << data->args.hipMemcpy2D.kind - << ")"; - break; - case HIP_API_ID_hipIpcCloseMemHandle: - oss << "hipIpcCloseMemHandle(" - << " devPtr=" << data->args.hipIpcCloseMemHandle.devPtr - << ")"; - break; - case HIP_API_ID_hipChooseDevice: - oss << "hipChooseDevice(" - << " device=" << data->args.hipChooseDevice.device << "," - << " prop=" << data->args.hipChooseDevice.prop - << ")"; - break; - case HIP_API_ID_hipDeviceSetSharedMemConfig: - oss << "hipDeviceSetSharedMemConfig(" - << " config=" << data->args.hipDeviceSetSharedMemConfig.config - << ")"; - break; - case HIP_API_ID_hipDeviceComputeCapability: - oss << "hipDeviceComputeCapability(" - << " major=" << data->args.hipDeviceComputeCapability.major << "," - << " minor=" << data->args.hipDeviceComputeCapability.minor << "," - << " device=" << data->args.hipDeviceComputeCapability.device - << ")"; - break; - case HIP_API_ID_hipDeviceGet: - oss << "hipDeviceGet(" - << " device=" << data->args.hipDeviceGet.device << "," - << " ordinal=" << data->args.hipDeviceGet.ordinal - << ")"; - break; - case HIP_API_ID_hipProfilerStart: - oss << "hipProfilerStart(" - << ")"; - break; - case HIP_API_ID_hipCtxSetCacheConfig: - oss << "hipCtxSetCacheConfig(" - << " cacheConfig=" << data->args.hipCtxSetCacheConfig.cacheConfig - << ")"; - break; - case HIP_API_ID_hipFuncSetCacheConfig: - oss << "hipFuncSetCacheConfig(" - << " func=" << data->args.hipFuncSetCacheConfig.func << "," - << " config=" << data->args.hipFuncSetCacheConfig.config - << ")"; - break; - case HIP_API_ID_hipMemcpyPeerAsync: - oss << "hipMemcpyPeerAsync(" - << " dst=" << data->args.hipMemcpyPeerAsync.dst << "," - << " dstDeviceId=" << data->args.hipMemcpyPeerAsync.dstDeviceId << "," - << " src=" << data->args.hipMemcpyPeerAsync.src << "," - << " srcDevice=" << data->args.hipMemcpyPeerAsync.srcDevice << "," - << " sizeBytes=" << data->args.hipMemcpyPeerAsync.sizeBytes << "," - << " stream=" << data->args.hipMemcpyPeerAsync.stream - << ")"; - break; - case HIP_API_ID_hipEventElapsedTime: - oss << "hipEventElapsedTime(" - << " ms=" << data->args.hipEventElapsedTime.ms << "," - << " start=" << data->args.hipEventElapsedTime.start << "," - << " stop=" << data->args.hipEventElapsedTime.stop - << ")"; - break; - case HIP_API_ID_hipDevicePrimaryCtxReset: - oss << "hipDevicePrimaryCtxReset(" - << " dev=" << data->args.hipDevicePrimaryCtxReset.dev - << ")"; - break; - case HIP_API_ID_hipEventDestroy: - oss << "hipEventDestroy(" - << " event=" << data->args.hipEventDestroy.event - << ")"; - break; - case HIP_API_ID_hipCtxPopCurrent: - oss << "hipCtxPopCurrent(" - << " ctx=" << data->args.hipCtxPopCurrent.ctx - << ")"; - break; - case HIP_API_ID_hipHostGetFlags: - oss << "hipHostGetFlags(" - << " flagsPtr=" << data->args.hipHostGetFlags.flagsPtr << "," - << " hostPtr=" << data->args.hipHostGetFlags.hostPtr - << ")"; - break; - case HIP_API_ID_hipHostMalloc: - oss << "hipHostMalloc(" - << " ptr=" << data->args.hipHostMalloc.ptr << "," - << " size=" << data->args.hipHostMalloc.size << "," - << " flags=" << data->args.hipHostMalloc.flags - << ")"; - break; - case HIP_API_ID_hipDriverGetVersion: - oss << "hipDriverGetVersion(" - << " driverVersion=" << data->args.hipDriverGetVersion.driverVersion - << ")"; - break; - case HIP_API_ID_hipMemGetInfo: - oss << "hipMemGetInfo(" - << " free=" << data->args.hipMemGetInfo.free << "," - << " total=" << data->args.hipMemGetInfo.total - << ")"; - break; - case HIP_API_ID_hipDeviceReset: - oss << "hipDeviceReset(" - << ")"; - break; - case HIP_API_ID_hipMemset: - oss << "hipMemset(" - << " dst=" << data->args.hipMemset.dst << "," - << " value=" << data->args.hipMemset.value << "," - << " sizeBytes=" << data->args.hipMemset.sizeBytes - << ")"; - break; - case HIP_API_ID_hipMemsetD8: - oss << "hipMemsetD8(" - << " dest=" << data->args.hipMemsetD8.dest << "," - << " value=" << data->args.hipMemsetD8.value << "," - << " sizeBytes=" << data->args.hipMemsetD8.sizeBytes - << ")"; - break; - case HIP_API_ID_hipHostRegister: - oss << "hipHostRegister(" - << " hostPtr=" << data->args.hipHostRegister.hostPtr << "," - << " sizeBytes=" << data->args.hipHostRegister.sizeBytes << "," - << " flags=" << data->args.hipHostRegister.flags - << ")"; - break; - case HIP_API_ID_hipCtxSetSharedMemConfig: - oss << "hipCtxSetSharedMemConfig(" - << " config=" << data->args.hipCtxSetSharedMemConfig.config - << ")"; - break; - case HIP_API_ID_hipArray3DCreate: - oss << "hipArray3DCreate(" - << " array=" << data->args.hipArray3DCreate.array << "," - << " pAllocateArray=" << data->args.hipArray3DCreate.pAllocateArray - << ")"; - break; - case HIP_API_ID_hipIpcOpenMemHandle: - oss << "hipIpcOpenMemHandle(" - << " devPtr=" << data->args.hipIpcOpenMemHandle.devPtr << "," - << " handle=" << data->args.hipIpcOpenMemHandle.handle << "," - << " flags=" << data->args.hipIpcOpenMemHandle.flags - << ")"; - break; - case HIP_API_ID_hipGetLastError: - oss << "hipGetLastError(" - << ")"; - break; - case HIP_API_ID_hipCtxDestroy: - oss << "hipCtxDestroy(" - << " ctx=" << data->args.hipCtxDestroy.ctx - << ")"; - break; - case HIP_API_ID_hipDeviceGetSharedMemConfig: - oss << "hipDeviceGetSharedMemConfig(" - << " pConfig=" << data->args.hipDeviceGetSharedMemConfig.pConfig - << ")"; - break; - case HIP_API_ID_hipSetDeviceFlags: - oss << "hipSetDeviceFlags(" - << " flags=" << data->args.hipSetDeviceFlags.flags - << ")"; - break; - case HIP_API_ID_hipHccModuleLaunchKernel: - oss << "hipHccModuleLaunchKernel(" - << " f=" << data->args.hipHccModuleLaunchKernel.f << "," - << " globalWorkSizeX=" << data->args.hipHccModuleLaunchKernel.globalWorkSizeX << "," - << " globalWorkSizeY=" << data->args.hipHccModuleLaunchKernel.globalWorkSizeY << "," - << " globalWorkSizeZ=" << data->args.hipHccModuleLaunchKernel.globalWorkSizeZ << "," - << " localWorkSizeX=" << data->args.hipHccModuleLaunchKernel.localWorkSizeX << "," - << " localWorkSizeY=" << data->args.hipHccModuleLaunchKernel.localWorkSizeY << "," - << " localWorkSizeZ=" << data->args.hipHccModuleLaunchKernel.localWorkSizeZ << "," - << " sharedMemBytes=" << data->args.hipHccModuleLaunchKernel.sharedMemBytes << "," - << " hStream=" << data->args.hipHccModuleLaunchKernel.hStream << "," - << " kernelParams=" << data->args.hipHccModuleLaunchKernel.kernelParams << "," - << " extra=" << data->args.hipHccModuleLaunchKernel.extra << "," - << " startEvent=" << data->args.hipHccModuleLaunchKernel.startEvent << "," - << " stopEvent=" << data->args.hipHccModuleLaunchKernel.stopEvent - << ")"; - break; - case HIP_API_ID_hipFree: - oss << "hipFree(" - << " ptr=" << data->args.hipFree.ptr - << ")"; - break; - case HIP_API_ID_hipDeviceGetAttribute: - oss << "hipDeviceGetAttribute(" - << " pi=" << data->args.hipDeviceGetAttribute.pi << "," - << " attr=" << data->args.hipDeviceGetAttribute.attr << "," - << " deviceId=" << data->args.hipDeviceGetAttribute.deviceId - << ")"; - break; - case HIP_API_ID_hipMemcpyDtoH: - oss << "hipMemcpyDtoH(" - << " dst=" << data->args.hipMemcpyDtoH.dst << "," - << " src=" << data->args.hipMemcpyDtoH.src << "," - << " sizeBytes=" << data->args.hipMemcpyDtoH.sizeBytes - << ")"; - break; - case HIP_API_ID_hipCtxDisablePeerAccess: - oss << "hipCtxDisablePeerAccess(" - << " peerCtx=" << data->args.hipCtxDisablePeerAccess.peerCtx - << ")"; - break; - case HIP_API_ID_hipDeviceGetByPCIBusId: - oss << "hipDeviceGetByPCIBusId(" - << " device=" << data->args.hipDeviceGetByPCIBusId.device << "," - << " pciBusId=" << data->args.hipDeviceGetByPCIBusId.pciBusId - << ")"; - break; - case HIP_API_ID_hipIpcGetMemHandle: - oss << "hipIpcGetMemHandle(" - << " handle=" << data->args.hipIpcGetMemHandle.handle << "," - << " devPtr=" << data->args.hipIpcGetMemHandle.devPtr - << ")"; - break; - case HIP_API_ID_hipMemcpyHtoDAsync: - oss << "hipMemcpyHtoDAsync(" - << " dst=" << data->args.hipMemcpyHtoDAsync.dst << "," - << " src=" << data->args.hipMemcpyHtoDAsync.src << "," - << " sizeBytes=" << data->args.hipMemcpyHtoDAsync.sizeBytes << "," - << " stream=" << data->args.hipMemcpyHtoDAsync.stream - << ")"; - break; - case HIP_API_ID_hipCtxGetDevice: - oss << "hipCtxGetDevice(" - << " device=" << data->args.hipCtxGetDevice.device - << ")"; - break; - case HIP_API_ID_hipMemset3D: - oss << "hipMemset3D(" - << " pitchedDevPtr=" << data->args.hipMemset3D.pitchedDevPtr << "," - << " value=" << data->args.hipMemset3D.value << "," - << " extent=" << data->args.hipMemset3D.extent - << ")"; - break; - case HIP_API_ID_hipModuleLoadData: - oss << "hipModuleLoadData(" - << " module=" << data->args.hipModuleLoadData.module << "," - << " image=" << data->args.hipModuleLoadData.image - << ")"; - break; - case HIP_API_ID_hipDeviceTotalMem: - oss << "hipDeviceTotalMem(" - << " bytes=" << data->args.hipDeviceTotalMem.bytes << "," - << " device=" << data->args.hipDeviceTotalMem.device - << ")"; - break; - case HIP_API_ID_hipCtxSetCurrent: - oss << "hipCtxSetCurrent(" - << " ctx=" << data->args.hipCtxSetCurrent.ctx - << ")"; - break; - case HIP_API_ID_hipGetErrorString: - oss << "hipGetErrorString(" - << " hipError=" << data->args.hipGetErrorString.hipError - << ")"; - break; - case HIP_API_ID_hipDevicePrimaryCtxRetain: - oss << "hipDevicePrimaryCtxRetain(" - << " pctx=" << data->args.hipDevicePrimaryCtxRetain.pctx << "," - << " dev=" << data->args.hipDevicePrimaryCtxRetain.dev - << ")"; - break; - case HIP_API_ID_hipDeviceDisablePeerAccess: - oss << "hipDeviceDisablePeerAccess(" - << " peerDeviceId=" << data->args.hipDeviceDisablePeerAccess.peerDeviceId - << ")"; - break; - case HIP_API_ID_hipStreamCreateWithFlags: - oss << "hipStreamCreateWithFlags(" - << " stream=" << data->args.hipStreamCreateWithFlags.stream << "," - << " flags=" << data->args.hipStreamCreateWithFlags.flags - << ")"; - break; - case HIP_API_ID_hipMemcpyFromArray: - oss << "hipMemcpyFromArray(" - << " dst=" << data->args.hipMemcpyFromArray.dst << "," - << " srcArray=" << data->args.hipMemcpyFromArray.srcArray << "," - << " wOffset=" << data->args.hipMemcpyFromArray.wOffset << "," - << " hOffset=" << data->args.hipMemcpyFromArray.hOffset << "," - << " count=" << data->args.hipMemcpyFromArray.count << "," - << " kind=" << data->args.hipMemcpyFromArray.kind - << ")"; - break; - case HIP_API_ID_hipMemcpy2DAsync: - oss << "hipMemcpy2DAsync(" - << " dst=" << data->args.hipMemcpy2DAsync.dst << "," - << " dpitch=" << data->args.hipMemcpy2DAsync.dpitch << "," - << " src=" << data->args.hipMemcpy2DAsync.src << "," - << " spitch=" << data->args.hipMemcpy2DAsync.spitch << "," - << " width=" << data->args.hipMemcpy2DAsync.width << "," - << " height=" << data->args.hipMemcpy2DAsync.height << "," - << " kind=" << data->args.hipMemcpy2DAsync.kind << "," - << " stream=" << data->args.hipMemcpy2DAsync.stream - << ")"; - break; - case HIP_API_ID_hipEventCreateWithFlags: - oss << "hipEventCreateWithFlags(" - << " event=" << data->args.hipEventCreateWithFlags.event << "," - << " flags=" << data->args.hipEventCreateWithFlags.flags - << ")"; - break; - case HIP_API_ID_hipStreamQuery: - oss << "hipStreamQuery(" - << " stream=" << data->args.hipStreamQuery.stream - << ")"; - break; - case HIP_API_ID_hipDeviceGetPCIBusId: - oss << "hipDeviceGetPCIBusId(" - << " pciBusId=" << data->args.hipDeviceGetPCIBusId.pciBusId << "," - << " len=" << data->args.hipDeviceGetPCIBusId.len << "," - << " device=" << data->args.hipDeviceGetPCIBusId.device - << ")"; - break; - case HIP_API_ID_hipMemcpy: - oss << "hipMemcpy(" - << " dst=" << data->args.hipMemcpy.dst << "," - << " src=" << data->args.hipMemcpy.src << "," - << " sizeBytes=" << data->args.hipMemcpy.sizeBytes << "," - << " kind=" << data->args.hipMemcpy.kind - << ")"; - break; - case HIP_API_ID_hipPeekAtLastError: - oss << "hipPeekAtLastError(" - << ")"; - break; - case HIP_API_ID_hipStreamAddCallback: - oss << "hipStreamAddCallback(" - << " stream=" << data->args.hipStreamAddCallback.stream << "," - << " callback=" << data->args.hipStreamAddCallback.callback << "," - << " userData=" << data->args.hipStreamAddCallback.userData << "," - << " flags=" << data->args.hipStreamAddCallback.flags - << ")"; - break; - case HIP_API_ID_hipMemcpyToArray: - oss << "hipMemcpyToArray(" - << " dst=" << data->args.hipMemcpyToArray.dst << "," - << " wOffset=" << data->args.hipMemcpyToArray.wOffset << "," - << " hOffset=" << data->args.hipMemcpyToArray.hOffset << "," - << " src=" << data->args.hipMemcpyToArray.src << "," - << " count=" << data->args.hipMemcpyToArray.count << "," - << " kind=" << data->args.hipMemcpyToArray.kind - << ")"; - break; - case HIP_API_ID_hipMemsetD32: - oss << "hipMemsetD32(" - << " dest=" << data->args.hipMemsetD32.dest << "," - << " value=" << data->args.hipMemsetD32.value << "," - << " count=" << data->args.hipMemsetD32.count - << ")"; - break; - case HIP_API_ID_hipDeviceSynchronize: - oss << "hipDeviceSynchronize(" - << ")"; - break; - case HIP_API_ID_hipDeviceGetCacheConfig: - oss << "hipDeviceGetCacheConfig(" - << " cacheConfig=" << data->args.hipDeviceGetCacheConfig.cacheConfig - << ")"; - break; - case HIP_API_ID_hipMalloc3D: - oss << "hipMalloc3D(" - << " pitchedDevPtr=" << data->args.hipMalloc3D.pitchedDevPtr << "," - << " extent=" << data->args.hipMalloc3D.extent - << ")"; - break; - case HIP_API_ID_hipPointerGetAttributes: - oss << "hipPointerGetAttributes(" - << " attributes=" << data->args.hipPointerGetAttributes.attributes << "," - << " ptr=" << data->args.hipPointerGetAttributes.ptr - << ")"; - break; - case HIP_API_ID_hipMemsetAsync: - oss << "hipMemsetAsync(" - << " dst=" << data->args.hipMemsetAsync.dst << "," - << " value=" << data->args.hipMemsetAsync.value << "," - << " sizeBytes=" << data->args.hipMemsetAsync.sizeBytes << "," - << " stream=" << data->args.hipMemsetAsync.stream - << ")"; - break; - case HIP_API_ID_hipDeviceGetName: - oss << "hipDeviceGetName(" - << " name=" << data->args.hipDeviceGetName.name << "," - << " len=" << data->args.hipDeviceGetName.len << "," - << " device=" << data->args.hipDeviceGetName.device - << ")"; - break; - case HIP_API_ID_hipCtxPushCurrent: - oss << "hipCtxPushCurrent(" - << " ctx=" << data->args.hipCtxPushCurrent.ctx - << ")"; - break; - case HIP_API_ID_hipMemcpyPeer: - oss << "hipMemcpyPeer(" - << " dst=" << data->args.hipMemcpyPeer.dst << "," - << " dstDeviceId=" << data->args.hipMemcpyPeer.dstDeviceId << "," - << " src=" << data->args.hipMemcpyPeer.src << "," - << " srcDeviceId=" << data->args.hipMemcpyPeer.srcDeviceId << "," - << " sizeBytes=" << data->args.hipMemcpyPeer.sizeBytes - << ")"; - break; - case HIP_API_ID_hipEventSynchronize: - oss << "hipEventSynchronize(" - << " event=" << data->args.hipEventSynchronize.event - << ")"; - break; - case HIP_API_ID_hipMemcpyDtoDAsync: - oss << "hipMemcpyDtoDAsync(" - << " dst=" << data->args.hipMemcpyDtoDAsync.dst << "," - << " src=" << data->args.hipMemcpyDtoDAsync.src << "," - << " sizeBytes=" << data->args.hipMemcpyDtoDAsync.sizeBytes << "," - << " stream=" << data->args.hipMemcpyDtoDAsync.stream - << ")"; - break; - case HIP_API_ID_hipCtxEnablePeerAccess: - oss << "hipCtxEnablePeerAccess(" - << " peerCtx=" << data->args.hipCtxEnablePeerAccess.peerCtx << "," - << " flags=" << data->args.hipCtxEnablePeerAccess.flags - << ")"; - break; - case HIP_API_ID_hipMemcpyDtoHAsync: - oss << "hipMemcpyDtoHAsync(" - << " dst=" << data->args.hipMemcpyDtoHAsync.dst << "," - << " src=" << data->args.hipMemcpyDtoHAsync.src << "," - << " sizeBytes=" << data->args.hipMemcpyDtoHAsync.sizeBytes << "," - << " stream=" << data->args.hipMemcpyDtoHAsync.stream - << ")"; - break; - case HIP_API_ID_hipModuleLaunchKernel: - oss << "hipModuleLaunchKernel(" - << " f=" << data->args.hipModuleLaunchKernel.f << "," - << " gridDimX=" << data->args.hipModuleLaunchKernel.gridDimX << "," - << " gridDimY=" << data->args.hipModuleLaunchKernel.gridDimY << "," - << " gridDimZ=" << data->args.hipModuleLaunchKernel.gridDimZ << "," - << " blockDimX=" << data->args.hipModuleLaunchKernel.blockDimX << "," - << " blockDimY=" << data->args.hipModuleLaunchKernel.blockDimY << "," - << " blockDimZ=" << data->args.hipModuleLaunchKernel.blockDimZ << "," - << " sharedMemBytes=" << data->args.hipModuleLaunchKernel.sharedMemBytes << "," - << " stream=" << data->args.hipModuleLaunchKernel.stream << "," - << " kernelParams=" << data->args.hipModuleLaunchKernel.kernelParams << "," - << " extra=" << data->args.hipModuleLaunchKernel.extra - << ")"; - break; - case HIP_API_ID_hipModuleGetTexRef: - oss << "hipModuleGetTexRef(" - << " texRef=" << data->args.hipModuleGetTexRef.texRef << "," - << " hmod=" << data->args.hipModuleGetTexRef.hmod << "," - << " name=" << data->args.hipModuleGetTexRef.name - << ")"; - break; - case HIP_API_ID_hipDeviceGetLimit: - oss << "hipDeviceGetLimit(" - << " pValue=" << data->args.hipDeviceGetLimit.pValue << "," - << " limit=" << data->args.hipDeviceGetLimit.limit - << ")"; - break; - case HIP_API_ID_hipModuleLoadDataEx: - oss << "hipModuleLoadDataEx(" - << " module=" << data->args.hipModuleLoadDataEx.module << "," - << " image=" << data->args.hipModuleLoadDataEx.image << "," - << " numOptions=" << data->args.hipModuleLoadDataEx.numOptions << "," - << " options=" << data->args.hipModuleLoadDataEx.options << "," - << " optionValues=" << data->args.hipModuleLoadDataEx.optionValues - << ")"; - break; - case HIP_API_ID_hipRuntimeGetVersion: - oss << "hipRuntimeGetVersion(" - << " runtimeVersion=" << data->args.hipRuntimeGetVersion.runtimeVersion - << ")"; - break; - case HIP_API_ID_hipGetDeviceProperties: - oss << "hipGetDeviceProperties(" - << " prop=" << data->args.hipGetDeviceProperties.prop << "," - << " deviceId=" << data->args.hipGetDeviceProperties.deviceId - << ")"; - break; - case HIP_API_ID_hipFreeArray: - oss << "hipFreeArray(" - << " array=" << data->args.hipFreeArray.array - << ")"; - break; - case HIP_API_ID_hipDevicePrimaryCtxRelease: - oss << "hipDevicePrimaryCtxRelease(" - << " dev=" << data->args.hipDevicePrimaryCtxRelease.dev - << ")"; - break; - case HIP_API_ID_hipHostGetDevicePointer: - oss << "hipHostGetDevicePointer(" - << " devPtr=" << data->args.hipHostGetDevicePointer.devPtr << "," - << " hstPtr=" << data->args.hipHostGetDevicePointer.hstPtr << "," - << " flags=" << data->args.hipHostGetDevicePointer.flags - << ")"; - break; - case HIP_API_ID_hipMemcpyParam2D: - oss << "hipMemcpyParam2D(" - << " pCopy=" << data->args.hipMemcpyParam2D.pCopy - << ")"; - break; - case HIP_API_ID_hipModuleGetFunction: - oss << "hipModuleGetFunction(" - << " function=" << data->args.hipModuleGetFunction.function << "," - << " module=" << data->args.hipModuleGetFunction.module << "," - << " kname=" << data->args.hipModuleGetFunction.kname - << ")"; - break; - case HIP_API_ID_hipMemsetD32Async: - oss << "hipMemsetD32Async(" - << " dst=" << data->args.hipMemsetD32Async.dst << "," - << " value=" << data->args.hipMemsetD32Async.value << "," - << " count=" << data->args.hipMemsetD32Async.count << "," - << " stream=" << data->args.hipMemsetD32Async.stream - << ")"; - break; - case HIP_API_ID_hipGetDevice: - oss << "hipGetDevice(" - << " deviceId=" << data->args.hipGetDevice.deviceId - << ")"; - break; - case HIP_API_ID_hipGetDeviceCount: - oss << "hipGetDeviceCount(" - << " count=" << data->args.hipGetDeviceCount.count - << ")"; - break; - default: oss << "unknown"; - }; - return strdup(oss.str().c_str()); -}; -#endif -#endif // _HIP_PROF_STR_H From f167e4daf66e283b5a5b783a82647a9b89205275 Mon Sep 17 00:00:00 2001 From: eshcherb <33529668+eshcherb@users.noreply.github.com> Date: Thu, 21 Mar 2019 00:28:18 -0500 Subject: [PATCH 9/9] adding hip_prof_gen verbose log (#977) * adding hip_prof_gen verbose log * adding stderr fatal error * adding no error exit by default * adding hip_prof_str regeneration dependencies * adding more informative messages * fixing error mesage --- CMakeLists.txt | 11 +++++------ hip_prof_gen.py | 50 ++++++++++++++++++++++++++++++++++++------------- 2 files changed, 42 insertions(+), 19 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 115f746cd9..78d2bc7a30 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -168,12 +168,11 @@ set(PROF_API_STR "${CMAKE_CURRENT_SOURCE_DIR}/include/hip/hcc_detail/hip_prof_st set(PROF_API_HDR "${CMAKE_CURRENT_SOURCE_DIR}/include/hip/hcc_detail/hip_runtime_api.h") set(PROF_API_SRC "${CMAKE_CURRENT_SOURCE_DIR}/src") set(PROF_API_GEN "${CMAKE_CURRENT_SOURCE_DIR}/hip_prof_gen.py") -add_custom_command( - OUTPUT ${PROF_API_STR} - COMMAND ${PROF_API_GEN} ${OPT_PROF_API} ${PROF_API_HDR} ${PROF_API_SRC} ${PROF_API_STR} - DEPENDS ${PROF_API_GEN} ${PROF_API_HDR} -) -add_custom_target(prof ALL DEPENDS ${PROF_API_STR}) +set(PROF_API_LOG "${PROJECT_BINARY_DIR}/hip_prof_gen.log.txt") +set(PROF_API_CMD "${PROF_API_GEN} -v ${OPT_PROF_API} ${PROF_API_HDR} ${PROF_API_SRC} ${PROF_API_STR} >${PROF_API_LOG}") +MESSAGE(STATUS "Generating profiling promitives: ${PROF_API_STR}") +execute_process(COMMAND sh -c "rm -f ${PROF_API_STR}; ${PROF_API_CMD}") +set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${PROF_API_GEN} ${PROF_API_HDR} ${PROF_API_STR}) # Enable profiling API if(USE_PROF_API EQUAL 1) diff --git a/hip_prof_gen.py b/hip_prof_gen.py index 81df5132dd..09bf65ebcd 100755 --- a/hip_prof_gen.py +++ b/hip_prof_gen.py @@ -1,25 +1,36 @@ #!/usr/bin/python import os, sys, re -verbose = 0 PROF_HEADER = "hip_prof_str.h" OUTPUT = PROF_HEADER REC_MAX_LEN = 1024 -# Fatal error termination +# Messages and errors controll +verbose = 0 +errexit = 0 inp_file = 'none' line_num = -1 -def fatal(msg): - if line_num != -1: - print >>sys.stderr, "Error: " + msg + ", file '" + inp_file + "', line (" + str(line_num) + ")" - else: - print >>sys.stderr, "Error: " + msg - sys.exit(1) # Verbose message def message(msg): if verbose: print >>sys.stdout, msg +# Fatal error termination +def error(msg): + if line_num != -1: + msg += ", file '" + inp_file + "', line (" + str(line_num) + ")" + if errexit: + msg = " Error: " + msg + else: + msg = " Warning: " + msg + + print >>sys.stdout, msg + print >>sys.stderr, sys.argv[0] + msg + +def fatal(msg): + error(msg) + if errexit: sys.exit(1) + ############################################################# # Normalizing API arguments def filtr_api_args(args_str): @@ -77,7 +88,7 @@ def parse_api(inp_file_p, out): global line_num inp_file = inp_file_p - beg_pattern = re.compile("^(hipError_t|const char\s*\*)\s+[^\(]+\("); + beg_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\("); api_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)"); end_pattern = re.compile("Texture"); hidden_pattern = re.compile(r'__attribute__\(\(visibility\("hidden"\)\)\)') @@ -99,7 +110,16 @@ def parse_api(inp_file_p, out): if len(record) > REC_MAX_LEN: fatal("bad record \"" + record + "\"") - if beg_pattern.match(record) and (hidden == 0) and (nms_level == 0): found = 1 + m = beg_pattern.match(line) + if m: + name = m.group(2) + if hidden != 0: + message("api: " + name + " - hidden") + elif nms_level != 0: + message("api: " + name + " - hip_impl") + else: + message("api: " + name) + found = 1 if found != 0: record = re.sub("\s__dparm\([^\)]*\)", '', record); @@ -206,7 +226,7 @@ def parse_content(inp_file_p, api_map, out): else: # Warning about mismatched API, possible non public overloaded version api_diff = '\t\t' + inp_file + " line(" + str(line_num) + ")\n\t\tapi: " + api_types + "\n\t\teta: " + eta_types - message("\t" + api_name + ':\n' + api_diff + '\n') + message("\t" + api_name + ' args mismatch:\n' + api_diff + '\n') # API found action if found == 2: @@ -394,6 +414,10 @@ if (len(sys.argv) > 1) and (sys.argv[1] == '-v'): verbose = 1 sys.argv.pop(1) +if (len(sys.argv) > 1) and (sys.argv[1] == '-e'): + errexit = 1 + sys.argv.pop(1) + if (len(sys.argv) < 3): fatal ("Usage: " + sys.argv[0] + " [-v] \n" + " -v - verbose messages\n" + @@ -441,10 +465,10 @@ if len(opts_map) != 0: args_str = api_map[name]; api_map[name] = list_api_args(args_str) if not name in opts_map: - fatal("not found: " + name) + error("implementation not found: " + name) not_found += 1 if not_found != 0: - fatal(not_found + " API calls not found") + fatal(str(not_found) + " API calls missing in interception layer") # Generating output header file with open(OUTPUT, 'w') as f: