Mempool may reuse memory without a wait. Hence, the timestamp has
to be preserved and can't be destroyed.
Change-Id: I6f095f44afa69887a4b7aeb3b329804aedd96f3e
For all windows allocation on SVM memory tagged with flag ROCCLR_MEM_INTERPROCESS.
hipHostMalloc validation is based on flag. So remove ROCCLR_MEM_INTERPROCESS before check.
Change-Id: I823bbf228d9a4a9acb4abffc01ac6b3f544c6e12
- Application is passing null for parameter stream in API hipStreamWaitEvent
- When event stream isn't capturing and event is not recorded, causes segfault because we are accessing deviceId() from waitStream
Change-Id: I8b87ffd6f234677f68b66dcb7ef44b2ff04a7c91
=> hipDeviceSynchronize is not allowed during capture.
=> hipEventSynchronize during capture should return hipErrorCapturedEvent error
=> hipEventQuery during capture should return hipErrorCapturedEvent error
hipStreamSynchronize, hipEventSynchronize, hipStreamWaitEvent, hipStreamQuery
For Side Stream(Stream that is not currently under capture):
=> If current thread is capturing in relaxed mode, calls are allowed
=> If any stream in current/concurrent thread is capturing in global mode, calls are not allowed
=> If any stream in current thread is capturing in ThreadLocal mode, calls are not allowed
For Stream that is currently under capture
=> calls are not allowed
=> Any call that is not allowed during capture invalidates the capture sequence
=> It is invalid to call synchronous APIs during capture. Synchronous APIs,
such as hipMemcpy(), enqueue work to the legacy stream and synchronize it before returning.
Change-Id: I201c6e63e1a5d93fd416a3b520264c0fdbe31237
Under ROCr physical allocations don't have initial VA and require extra
flag in ROCclr. Add an option to have a mempool of physical allocations.
Change-Id: I4d062fe0dd8113d4eaf6e8b51749ed56d8701d1e
Added call to hipDeviceSynchronize in __hipUnregisterFatBinary
to ensure that all HSA async signal handlers complete before removeFatBinary
Change-Id: I756fecca1c2a5eae092613d8079de266399e5685
Mempool has capability to track dependency between streams for
faster memory reuse. Enable that capability.
Change-Id: I28266a7e38d0fc4c5d027b9542d3719653840821
- Using runtime unbundler, no any gfx device can load fat binary,
if there is any device without available code object.
- Extract available code object to corresponding gfx devices. So
users can work ROCm with those ready devices without segmentation
fault.
Change-Id: I9f14c65ecebf2d3c4b127a007cb434a3ae98c450
Node can be enabled/disabled only for kernel, memcpy and memset nodes.
If the node is disabled it becomes empty node.
To maintain ordering just enqueue marker with respective node dependencies.
Change-Id: I710f3e88ab4e76c81f6f86a40a7dc61fd4c7e440
- enforce incrementing the table versioning number when a table size changes outside of ifdef for ROCPROFILER_REGISTER
- add new HIP_ENFORCE_ABI entries
- update the HipDispatchTable size and bump HIP_RUNTIME_API_TABLE_STEP_VERSION to 1
- re-enable rocprofiler-register
Change-Id: Ie0cc1d8491c5640056e5dd393ea243e4dce4e8a9
When kernel does device side malloc, initial heap is allocated with __amd_rocclr_initHeap.
During graph launch kernel __amd_rocclr_initHeap is enqueued followed by actual kernel . So kernel will execute after initHeap kernel.
But with graph optimizations during capture initHeap gets enqueued on device null stream and actual kernel on graph launch stream.
So no proper synchronization. Switch to command creation and enqueue during launch for kernel node with hidden heap.
Change-Id: Iaf600251faef9a448853f19429023c118aa760b9
Check the pointer if its present in the arrayset before trying to dereference
it as it can cause access violation if the pointer is allocated using malloc
Change-Id: Ida72b9015dc22269fc1fbe0728e66e3de29fda3d
German Andryeyev
Satyanvesh Dittakavi
Ioannis Assiouras
Anusha GodavarthySurya
kjayapra-amd
Jaydeep Patel
sdashmiz
Sourabh Betigeri
Rakesh Roy
jiabaxie
taosang2
Jatin Chaudhary
ChingShihLi
Rahul Manocha
Jonathan R. Madsen