Commit Graph

98 Commits

Author SHA1 Message Date
German Andryeyev 8e4101b4fd SWDEV-294669 - Avoid queue drain
Use slot wait logic for direct dispatch

Change-Id: I431ba1418eb4aa066b9881934f4055b3d338ce3a
2021-11-18 13:06:12 -05:00
kjayapra-amd 7e32d6d909 SWDEV-309657 - Align Virtual queue size to sizeof(uint64_t).
Change-Id: Ia55d7316693bd13938875ce53f7849d5eb658e8c
2021-11-12 10:35:36 -05:00
German Andryeyev 7e12cf6318 SWDEV-257789 - Initial change to skip kernel arg copy
The optimization is controlled with ROCR_SKIP_KERNEL_ARG_COPY.
This is initial check-in for experiments. Extra changes are
necessary for full support:
- handle graph capture with the original sysmem alloc
- avoid memobject references, otherwise there is a race condition with
reusage of the arg buffer
- Remove arg setup from hip

Change-Id: Ib0af710f93e79834711fa4049a7c66093711e68b
2021-10-28 20:35:35 -04:00
German Andryeyev f78b3a8919 SWDEV-303567 - Add chunks for the pool of kernel arguments
The kernel arg pool will be divided into 8 chunks to avoid long stalls,
when the pool will be reused.

Change-Id: I228e6ca1c09e428c1775f1e5b685220a9a5d71af
2021-10-26 16:31:37 -04:00
Sarbojit Sarkar c06c9f7b93 SWDEV-306302 - Fix for OCLCreateImage test failure
Change-Id: I781504bd1ff599ed75c5ea730be03b71f69761b2
2021-10-07 19:52:58 +00:00
German Andryeyev 7fe696b6ef SWDEV-303567 - Increase the size of AQL queue
ROC_AQL_QUEUE_SIZE will control the size of AQL queue.
The current sefault value is 4096.

Change-Id: Icd2a4ee3ba554c06aa05b08defd922d2c63e43fd
2021-10-06 08:27:36 -04:00
Sarbojit Sarkar 22a847f3ce SWDEV-301823 - Optimize hipMemset2D/3D
Change-Id: Ibe560149a263c2ac6b08e4eb1a1d331d2aeac78c
2021-09-27 14:10:06 -04:00
Sourabh cbb8d82bdb SWDEV-292525 - [vdi] Path to streamOps shaders
Implementation to use a blit kernel to perform
a hipStreamWait/write instead of an AQL packet.

Change-Id: I462671ed5cec37144dfe97ff66439249196117c1
2021-09-27 13:59:35 -04:00
German Andryeyev 65ddfcc6a8 SWDEV-294669 - Keep one more slot for HW processing
The original logic left only one slot for HW processing in the queue.
For some reason there is a race condition on CPU overwrite of the slot
before the current active. The workaround is to avoid the previous to
the current active slot for possible unfinished HW processing.

Change-Id: I565495a8feeaedffc9fc8a505edbee5ff5816975
2021-09-13 13:56:05 -04:00
Jason Tang 73967c3b17 SWDEV-1 - Some 'delete' clean up
Change-Id: I02564f0f0e349375bde1471e9f82df268703367b
2021-09-09 12:12:40 -04:00
Sarbojit Sarkar 42d33029dc SWDEV-300655 - Added thread ID to hip trace
Change-Id: I9234d4ec93e7687cd0a5d1bd930bd4f80936311b
2021-09-06 00:22:42 -04:00
Satyanvesh Dittakavi 169cc857fd SWDEV-298985 - hipMemPrefetchAsync should prefetch the data to the specified destination device
Pass the device agent specified by the user to the ROCr api instead of passing the device agent attached to the specified stream

Change-Id: I86c98935b9dc404eaa6d47ccdd082a8c3678fb36
2021-08-27 05:12:07 -04:00
Saleel Kudchadker 75fea4dca6 SWDEV-297448 - Improve logging
Print non pointer kernel args
Change-Id: Ice0dbc894aae1430ac085df319f4b91dfa21665a
2021-08-25 15:46:06 -07:00
Saleel Kudchadker 24442be35a SWDEV-297448 - Refactor streamOperations code
Reuse FillMemory function, that should fix the cache syncs from the host

Change-Id: Ieebec5fc3ed3a322b88d5187c8dca4805ec6f84b
2021-08-25 12:33:02 -04:00
German Andryeyev ec89348291 SWDEV-296329 - Add lock protection for Timestamp update
There is a possible race condition when signal reuse can have
access to a destroyed Timestamp object, because the callback
was running asynchronously. Use reference counter and lock
to allow asynchronous timestamp update

Change-Id: I6224f7c62cb0a03a7466fcc512e5e5afb06736fa
2021-08-22 23:56:08 -07:00
German Andryeyev 2babcfbdbb SWDEV-290384 - Enable active wait on CPU if HIP requested
Change-Id: Idea5adf7a4705cb999da6785e6229fe3200dce17
2021-08-22 23:56:08 -07:00
German Andryeyev c144383971 SWDEV-290384 - Disable HSA callback for any host wait
Change-Id: Ie876deb62859f5551f4ed69eb8187ac3fa35f42a
2021-08-22 23:56:08 -07:00
German Andryeyev ff15c0893e SWDEV-292018 - Switch to internal signals for markers
Add ref counting to ProfilingSignal class to track the last release.
If a signal was used in the marker, then don't reuse it,
but create a new one for internal usage.
Don't rely on HSA callback for the command status update if there
are no pending dispatches.

Change-Id: I19f14ed9d80acfe79993b343b2187635f8428a20
2021-08-22 23:56:07 -07:00
Sourabh Betigeri 6408b9b906 SWDEV-290685 - Relacing release fence instead of acquire for hipStreamWrite()
Change-Id: Ic2946b68c427d3e058948c0813863a27c21b903d
2021-08-22 23:56:07 -07:00
Sourabh Betigeri e4ddb5e08a SWDEV-290685 - Dispatch a barrier packet with acquire system scope to ensure ordering before a hipStreamWrite()
Change-Id: I8853ad86a6634d55a98173ca3f79d93b85c08f85
2021-08-22 23:56:07 -07:00
agunashe d96481fb36 SWDEV-293742 - Update copyright end year VDI repo
Change-Id: I69d2fea4a7a43adf96ccea794270e4af991c5261
2021-08-22 23:56:07 -07:00
Saleel Kudchadker cd21af757e SWDEV-260448 - Honor NUMACTL for Direct Dispatch
Setting AMD_CPU_AFFINITY=1 will keep Async Handler thread within the
bounds set by numactl.

Change-Id: Id01b30df5127d65c29ac072bf74a04986b7128de
2021-08-22 23:56:07 -07:00
German Andryeyev ce8dad2ecc SWDEV-290160 - Switch to global HSA signals
Runtime can't assign internal HSA signals for HIP events, because
HIP application can destroy the HIP stream or signal reuse may
occur internally. Switch to global HSA signals for HIP events.

Change-Id: Ieaea2d6b039e492b2e7c5112782a8f4e601e50a1
2021-08-22 23:56:07 -07:00
German Andryeyev c49f1069ab SWDEV-290160 - Don't send notification for batch markers
Batch marker already has a barrier with HSA signal callback

Change-Id: I69fc63d72320c2e9cc2d2e59ebd3f07c0bd0e3b5
2021-08-22 23:56:07 -07:00
Saleel Kudchadker d3213eca90 SWDEV-247372 - Reset hasPendingDispatch
Reset hasPendingDispatch_ if we insert barrier for time marker.

Change-Id: Id038fd4e1c910c0a657978fee00630e49c372321
2021-08-22 23:56:07 -07:00
German Andryeyev a1629cad26 SWDEV-290371 - Add lock protection for signal
Add lock protection for signal processing
If signal is reused, then disable reference to it from HIP
Increase the pool signal size to 32

Change-Id: I7d529b35910f83ce577c9eca6d3386759611ccc0
2021-08-22 23:56:07 -07:00
Saleel Kudchadker 0af6ba9428 SWDEV-286092 - Use Barrier Header for event
Change-Id: I9701fbab587e2ea31e58449e8c8b07341a7aa161
2021-08-22 23:56:07 -07:00
Saleel Kudchadker b416ad7b9d SWDEV-247372 - Active wait timeout env var
- Create an env var ROC_ACTIVE_WAIT_TIMEOUT to set active wait timeout
- Record profiling informaion if marker_ts_ property is valid.

Change-Id: If0d8aec8d9b0715027cf0f7c3dc8a4c722a6bae6
2021-08-22 23:56:07 -07:00
German Andryeyev d93df7037c SWDEV-287137 - Add blocking signal logic
With HIP API callback runtime has to stall the queue until the
callback is done. Rocclr will introduce SW blocking HSA signal,
which will be released after the callback is done.

Change-Id: I6411f3efab31b468e3b87ebb5c8d155e116b613d
2021-08-22 23:56:07 -07:00
German Andryeyev 148a5cac39 SWDEV-287137 - Fix regression with dependent signals
- Make sure barrier with dependent signals issues before queue
index reservation
- Don't issue extra barrier if it's already a barrier command
with dependent signals

Change-Id: I179a8b7adac79eed698f4a4d9eca2606d8e913aa
2021-05-26 12:36:56 -04:00
German Andryeyev fa2e154a8b SWDEV-278894 - Use GPU waits for HIP events
Save HW events in amd::Event.
Use HW events for synchronization

Change-Id: I98cf9c2d0ec3c7fcaf254b749ac6c568d7270ae0
2021-05-25 13:41:15 -04:00
German Andryeyev 3f7a6b01e3 SWDEV-240804 - Remove AMD_HMM_SUPPORT define
Use dynamic logic for HMM based on it's availability

Change-Id: I63751d94571d5af6eb57bef2cb0e071120bfa103
2021-05-14 17:41:06 -04:00
Saleel Kudchadker e5e635f9bf SWDEV-283726 - Workaround for rocprof hang
This addresses the rocprof hang seen with direct dispatch. The
workaround queues the handler back if any of the signal value in the batch
is not decremented. To rememmber the last position in the list, we save
the parsed command in the current timestamp struct.

Change-Id: I02959e463cfe3cee83c54808ffd6e6f48f43b4e8
2021-05-13 11:58:03 -04:00
Saleel Kudchadker 42b8236f93 SWDEV-280773 - Additional logging for signals
Cleanup new lines in debug log

Change-Id: I6862c332eb9457b51e23cf4e9db9ba3f870d0c39
2021-04-30 15:05:57 -07:00
German Andryeyev ca2ea70a6c SWDEV-278896 - Increase thresholds to match MT behavior
MT doesn't use GPU waits, but CPU for sync between engines.
Change the threshold values for CPU waits for direct dispatch.
That will bring behavior closer to MT.

Change-Id: Ia41c3cb812614962aff2746b6cf858f1bf77dda2
2021-04-16 17:47:57 -04:00
Saleel Kudchadker aa38af8c96 SWDEV-276120 - Remove support for barrier sync
ROC_BARRIER_SYNC will not work with direct dispatch.
Remove and cleanup.

Change-Id: I81368b2e65039477bd0343bb92708dab48867db6
2021-04-07 17:08:39 -04:00
German Andryeyev 2d41031aec SWDEV-279288 - Move the batch condition
The check has to be performed inside the signal loop, because
active signals need to be processed to avoid a stale timestamp
class.

Change-Id: I26af8287aae18eb19c096d9358cd0b86cfd1c561
2021-04-06 10:22:31 -04:00
German Andryeyev a71f7f931e SWDEV-279288 - Avoid profiling info for the sync barrier
- With direct disaptch profiling state is enabled to trigger the
callback on HSA signal. However ROCr has very low peformance on
the first call to get the profiling info. That impacts some tiny
performance tests.

Change-Id: Idacd1b10a473fcfb5feef3074b7191d35743f769
2021-04-05 12:54:23 -04:00
Alex Xie 2bd1836049 SWDEV-276304 - [Navi][OpenCL][Linux]AMF Converter generate corrupted Frames
The copy image workaround could be recursively used by ROCclr blit kernel.
Avoid such situation by using stack variable.

Change-Id: Iadaa8cad9216220194760dd461a9533bb236aea0
2021-03-30 12:07:00 -04:00
Saleel Kudchadker d034c48405 SWDEV-271010 - Increase active wait time
Increase wait time for active wait to 50us

Change-Id: I8f269ab25ecc6775e655b9eb36fafc5f41a59c95
2021-03-11 13:16:59 -08:00
German Andryeyev 7f32d0b425 SWDEV-272496 - Detect callbacks and force AQL barrier
HIP tests require HIP callbacks to be processed in another thread.
This change will use a thread from HSA signal callbacks to make
sure a HIP callback was done asynchronously.
Also process the callback before changing the status of command

Change-Id: Icef85d0e0f808663882cf6881ff1be3e5eca29ac
2021-03-05 11:33:51 -05:00
German Andryeyev 24299e25bd SWDEV-272496 - Fix multiple timing issues
- Don't notify if the batch is empty, because that means
the current command was processed already.
- Disable pinning optimization to avoid a race condition on stall.
- TS marker submition requires extra AQL barrier
to track the status.

Change-Id: I17eff4ad12ac66cfe1bb44048bebb1891805279d
2021-03-01 12:46:57 -05:00
German Andryeyev ac387f9b03 SWDEV-272496 - Wait on CPU before switching to GPU wait
GPU waits have noticeable overheads on compute with extra
AQL barrier packet and on SDMA with power saving features. This
change introduces a wait on CPU for 30 us in case the app has tiny
operations.

Change-Id: I761ba3af595f3f48544980058a9077dda15aa5f9
2021-02-17 15:19:17 -05:00
Vladislav Sytchenko 671778bdd3 SWDEV-232428 - Push hostcall implementation to the device layer
This change unifies the hostcall implementation for all the backends,
by pushing the common logic to the device layer. This is done by
replacing the use of hsa_signal_t with device::Signal (a light wrapper
around it).

Change-Id: I7b6fca7930b5a0b199da5d85e2e048354cc04e7b
2021-02-16 17:19:57 -05:00
Ravi C Akkenapally 0a5f9a3b10 SWDEV-179105 - Stream Operations: Add support for Wait and Write
Change-Id: Ibffa1d6d573826b64763da280074a77271d66808
2021-02-15 17:02:38 -08:00
kjayapra-amd 2df099df9e SWDEV-270013 - Allocate kernel_arguments from kern_arg & finegrain pool instead of coarse grain.
Change-Id: Id4c6977934fdd6ef2311f6e75593801f1e51983c
2021-02-15 18:20:08 -05:00
Payam a2e0b0495c SWDEV-257937 - Updated fix for ROC_BARRIER_SYNC=0
Change-Id: I7e28e541b654db57fb0890d7dbb7519cfb2d93db
2021-02-11 14:01:45 -05:00
German Andryeyev 6966d8098e SWDEV-269654 - Fix HIP stream busy query
- Avoid GPU wait on the marker submission and update the command
batch after HSA signal callback upon HSA barrier completion.

Change-Id: I5c1c97212aefc2ae4b99aa9e2a81627ee9a38c1c
2021-02-09 12:57:12 -05:00
Alex Xie 5330679473 SWDEV-268914 - Linux, AMF, tests fail to scale images down
In pitch workaround, we need to copy the image to copy buffer
when application wants to read image to buffer. After this
patch, we unconditionally copy the image data to the copy buffer.

Change-Id: I71b0d19459542dfbb3ca51a2c8a3a81367fa2fb5
2021-02-01 15:22:15 -05:00
Joseph Greathouse 54d1d69c0a SWDEV-270773 - Fix workgroup calculation logic for GWS initialization
The existing workgroup calculation logic for GWS initialization is
incorrect. It tries to add together workgroups across dimensions,
leading to major under-count in 2D and 3D kernels. An (x,y,z) kernel
uses x * y * z blocks, not x + y + z.

In addition, the previous logic was incorrect for the case of launching
a single-threaded kernel. It calculated 0 workgroups, leading to
initializing GWS to -1.

Change-Id: I1bb20a0d5b6e0cc10ac55901c28d8f93aac61c09
2021-01-31 01:16:48 -06:00