diff --git a/projects/clr/CHANGELOG.md b/projects/clr/CHANGELOG.md index ec19010a72..6a4fe5852f 100644 --- a/projects/clr/CHANGELOG.md +++ b/projects/clr/CHANGELOG.md @@ -8,6 +8,12 @@ Full documentation for HIP is available at [rocm.docs.amd.com](https://rocm.docs * New HIP APIs - `hipModuleGetFunctionCount` returns the number of functions within a module + - `hipMemsetD2D8` Used for setting 2D memory range with specified 8-bit values + - `hipMemsetD2D8Async` Used for setting 2D memory range with specified 8-bit values asynchronously + - `hipMemsetD2D16` Used for setting 2D memory range with specified 16-bit values + - `hipMemsetD2D16Async` Used for setting 2D memory range with specified 16-bit values asynchronously + - `hipMemsetD2D32` Used for setting 2D memory range with specified 32-bit values + - `hipMemsetD2D32Async` Used for setting 2D memory range with specified 32-bit values asynchronously ## HIP 7.0 for ROCm 7.0 diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp index cb93a220e9..de4273250a 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp @@ -1046,7 +1046,19 @@ typedef hipError_t (*t_hipMemGetHandleForAddressRange)(void* handle, hipDevicept size_t size, hipMemRangeHandleType handleType, unsigned long long flags); - +typedef hipError_t (*t_hipMemsetD2D8)(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, + size_t width, size_t height); +typedef hipError_t (*t_hipMemsetD2D8Async)(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, + size_t width, size_t height, hipStream_t stream); +typedef hipError_t (*t_hipMemsetD2D16)(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, + size_t width, size_t height); +typedef hipError_t (*t_hipMemsetD2D16Async)(hipDeviceptr_t dst, size_t dstPitch, + unsigned short value, size_t width, size_t height, + hipStream_t stream); +typedef hipError_t (*t_hipMemsetD2D32)(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, + size_t width, size_t height); +typedef hipError_t (*t_hipMemsetD2D32Async)(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, + size_t width, size_t height, hipStream_t stream); // HIP Compiler dispatch table struct HipCompilerDispatchTable { // HIP_COMPILER_API_TABLE_STEP_VERSION == 0 @@ -1590,6 +1602,12 @@ struct HipDispatchTable { // HIP_RUNTIME_API_TABLE_STEP_VERSION = 13 t_hipModuleGetFunctionCount hipModuleGetFunctionCount_fn; + t_hipMemsetD2D8 hipMemsetD2D8_fn; + t_hipMemsetD2D8Async hipMemsetD2D8Async_fn; + t_hipMemsetD2D16 hipMemsetD2D16_fn; + t_hipMemsetD2D16Async hipMemsetD2D16Async_fn; + t_hipMemsetD2D32 hipMemsetD2D32_fn; + t_hipMemsetD2D32Async hipMemsetD2D32Async_fn; // HIP_RUNTIME_API_TABLE_STEP_VERSION = 14 // removed HIP_MEMSET_NODE_PARAMS replaced by hipMemsetParams diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index 38d35343d6..4179556414 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -439,7 +439,13 @@ enum hip_api_id_t { HIP_API_ID_hipLaunchKernelExC = 419, HIP_API_ID_hipDrvLaunchKernelEx = 420, HIP_API_ID_hipModuleGetFunctionCount = 421, - HIP_API_ID_LAST = 421, + HIP_API_ID_hipMemsetD2D16 = 422, + HIP_API_ID_hipMemsetD2D16Async = 423, + HIP_API_ID_hipMemsetD2D32 = 424, + HIP_API_ID_hipMemsetD2D32Async = 425, + HIP_API_ID_hipMemsetD2D8 = 426, + HIP_API_ID_hipMemsetD2D8Async = 427, + HIP_API_ID_LAST = 427, HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice), HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties), @@ -802,6 +808,12 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipMemsetAsync: return "hipMemsetAsync"; case HIP_API_ID_hipMemsetD16: return "hipMemsetD16"; case HIP_API_ID_hipMemsetD16Async: return "hipMemsetD16Async"; + case HIP_API_ID_hipMemsetD2D16: return "hipMemsetD2D16"; + case HIP_API_ID_hipMemsetD2D16Async: return "hipMemsetD2D16Async"; + case HIP_API_ID_hipMemsetD2D32: return "hipMemsetD2D32"; + case HIP_API_ID_hipMemsetD2D32Async: return "hipMemsetD2D32Async"; + case HIP_API_ID_hipMemsetD2D8: return "hipMemsetD2D8"; + case HIP_API_ID_hipMemsetD2D8Async: return "hipMemsetD2D8Async"; case HIP_API_ID_hipMemsetD32: return "hipMemsetD32"; case HIP_API_ID_hipMemsetD32Async: return "hipMemsetD32Async"; case HIP_API_ID_hipMemsetD8: return "hipMemsetD8"; @@ -1218,6 +1230,12 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipMemsetAsync", name) == 0) return HIP_API_ID_hipMemsetAsync; if (strcmp("hipMemsetD16", name) == 0) return HIP_API_ID_hipMemsetD16; if (strcmp("hipMemsetD16Async", name) == 0) return HIP_API_ID_hipMemsetD16Async; + if (strcmp("hipMemsetD2D16", name) == 0) return HIP_API_ID_hipMemsetD2D16; + if (strcmp("hipMemsetD2D16Async", name) == 0) return HIP_API_ID_hipMemsetD2D16Async; + if (strcmp("hipMemsetD2D32", name) == 0) return HIP_API_ID_hipMemsetD2D32; + if (strcmp("hipMemsetD2D32Async", name) == 0) return HIP_API_ID_hipMemsetD2D32Async; + if (strcmp("hipMemsetD2D8", name) == 0) return HIP_API_ID_hipMemsetD2D8; + if (strcmp("hipMemsetD2D8Async", name) == 0) return HIP_API_ID_hipMemsetD2D8Async; if (strcmp("hipMemsetD32", name) == 0) return HIP_API_ID_hipMemsetD32; if (strcmp("hipMemsetD32Async", name) == 0) return HIP_API_ID_hipMemsetD32Async; if (strcmp("hipMemsetD8", name) == 0) return HIP_API_ID_hipMemsetD8; @@ -3247,6 +3265,51 @@ typedef struct hip_api_data_s { size_t count; hipStream_t stream; } hipMemsetD16Async; + struct { + hipDeviceptr_t dst; + size_t dstPitch; + unsigned short value; + size_t width; + size_t height; + } hipMemsetD2D16; + struct { + hipDeviceptr_t dst; + size_t dstPitch; + unsigned short value; + size_t width; + size_t height; + hipStream_t stream; + } hipMemsetD2D16Async; + struct { + hipDeviceptr_t dst; + size_t dstPitch; + unsigned int value; + size_t width; + size_t height; + } hipMemsetD2D32; + struct { + hipDeviceptr_t dst; + size_t dstPitch; + unsigned int value; + size_t width; + size_t height; + hipStream_t stream; + } hipMemsetD2D32Async; + struct { + hipDeviceptr_t dst; + size_t dstPitch; + unsigned char value; + size_t width; + size_t height; + } hipMemsetD2D8; + struct { + hipDeviceptr_t dst; + size_t dstPitch; + unsigned char value; + size_t width; + size_t height; + hipStream_t stream; + } hipMemsetD2D8Async; struct { hipDeviceptr_t dest; int value; @@ -5609,16 +5672,16 @@ typedef struct hip_api_data_s { cb_data.args.hipMemcpyHtoAAsync.ByteCount = (size_t)ByteCount; \ cb_data.args.hipMemcpyHtoAAsync.stream = (hipStream_t)stream; \ }; -// hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes')] +// hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes')] #define INIT_hipMemcpyHtoD_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMemcpyHtoD.dst = (hipDeviceptr_t)dstDevice; \ - cb_data.args.hipMemcpyHtoD.src = (void*)srcHost; \ + cb_data.args.hipMemcpyHtoD.src = (const void*)srcHost; \ cb_data.args.hipMemcpyHtoD.sizeBytes = (size_t)ByteCount; \ }; -// hipMemcpyHtoDAsync[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes'), ('hipStream_t', 'stream')] +// hipMemcpyHtoDAsync[('hipDeviceptr_t', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes'), ('hipStream_t', 'stream')] #define INIT_hipMemcpyHtoDAsync_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMemcpyHtoDAsync.dst = (hipDeviceptr_t)dstDevice; \ - cb_data.args.hipMemcpyHtoDAsync.src = (void*)srcHost; \ + cb_data.args.hipMemcpyHtoDAsync.src = (const void*)srcHost; \ cb_data.args.hipMemcpyHtoDAsync.sizeBytes = (size_t)ByteCount; \ cb_data.args.hipMemcpyHtoDAsync.stream = (hipStream_t)stream; \ }; @@ -5738,6 +5801,57 @@ typedef struct hip_api_data_s { cb_data.args.hipMemsetD16Async.count = (size_t)count; \ cb_data.args.hipMemsetD16Async.stream = (hipStream_t)stream; \ }; +// hipMemsetD2D16[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned short', 'value'), ('size_t', 'width'), ('size_t', 'height')] +#define INIT_hipMemsetD2D16_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemsetD2D16.dst = (hipDeviceptr_t)dst; \ + cb_data.args.hipMemsetD2D16.dstPitch = (size_t)dstPitch; \ + cb_data.args.hipMemsetD2D16.value = (unsigned short)value; \ + cb_data.args.hipMemsetD2D16.width = (size_t)width; \ + cb_data.args.hipMemsetD2D16.height = (size_t)height; \ +}; +// hipMemsetD2D16Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned short', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')] +#define INIT_hipMemsetD2D16Async_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemsetD2D16Async.dst = (hipDeviceptr_t)dst; \ + cb_data.args.hipMemsetD2D16Async.dstPitch = (size_t)dstPitch; \ + cb_data.args.hipMemsetD2D16Async.value = (unsigned short)value; \ + cb_data.args.hipMemsetD2D16Async.width = (size_t)width; \ + cb_data.args.hipMemsetD2D16Async.height = (size_t)height; \ + cb_data.args.hipMemsetD2D16Async.stream = (hipStream_t)stream; \ +}; +// hipMemsetD2D32[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned int', 'value'), ('size_t', 'width'), ('size_t', 'height')] +#define INIT_hipMemsetD2D32_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemsetD2D32.dst = (hipDeviceptr_t)dst; \ + cb_data.args.hipMemsetD2D32.dstPitch = (size_t)dstPitch; \ + cb_data.args.hipMemsetD2D32.value = (unsigned int)value; \ + cb_data.args.hipMemsetD2D32.width = (size_t)width; \ + cb_data.args.hipMemsetD2D32.height = (size_t)height; \ +}; +// hipMemsetD2D32Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned int', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')] +#define INIT_hipMemsetD2D32Async_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemsetD2D32Async.dst = (hipDeviceptr_t)dst; \ + cb_data.args.hipMemsetD2D32Async.dstPitch = (size_t)dstPitch; \ + cb_data.args.hipMemsetD2D32Async.value = (unsigned int)value; \ + cb_data.args.hipMemsetD2D32Async.width = (size_t)width; \ + cb_data.args.hipMemsetD2D32Async.height = (size_t)height; \ + cb_data.args.hipMemsetD2D32Async.stream = (hipStream_t)stream; \ +}; +// hipMemsetD2D8[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned char', 'value'), ('size_t', 'width'), ('size_t', 'height')] +#define INIT_hipMemsetD2D8_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemsetD2D8.dst = (hipDeviceptr_t)dst; \ + cb_data.args.hipMemsetD2D8.dstPitch = (size_t)dstPitch; \ + cb_data.args.hipMemsetD2D8.value = (unsigned char)value; \ + cb_data.args.hipMemsetD2D8.width = (size_t)width; \ + cb_data.args.hipMemsetD2D8.height = (size_t)height; \ +}; +// hipMemsetD2D8Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned char', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')] +#define INIT_hipMemsetD2D8Async_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipMemsetD2D8Async.dst = (hipDeviceptr_t)dst; \ + cb_data.args.hipMemsetD2D8Async.dstPitch = (size_t)dstPitch; \ + cb_data.args.hipMemsetD2D8Async.value = (unsigned char)value; \ + cb_data.args.hipMemsetD2D8Async.width = (size_t)width; \ + cb_data.args.hipMemsetD2D8Async.height = (size_t)height; \ + cb_data.args.hipMemsetD2D8Async.stream = (hipStream_t)stream; \ +}; // hipMemsetD32[('hipDeviceptr_t', 'dest'), ('int', 'value'), ('size_t', 'count')] #define INIT_hipMemsetD32_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipMemsetD32.dest = (hipDeviceptr_t)dst; \ @@ -6380,7 +6494,7 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipCtxEnablePeerAccess[('hipCtx_t', 'peerCtx'), ('unsigned int', 'flags')] case HIP_API_ID_hipCtxEnablePeerAccess: break; -// hipCtxGetApiVersion[('hipCtx_t', 'ctx'), ('int*', 'apiVersion')] +// hipCtxGetApiVersion[('hipCtx_t', 'ctx'), ('unsigned int*', 'apiVersion')] case HIP_API_ID_hipCtxGetApiVersion: if (data->args.hipCtxGetApiVersion.apiVersion) data->args.hipCtxGetApiVersion.apiVersion__val = *(data->args.hipCtxGetApiVersion.apiVersion); break; @@ -7521,10 +7635,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipMemcpyHtoAAsync[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('const void*', 'srcHost'), ('size_t', 'ByteCount'), ('hipStream_t', 'stream')] case HIP_API_ID_hipMemcpyHtoAAsync: break; -// hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes')] +// hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes')] case HIP_API_ID_hipMemcpyHtoD: break; -// hipMemcpyHtoDAsync[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes'), ('hipStream_t', 'stream')] +// hipMemcpyHtoDAsync[('hipDeviceptr_t', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes'), ('hipStream_t', 'stream')] case HIP_API_ID_hipMemcpyHtoDAsync: break; // hipMemcpyParam2D[('const hip_Memcpy2D*', 'pCopy')] @@ -7577,6 +7691,24 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipMemsetD16Async[('hipDeviceptr_t', 'dest'), ('unsigned short', 'value'), ('size_t', 'count'), ('hipStream_t', 'stream')] case HIP_API_ID_hipMemsetD16Async: break; +// hipMemsetD2D16[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned short', 'value'), ('size_t', 'width'), ('size_t', 'height')] + case HIP_API_ID_hipMemsetD2D16: + break; +// hipMemsetD2D16Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned short', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')] + case HIP_API_ID_hipMemsetD2D16Async: + break; +// hipMemsetD2D32[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned int', 'value'), ('size_t', 'width'), ('size_t', 'height')] + case HIP_API_ID_hipMemsetD2D32: + break; +// hipMemsetD2D32Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned int', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')] + case HIP_API_ID_hipMemsetD2D32Async: + break; +// hipMemsetD2D8[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned char', 'value'), ('size_t', 'width'), ('size_t', 'height')] + case HIP_API_ID_hipMemsetD2D8: + break; +// hipMemsetD2D8Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned char', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')] + case HIP_API_ID_hipMemsetD2D8Async: + break; // hipMemsetD32[('hipDeviceptr_t', 'dest'), ('int', 'value'), ('size_t', 'count')] case HIP_API_ID_hipMemsetD32: break; @@ -10531,6 +10663,63 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD16Async.stream); oss << ")"; break; + case HIP_API_ID_hipMemsetD2D16: + oss << "hipMemsetD2D16("; + oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16.dst); + oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16.dstPitch); + oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16.value); + oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16.width); + oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16.height); + oss << ")"; + break; + case HIP_API_ID_hipMemsetD2D16Async: + oss << "hipMemsetD2D16Async("; + oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.dst); + oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.dstPitch); + oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.value); + oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.width); + oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.height); + oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.stream); + oss << ")"; + break; + case HIP_API_ID_hipMemsetD2D32: + oss << "hipMemsetD2D32("; + oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32.dst); + oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32.dstPitch); + oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32.value); + oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32.width); + oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32.height); + oss << ")"; + break; + case HIP_API_ID_hipMemsetD2D32Async: + oss << "hipMemsetD2D32Async("; + oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.dst); + oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.dstPitch); + oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.value); + oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.width); + oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.height); + oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.stream); + oss << ")"; + break; + case HIP_API_ID_hipMemsetD2D8: + oss << "hipMemsetD2D8("; + oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8.dst); + oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8.dstPitch); + oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8.value); + oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8.width); + oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8.height); + oss << ")"; + break; + case HIP_API_ID_hipMemsetD2D8Async: + oss << "hipMemsetD2D8Async("; + oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.dst); + oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.dstPitch); + oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.value); + oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.width); + oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.height); + oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.stream); + oss << ")"; + break; case HIP_API_ID_hipMemsetD32: oss << "hipMemsetD32("; oss << "dest="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD32.dest); diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 77b2173a87..a93edcca5d 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -493,3 +493,9 @@ hipLinkDestroy hipLaunchKernelExC hipDrvLaunchKernelEx hipModuleGetFunctionCount +hipMemsetD2D8 +hipMemsetD2D8Async +hipMemsetD2D16 +hipMemsetD2D16Async +hipMemsetD2D32 +hipMemsetD2D32Async diff --git a/projects/clr/hipamd/src/hip_api_trace.cpp b/projects/clr/hipamd/src/hip_api_trace.cpp index edfc688d41..1e4f4ca22d 100644 --- a/projects/clr/hipamd/src/hip_api_trace.cpp +++ b/projects/clr/hipamd/src/hip_api_trace.cpp @@ -831,6 +831,18 @@ hipError_t hipDrvLaunchKernelEx(const HIP_LAUNCH_CONFIG* config, hipFunction_t f hipError_t hipMemGetHandleForAddressRange(void* handle, hipDeviceptr_t dptr, size_t size, hipMemRangeHandleType handleType, unsigned long long flags); +hipError_t hipMemsetD2D8(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width, + size_t height); +hipError_t hipMemsetD2D8Async(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width, + size_t height, hipStream_t stream); +hipError_t hipMemsetD2D16(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, size_t width, + size_t height); +hipError_t hipMemsetD2D16Async(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, + size_t width, size_t height, hipStream_t stream); +hipError_t hipMemsetD2D32(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, size_t width, + size_t height); +hipError_t hipMemsetD2D32Async(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, + size_t width, size_t height, hipStream_t stream); } // namespace hip @@ -1348,6 +1360,12 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) { ptrDispatchTable->hipLaunchKernelExC_fn = hip::hipLaunchKernelExC; ptrDispatchTable->hipDrvLaunchKernelEx_fn = hip::hipDrvLaunchKernelEx; ptrDispatchTable->hipMemGetHandleForAddressRange_fn = hip::hipMemGetHandleForAddressRange; + ptrDispatchTable->hipMemsetD2D8_fn = hip::hipMemsetD2D8; + ptrDispatchTable->hipMemsetD2D8Async_fn = hip::hipMemsetD2D8Async; + ptrDispatchTable->hipMemsetD2D16_fn = hip::hipMemsetD2D16; + ptrDispatchTable->hipMemsetD2D16Async_fn = hip::hipMemsetD2D16Async; + ptrDispatchTable->hipMemsetD2D32_fn = hip::hipMemsetD2D32; + ptrDispatchTable->hipMemsetD2D32Async_fn = hip::hipMemsetD2D32Async; } #if HIP_ROCPROFILER_REGISTER > 0 @@ -1993,13 +2011,20 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipDrvLaunchKernelEx_fn, 475); HIP_ENFORCE_ABI(HipDispatchTable, hipMemGetHandleForAddressRange_fn, 476); // HIP_RUNTIME_API_TABLE_STEP_VERSION == 13 HIP_ENFORCE_ABI(HipDispatchTable, hipModuleGetFunctionCount_fn, 477); +HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D8_fn, 478); +HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D8Async_fn, 479); +HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D16_fn, 480); +HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D16Async_fn, 481); +HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D32_fn, 482); +HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D32Async_fn, 483); + // if HIP_ENFORCE_ABI entries are added for each new function pointer in the table, the number below // will be +1 of the number in the last HIP_ENFORCE_ABI line. E.g.: // // HIP_ENFORCE_ABI(, , 8) // // HIP_ENFORCE_ABI_VERSIONING(
, 9) <- 8 + 1 = 9 -HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 478) +HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 484) static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 14, "If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function " diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index 71ea9938a5..f4ad7dba11 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -609,6 +609,12 @@ local: hip_7.1 { global: hipModuleGetFunctionCount; + hipMemsetD2D8; + hipMemsetD2D8Async; + hipMemsetD2D16; + hipMemsetD2D16Async; + hipMemsetD2D32; + hipMemsetD2D32Async; local: *; -} hip_6.5; +} hip_6.5; \ No newline at end of file diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 280e6252a3..318dfb850e 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -3117,7 +3117,7 @@ hipError_t ihipMemset3DCommand(std::vector &commands, hipPitchedP hipError_t ihipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, - hipStream_t stream, bool isAsync = false) { + hipStream_t stream, bool isAsync = false, size_t elementSize = 1) { auto sizeBytes = extent.width * extent.height * extent.depth; if (sizeBytes == 0) { @@ -3142,7 +3142,7 @@ hipError_t ihipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent } hip::Stream* hip_stream = hip::getStream(stream); std::vector commands; - status = ihipMemset3DCommand(commands, pitchedDevPtr, value, extent, hip_stream); + status = ihipMemset3DCommand(commands, pitchedDevPtr, value, extent, hip_stream, elementSize); if (status != hipSuccess) { return status; } @@ -3157,9 +3157,10 @@ hipError_t ihipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent } hipError_t hipMemset2D_common(void* dst, size_t pitch, int value, size_t width, - size_t height, hipStream_t stream=nullptr) { + size_t height, hipStream_t stream=nullptr, size_t elementSize = 1) { CHECK_STREAM_CAPTURING(); - return ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream); + return ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, false, + elementSize); } hipError_t hipMemset2D_spt(void* dst, size_t pitch, int value, size_t width, size_t height) { @@ -3173,11 +3174,11 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t HIP_RETURN(hipMemset2D_common(dst, pitch, value, width, height)); } -hipError_t hipMemset2DAsync_common(void* dst, size_t pitch, int value, - size_t width, size_t height, hipStream_t stream) { +hipError_t hipMemset2DAsync_common(void* dst, size_t pitch, int value, size_t width, size_t height, + hipStream_t stream, size_t elementSize = 1) { STREAM_CAPTURE(hipMemset2DAsync, stream, dst, pitch, value, width, height); - - return ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, true); + return ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, true, + elementSize); } hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, @@ -3193,6 +3194,48 @@ hipError_t hipMemset2DAsync_spt(void* dst, size_t pitch, int value, HIP_RETURN(hipMemset2DAsync_common(dst, pitch, value, width, height, stream)); } +hipError_t hipMemsetD2D8(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width, + size_t height) { + HIP_INIT_API(hipMemsetD2D8, dst, dstPitch, value, width, height); + HIP_RETURN(hipMemset2D_common(dst, dstPitch, value, width, height, nullptr, + sizeof(unsigned char))); +} + +hipError_t hipMemsetD2D8Async(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width, + size_t height, hipStream_t stream) { + HIP_INIT_API(hipMemsetD2D8Async, dst, dstPitch, value, width, height, stream); + HIP_RETURN(hipMemset2DAsync_common(dst, dstPitch, value, width, height, stream, + sizeof(unsigned char))); +} + +hipError_t hipMemsetD2D16(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, size_t width, + size_t height) { + HIP_INIT_API(hipMemsetD2D16, dst, dstPitch, value, width, height); + HIP_RETURN(hipMemset2D_common(dst, dstPitch, value, width, height, nullptr, + sizeof(unsigned short))); +} + +hipError_t hipMemsetD2D16Async(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, + size_t width, size_t height, hipStream_t stream) { + HIP_INIT_API(hipMemsetD2D16Async, dst, dstPitch, value, width, height, stream); + HIP_RETURN(hipMemset2DAsync_common(dst, dstPitch, value, width, height, stream, + sizeof(unsigned short))); +} + +hipError_t hipMemsetD2D32(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, size_t width, + size_t height) { + HIP_INIT_API(hipMemsetD2D32, dst, dstPitch, value, width, height); + HIP_RETURN(hipMemset2D_common(dst, dstPitch, value, width, height, nullptr, + sizeof(unsigned int))); +} + +hipError_t hipMemsetD2D32Async(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, + size_t width, size_t height, hipStream_t stream) { + HIP_INIT_API(hipMemsetD2D32Async, dst, dstPitch, value, width, height, stream); + HIP_RETURN(hipMemset2DAsync_common(dst, dstPitch, value, width, height, stream, + sizeof(unsigned int))); +} + // ================================================================================================ hipError_t hipMemset3D_common(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream=nullptr) { diff --git a/projects/clr/hipamd/src/hip_table_interface.cpp b/projects/clr/hipamd/src/hip_table_interface.cpp index 2aace90bdf..dd4b1996c6 100644 --- a/projects/clr/hipamd/src/hip_table_interface.cpp +++ b/projects/clr/hipamd/src/hip_table_interface.cpp @@ -1888,4 +1888,31 @@ hipError_t hipMemGetHandleForAddressRange(void* handle, hipDeviceptr_t dptr, siz unsigned long long flags) { return hip::GetHipDispatchTable()->hipMemGetHandleForAddressRange_fn(handle, dptr, size, handleType, flags); +} +hipError_t hipMemsetD2D8(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width, + size_t height) { + return hip::GetHipDispatchTable()->hipMemsetD2D8_fn(dst, dstPitch, value, width, height); +} +hipError_t hipMemsetD2D8Async(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width, + size_t height, hipStream_t stream) { + return hip::GetHipDispatchTable()->hipMemsetD2D8Async_fn(dst, dstPitch, value, width, height, + stream); +} +hipError_t hipMemsetD2D16(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, size_t width, + size_t height) { + return hip::GetHipDispatchTable()->hipMemsetD2D16_fn(dst, dstPitch, value, width, height); +} +hipError_t hipMemsetD2D16Async(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, size_t width, + size_t height, hipStream_t stream) { + return hip::GetHipDispatchTable()->hipMemsetD2D16Async_fn(dst, dstPitch, value, width, height, + stream); +} +hipError_t hipMemsetD2D32(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, size_t width, + size_t height) { + return hip::GetHipDispatchTable()->hipMemsetD2D32_fn(dst, dstPitch, value, width, height); +} +hipError_t hipMemsetD2D32Async(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, size_t width, + size_t height, hipStream_t stream) { + return hip::GetHipDispatchTable()->hipMemsetD2D32Async_fn(dst, dstPitch, value, width, height, + stream); } \ No newline at end of file