AI/rocm-systems
2
0
Fork 0
Code Issues Pull-Requests Actions Pakete Projekte Releases Wiki Aktivität

SWDEV-546328 - API definitions for hipMemsetD2D8/D16/D32 (#789)

Quellcode durchsuchen 
Co-authored-by: Rahul Manocha <rmanocha@amd.com>

[ROCm/clr commit: 5f86622adc]
Dieser Commit ist enthalten in:
Manocha, Rahul
2025-08-14 14:47:28 -07:00
committet von GitHub
Ursprung 208d124f54
Commit 2b79ca9435
8 geänderte Dateien mit 339 neuen und 19 gelöschten Zeilen
Patch-Datei herunterladen Vergleichs-Datei herunterladen Alle Dateien ausklappen Alle Dateien einklappen
projects/clr/CHANGELOG.md
+6
Datei anzeigen
@@ -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
projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp
+19 -1
Datei anzeigen
@@ -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
projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h
+197 -8
Datei anzeigen
@@ -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);
projects/clr/hipamd/src/amdhip.def
+6
Datei anzeigen
@@ -493,3 +493,9 @@ hipLinkDestroy
hipLaunchKernelExC
hipDrvLaunchKernelEx
hipModuleGetFunctionCount
hipMemsetD2D8
hipMemsetD2D8Async
hipMemsetD2D16
hipMemsetD2D16Async
hipMemsetD2D32
hipMemsetD2D32Async
projects/clr/hipamd/src/hip_api_trace.cpp
+26 -1
Datei anzeigen
@@ -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(<table>, <functor>, 8)
//
// HIP_ENFORCE_ABI_VERSIONING(<table>, 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 "
projects/clr/hipamd/src/hip_hcc.map.in
+7 -1
Datei anzeigen
@@ -609,6 +609,12 @@ local:
hip_7.1 {
global:
hipModuleGetFunctionCount;
hipMemsetD2D8;
hipMemsetD2D8Async;
hipMemsetD2D16;
hipMemsetD2D16Async;
hipMemsetD2D32;
hipMemsetD2D32Async;
local:
*;
} hip_6.5;
} hip_6.5;
projects/clr/hipamd/src/hip_memory.cpp
+51 -8
Datei anzeigen
@@ -3117,7 +3117,7 @@ hipError_t ihipMemset3DCommand(std::vector<amd::Command*> &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<amd::Command*> 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) {
projects/clr/hipamd/src/hip_table_interface.cpp
+27
Datei anzeigen
@@ -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);
}