Commit Graph

454 Commits

Author SHA1 Message Date
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
German Andryeyev e4dae85ade SWDEV-290495 - Add HSA_AMD_AGENT_INFO_SVM_DIRECT_HOST_ACCESS query
Change-Id: Ib18a53a9016eb5c5ffd51bf6835cba7299ec8421
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 f4032401cf SWDEV-289378 - Avoid CPU TS overwrite if GPU TS is available
Change-Id: I444ec284669b6a6fff31182e213b47c562c90192
2021-08-22 23:56:07 -07:00
German Andryeyev 4fd506cfe1 SWDEV-289378 - Avoid CPU TS overwrite if GPU TS is available
Change-Id: I2fc1a43c048b4fa1e42f27c4cf5c72fc7ae84e9c
2021-08-22 23:56:07 -07:00
pvellien c11c02f2c7 SWDEV-285333 - Introduce Address sanitizer hostcall service
Change-Id: Id29aacd09d0a9934a027446c57c7095804e1a454
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
Saleel Kudchadker 8e08880cc3 SWDEV-247372 - Add logging for debug
Change-Id: Id5a27034005a7deba37072d8a4c6f250104a96c8
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
Sourabh Betigeri b2a1dc26ba SWDEV-286446 - This patch enables stream operations on vega10, vega20, MI100 and MI200
Change-Id: I6f07036d8ee6e4c6b55196a13288f8107488d824
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 85c70a7495 SWDEV-284671 - Add HW event wait to improve hipDeviceSynchronize
If AMD event contains a reference to a HW event, then runtime
could check/wait for HW event. CPU status update will occur later
after HSA signal callback, but it's not important for the result.

Change-Id: I591391a953bbdba6a25ac07e2cd98aeb17cd4596
2021-08-22 23:56:07 -07:00
German Andryeyev 89b69638d1 SWDEV-240804 - Update ReadMostly attribute
Switch HSA_AMD_SVM_ATTRIB_READ_ONLY to
HSA_AMD_SVM_ATTRIB_READ_MOSTLY to match Cuda. The new attribute
was just exposed in ROCr/KFD.

Change-Id: I2ee522d33c347ba52a4e272d2cd7f67960490cf7
2021-08-22 23:56:07 -07:00
German Andryeyev a9a1e21445 SWDEV-240804 - Add coarse grain memory support
Add an extension to memory advise to disable cache coherency for
better performance

Change-Id: I283703d81d9c36ddfa2c8fffa15eef60e2195056
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
Vladislav Sytchenko 7f0c18457d SWDEV-273235 - Initial support for Windows CMake
This change refactors the current ROCclr cmake build to accomodate a
more modular approach. This allows easier support for multiple compiler
and/or multiple runtime backends.

Currently supported compilers:
    HSAIL - enabled by ROCCLR_ENABLE_HSAIL (defaults to OFF)
    LC    - enabled by ROCCLR_ENABLE_LC    (defaults to ON)

Currently supported runtimes:
    HSA - enabled by ROCCLR_ENABLE_HSA (defaults to ON)
    PAL - enabled by ROCCLR_ENABLE_PAL (defaults to OFF)

Any configuration is supported as long as at least one compiler and one
runtime is enabled.

Since ROCclr clients can configure it differently, one cannot reuse the
same ROCclr build artifacts between different clients. To assure this,
this patch assumes that ROCclr will be built as part of the clients
project.

Change-Id: Id4a5c43634296802b8ae87d1ad5984968391ccaf
2021-08-22 23:56:00 -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
kjayapra-amd 1c49d8816c SWDEV-286346 - Implement Arena Memory Object for externally created memory.
Change-Id: I8530602d89edf83ad367c52167e48a1559ee1e18
2021-05-18 10:59:52 -04:00
Jason Tang ed923eb12e SWDEV-269983 - Re-enable OpenCL Offline Compilation in ROCr path
Change-Id: I160c56a6964219c56c85ebeb5f475be535c39022
2021-05-18 10:46:46 -04:00
German Andryeyev 9b3072ae12 SWDEV-240804 - Switch SVM attribute for the first alloc
Use HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE flag for the initial
allocation instead of HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE_IN_PLACE.

Change-Id: Ia52fe205563df1ea916dc2dc81e749e11c16f83d
2021-05-18 09:32:29 -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
Ravi C Akkenapally 93ae30730a SWDEV-286446 - StreamOperations: Enable for all gfx9
Change-Id: I913a31c8bfdf1031d4cd491f1c689aa4df05c02c
2021-05-14 16:15:27 -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
Jason Tang 211ba25b4e SWDEV-277566 - Separate code object loading from building
Change-Id: I87b8178f55e8ef23762dfe11fab71665ba680f00
2021-05-13 08:41:35 -04:00
Brian Sumner 6d09a83b2d SWDEV-285332 - move common context into parent
Change-Id: I99ceb62ad948e1fa9d1dcaa5ede98626cc95bea7
2021-05-09 09:18:39 -07: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
pghafari 89168d4812 SWDEV-245532 - HIP - Vulkan interop semaphores
linuxpro syntax update

Change-Id: Id77e0d82c504fb7a1bd8edfac7dc962e428c87b4
2021-04-28 13:48:25 -04:00
pghafari 6e69bcef9e SWDEV-245532 - HIP - Vulkan interop semaphores
linuxpro syntax update

Change-Id: I947f24f67ed59c15fe5047348b0c005e19e93544
2021-04-28 07:44:22 -04:00
German Andryeyev ea3dba0832 SWDEV-240804 - Enable HMM build by default
Change-Id: Ia6175dff8eda8c18b7a7bb4ca87a90c1f3e4e6fb
2021-04-26 17:36:53 -04:00
pghafari 582d12b32f SWDEV-245532 - HIP - Vulkan interop semaphores
Change-Id: Ib79328ce4ec2f8ac3aade59fde4fd30e2d6e5cba
2021-04-26 13:20:58 -04:00
Satyanvesh Dittakavi 88fca7bf9e SWDEV-281062 - hipIpcOpenMemHandle should return the base ptr
hipIpcOpenMemHandle should return the device pointer which is
similar to the base ptr of the original allocation even if the offset
to the original pointer is passed to hipIpcGetMemHandle

Change-Id: I99c0553e8c67c15b5fed880b6a4c74bce39c3aee
2021-04-22 02:59:05 -04: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
Vladislav Sytchenko cbeb372e46 SWDEV-280473 - Remove HSAIL support from the ROCm backend
In adition to removing the HSAIL logic from the ROCm backend, guard all
of the HSAIL includes in the common layer behind the WITH_COMPILER_LIB
define. This is to avoid including HSAIL headers when building with
no support for it.

In common logic replace the use of the aclType enum with the new
Program::file_type_t enum. This is essentially a local copy of the HSAIL
enum to avoid including any HSAIL headers.

Change-Id: Ica0651d1b29dfccc255cc584eb82a5cb35e1b520
2021-04-12 14:55:06 -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
Sarbojit Sarkar d8d6c4e55d SWDEV-280089 - nullptr was getting added into MemObj
Change-Id: Ica56126156198140a143182de3d93d86730cadf1
2021-04-06 23:41:50 -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
German Andryeyev 8c513407af SWDEV-278346 - Replace assert with LogWarning()
ROCr returns some unexpected values and for now just return
invalid device back to the app

Change-Id: Ia6fb709ed2dc77d77bd5d64de80b41b3720bed76
2021-03-25 10:19:33 -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
Vladislav Sytchenko 99e8ac55cd SWDEV-232428 - PAL Hostcall support
Since the majority of the Hostcall implementation now sits in the
commmon layer, the PAL backend simply just needs to invoke it. One thing
that is missing though is HSA signal support.

The newly added pal::Signal class is a light emulaion of what HSA
signals provide. The current implementation is just enough to get
Hostcall working, but it can be expanded in the future if needed to
fully emulate HSA signals.

The major difference for now between PAL and ROCm hostcall
implemenations is that PAL doesn't support blocking signals. This will
be enabled in the near future. For now use active wait for PAL.

Change-Id: I746557354ab9d71a7d4a31f9320fcc2fee5aee7f
2021-03-09 10:44:34 -05:00
German Andryeyev e9c484d1ce SWDEV-274145 - process ACCESS_QUERY for all devices
HIP requires to return AccessedBy query for all device, but ROCr
can process one per query. Hence send the queries for all
available devices and then accumulate the results in runtime.

Change-Id: I082f9adb8e31c775a8ad1bf7a5af37440ef4bd16
2021-03-08 14:19:13 -05:00