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 7a52fca62c..40bac15efe 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 @@ -2139,27 +2139,27 @@ typedef struct hip_api_data_s { struct { hipStream_t stream; void* ptr; - int32_t value; + uint32_t value; unsigned int flags; unsigned int mask; } hipStreamWaitValue32; struct { hipStream_t stream; void* ptr; - int64_t value; + uint64_t value; unsigned int flags; uint64_t mask; } hipStreamWaitValue64; struct { hipStream_t stream; void* ptr; - int32_t value; + uint32_t value; unsigned int flags; } hipStreamWriteValue32; struct { hipStream_t stream; void* ptr; - int64_t value; + uint64_t value; unsigned int flags; } hipStreamWriteValue64; struct { @@ -3626,7 +3626,7 @@ typedef struct hip_api_data_s { #define INIT_hipStreamWaitValue32_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipStreamWaitValue32.stream = (hipStream_t)stream; \ cb_data.args.hipStreamWaitValue32.ptr = (void*)ptr; \ - cb_data.args.hipStreamWaitValue32.value = (int32_t)value; \ + cb_data.args.hipStreamWaitValue32.value = (uint32_t)value; \ cb_data.args.hipStreamWaitValue32.flags = (unsigned int)flags; \ cb_data.args.hipStreamWaitValue32.mask = (unsigned int)mask; \ }; @@ -3634,7 +3634,7 @@ typedef struct hip_api_data_s { #define INIT_hipStreamWaitValue64_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipStreamWaitValue64.stream = (hipStream_t)stream; \ cb_data.args.hipStreamWaitValue64.ptr = (void*)ptr; \ - cb_data.args.hipStreamWaitValue64.value = (int64_t)value; \ + cb_data.args.hipStreamWaitValue64.value = (uint64_t)value; \ cb_data.args.hipStreamWaitValue64.flags = (unsigned int)flags; \ cb_data.args.hipStreamWaitValue64.mask = (uint64_t)mask; \ }; @@ -3642,14 +3642,14 @@ typedef struct hip_api_data_s { #define INIT_hipStreamWriteValue32_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipStreamWriteValue32.stream = (hipStream_t)stream; \ cb_data.args.hipStreamWriteValue32.ptr = (void*)ptr; \ - cb_data.args.hipStreamWriteValue32.value = (int32_t)value; \ + cb_data.args.hipStreamWriteValue32.value = (uint32_t)value; \ cb_data.args.hipStreamWriteValue32.flags = (unsigned int)flags; \ }; // hipStreamWriteValue64[('hipStream_t', 'stream'), ('void*', 'ptr'), ('int64_t', 'value'), ('unsigned int', 'flags')] #define INIT_hipStreamWriteValue64_CB_ARGS_DATA(cb_data) { \ cb_data.args.hipStreamWriteValue64.stream = (hipStream_t)stream; \ cb_data.args.hipStreamWriteValue64.ptr = (void*)ptr; \ - cb_data.args.hipStreamWriteValue64.value = (int64_t)value; \ + cb_data.args.hipStreamWriteValue64.value = (uint64_t)value; \ cb_data.args.hipStreamWriteValue64.flags = (unsigned int)flags; \ }; // hipTexRefGetAddress[('hipDeviceptr_t*', 'dev_ptr'), ('const textureReference*', 'texRef')] @@ -4725,16 +4725,16 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { // hipStreamWaitEvent[('hipStream_t', 'stream'), ('hipEvent_t', 'event'), ('unsigned int', 'flags')] case HIP_API_ID_hipStreamWaitEvent: break; -// hipStreamWaitValue32[('hipStream_t', 'stream'), ('void*', 'ptr'), ('int32_t', 'value'), ('unsigned int', 'flags'), ('unsigned int', 'mask')] +// hipStreamWaitValue32[('hipStream_t', 'stream'), ('void*', 'ptr'), ('uint32_t', 'value'), ('unsigned int', 'flags'), ('unsigned int', 'mask')] case HIP_API_ID_hipStreamWaitValue32: break; -// hipStreamWaitValue64[('hipStream_t', 'stream'), ('void*', 'ptr'), ('int64_t', 'value'), ('unsigned int', 'flags'), ('uint64_t', 'mask')] +// hipStreamWaitValue64[('hipStream_t', 'stream'), ('void*', 'ptr'), ('uint64_t', 'value'), ('unsigned int', 'flags'), ('uint64_t', 'mask')] case HIP_API_ID_hipStreamWaitValue64: break; -// hipStreamWriteValue32[('hipStream_t', 'stream'), ('void*', 'ptr'), ('int32_t', 'value'), ('unsigned int', 'flags')] +// hipStreamWriteValue32[('hipStream_t', 'stream'), ('void*', 'ptr'), ('uint32_t', 'value'), ('unsigned int', 'flags')] case HIP_API_ID_hipStreamWriteValue32: break; -// hipStreamWriteValue64[('hipStream_t', 'stream'), ('void*', 'ptr'), ('int64_t', 'value'), ('unsigned int', 'flags')] +// hipStreamWriteValue64[('hipStream_t', 'stream'), ('void*', 'ptr'), ('uint64_t', 'value'), ('unsigned int', 'flags')] case HIP_API_ID_hipStreamWriteValue64: break; // hipTexRefGetAddress[('hipDeviceptr_t*', 'dev_ptr'), ('const textureReference*', 'texRef')] diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index 0b31709ff8..e0e37df664 100755 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -172,10 +172,6 @@ global: hipStreamQuery; hipStreamSynchronize; hipStreamWaitEvent; - hipStreamWaitValue32; - hipStreamWaitValue64; - hipStreamWriteValue32; - hipStreamWriteValue64; __hipPopCallConfiguration; __hipPushCallConfiguration; __hipRegisterFatBinary; @@ -326,6 +322,10 @@ global: hipGraphMemsetNodeGetParams; hipGraphMemsetNodeSetParams; hipGraphAddDependencies; + hipStreamWaitValue32; + hipStreamWaitValue64; + hipStreamWriteValue32; + hipStreamWriteValue64; local: *; -}hip_4.3; +} hip_4.3; diff --git a/projects/clr/hipamd/src/hip_stream_ops.cpp b/projects/clr/hipamd/src/hip_stream_ops.cpp index d982ee128b..88487358d4 100644 --- a/projects/clr/hipamd/src/hip_stream_ops.cpp +++ b/projects/clr/hipamd/src/hip_stream_ops.cpp @@ -23,7 +23,7 @@ #include "platform/command_utils.hpp" hipError_t ihipStreamOperation(hipStream_t stream, cl_command_type cmdType, void* ptr, - int64_t value, uint64_t mask, unsigned int flags, size_t sizeBytes) { + uint64_t value, uint64_t mask, unsigned int flags, size_t sizeBytes) { size_t offset = 0; unsigned int outFlags = 0; @@ -76,7 +76,7 @@ hipError_t ihipStreamOperation(hipStream_t stream, cl_command_type cmdType, void return hipSuccess; } -hipError_t hipStreamWaitValue32(hipStream_t stream, void* ptr, int32_t value, unsigned int flags, +hipError_t hipStreamWaitValue32(hipStream_t stream, void* ptr, uint32_t value, unsigned int flags, uint32_t mask) { HIP_INIT_API(hipStreamWaitValue32, stream, ptr, value, mask, flags); // NOTE: ptr corresponds to a HSA Signal memeory which is 64 bits. @@ -91,7 +91,7 @@ hipError_t hipStreamWaitValue32(hipStream_t stream, void* ptr, int32_t value, un 0)); // sizeBytes un-used for wait, set it to 0 } -hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, int64_t value, unsigned int flags, +hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, uint64_t value, unsigned int flags, uint64_t mask) { HIP_INIT_API(hipStreamWaitValue64, stream, ptr, value, mask, flags); HIP_RETURN_DURATION(ihipStreamOperation( @@ -104,7 +104,7 @@ hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, int64_t value, un 0)); // sizeBytes un-used for wait, set it to 0 } -hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, int32_t value, unsigned int flags) { +hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, uint32_t value, unsigned int flags) { HIP_INIT_API(hipStreamWriteValue32, stream, ptr, value, flags); HIP_RETURN_DURATION(ihipStreamOperation( stream, @@ -116,7 +116,7 @@ hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, int32_t value, u 4)); } -hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, int64_t value, unsigned int flags) { +hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, uint64_t value, unsigned int flags) { HIP_INIT_API(hipStreamWriteValue64, stream, ptr, value, flags); HIP_RETURN_DURATION(ihipStreamOperation( stream,