Recent changes disabled system memory allocation
in the abstraciton layer. That requires memory
allocation/destruction in ROCR. Add destruction logic.
Change-Id: I68fe6b0a620ca743fe5850052ea0efa8bb7931c2
[ROCm/clr commit: a6d480e098]
Extra CPU read back will be performed before every submission to make sure
previous writes over PCIE reached GPU. HDP flush is done by CP.
Change-Id: I402d28ca26c8cee4a3920feb3599af8c285d0889
[ROCm/clr commit: cfc07c88ee]
Persistent memory should use direct access for write map and
indirect for read map.
Change-Id: I9fc84836d60088b24012ed25f7ef8c16e33796a3
[ROCm/clr commit: c8b3253a24]
Original logic didn't use pitch because, abstraction layer had
a sysmem copy without pitch. Since extra sysmem copy was
disabled, the code has to accept pitch values from the app.
Change-Id: Ia9fba7b33ddff4e9109b4e63d0d6afa52f501c8f
[ROCm/clr commit: fb3dfcf889]
The kernel accepts uint64_t, but with 32bit OCL build size_t was 32 bit
Change-Id: I6fe37d2e5e69c7bd62d7b1bd4cace758758b3482
[ROCm/clr commit: b3171d08e6]
The new copy kernel can limit the number of launched workgoups.
It can copy in chunks of 16 bytes or 4 bytes.
Workgoup size is increased to 512 or 1024
Change-Id: Ic3fefa2d5bda6afebd1acc4d41ad310b138af6df
[ROCm/clr commit: ed4e1fec98]
Make hipDeviceAttributeHostNativeAtomicSupported return
right value in Windows
Change-Id: I149a8629769cd0d5848ffefb24f1752678a685b8
[ROCm/clr commit: 1faf1a4cb9]
It seems that due to removal of vdev()->isHandlerPending(),
Marker queued to ensure finish is not enqueued and that cause
hung at waiting event for kernel enqueue command.
Change-Id: I364abb2dcb4897b11a7eb61b5d85013b69292792
[ROCm/clr commit: eecbc2e436]
- Add the new fillBuffer kernel, which allows to launch a limited
number of workgroups for memory fill operation
- Switch fill memory to 16 bytes write by default
- Allow to limit the workgroups with DEBUG_CLR_LIMIT_BLIT_WG
Change-Id: Ibad1822f2d42b2fc71bcfc1917c31409c0623e8e
[ROCm/clr commit: f1dc81f427]
- Report kernel names for optimized graph path
- Refactor code so that we store profiling info in Accumulate command
Change-Id: Ib97735a0239aeb9fc3a50a4bb7126dd0bcadc8af
[ROCm/clr commit: b056686607]
The "long" type size seems to be platform dependent, causing hash value
overflow on implementations where "long" is 4 bytes. This addresses the
scenario.
Change-Id: I4e3c0df457e35b139dcc496d832210ba2cb849ba
[ROCm/clr commit: 1f8ead914a]
For avx build, the start address of values_ buffer in KernelParameters is not
correct as it is computed based on 16-byte alignment.
Change-Id: I3b28ae02d2c9c0517d4a348d95ae8c6721bec83d
[ROCm/clr commit: c4f773db0b]
- Do not use extra barrier to detect graph end. If its a kernel node we
can use a completion signal for the last packet. Saves roughly 6us for
Phantom testcase per graph launch.
Change-Id: I5e0c2479d9964fbeda86ed97533f6718f49a7f91
[ROCm/clr commit: c3bd229f4f]
Add support of HIP_FORCE_DEV_KERNARG under PAL.
Fix persistent memory detection for a resource view.
Change-Id: Ifb7db2db14e0c2205a9661cfa53887ec61ab26a4
[ROCm/clr commit: 5f297d75d9]
- Track all captured commands under a new AccumulateCommand
- Add begin() and end() methods to capture commands
- Explicit TS object now passed to certain methods because
profilingBegin() and profilingEnd() now happen separately and thus can
run into threading issues
Change-Id: I171106bdcad72b057836cb2f3fc398db3533119f
[ROCm/clr commit: 40f41f4d0b]
- Refactor code and cleanup logic for callback saving for event records
Change-Id: I5c56aa8e9c968a5bca70fb07ad1796da318e9e89
[ROCm/clr commit: 1338ff37e8]
Move context allocation into Device::init() method to simplify the logic and handle
HIP_VISIBLE_DEVICES properly
Change-Id: I0fc6f37c7ae39bedbdad0290295d6794c66d6c54
[ROCm/clr commit: a49d633883]
- Support graph with different types of nodes with single
branch when DEBUG_CLR_GRAPH_PACKET_CAPTURE flag is enabled
Change-Id: I149a8629769cd0d5849ffefb04f1352668a685b6
[ROCm/clr commit: 38d2c56784]
- Address corner cases that can arise with the new
hipMemcpyDeviceToDeviceNoCU enum
- Better log
Change-Id: I6035b901f8d616741054b7a5ff4f67956329ac57
[ROCm/clr commit: 5662d4037c]
If we submit a systemScope Barrier, we should reset this flag as there
is no need for dispatch AQL again to flush caches/HDP
Change-Id: I55710feb4ba6650852e785b5cadfa64c6b9ce14e
[ROCm/clr commit: de2b06a7a7]
Add hipMemcpyDeviceToDeviceNoCU to force a non blit copy path. This
helps in cases where an app may determine that CU may be busy and copies
with SDMA may be quicker.
Change-Id: I59b415dd8f6022c244e8d75f265464d5c635df1e
[ROCm/clr commit: f316a30e5d]
Fix wrong logic to get layer index;
Make layered image's layout match cuda spec;
Fix wrong comparision of element size.
Remove amd::BufferRect from ihipMemcpyAtoHCommand()
and ihipMemcpyHtoACommand().
Change-Id: Icc6a4233fbce2e9b2dc6feb79e6bfbd761684c7d
[ROCm/clr commit: 5a0085e516]
Pass LUID infromation from PAL to HIP device properties.
NodeMask is calculated based on the order of devices in PAL
Change-Id: Ic53707b68b16f392e5284a0105bf1f3941b46653
[ROCm/clr commit: 694fe4bedb]
If a WGL context wasn't attachhed to OGL, then OGL fails to return
the entry points for AMD extension.
Change-Id: I2bfa5f4f4d038785282e1c9af455e4593e75d3ff
[ROCm/clr commit: 00b58e2536]
- alias hipGetDeviceProperties to hipGetDevicePropertiesR0600
- alias hipDeviceProp_t to hipDeviceProp_tR0600
- remove gcnArch from new device property struct
- add new requested struct members
Change-Id: If3f5dbef3d608487d9f6f419285f4bf577ea9bf0
[ROCm/clr commit: 2989840511]
- Track last SDMA engine per queue, this results in better scheduling
- Reset last SDMA engine upon batch completion. That ensures we dont get
blocked if the same engine is used by another concurrent copy
Change-Id: Id53111980da7ee41d5c932fb44e4aab5b1e065a3
[ROCm/clr commit: bf8baeecb3]
This header file is used by ROCclr and might be used by HIP too.
ROCclr folder is a right place for this header file.
Change-Id: Ie2307331f9bd445f3c295d2ccbfc4924b0fa9510
[ROCm/clr commit: 124de0c519]
Some pytorch tests use a tracer plugin and rely on profiling information
to be reported right after hipDeviceSynchronize()
Change-Id: Ib021a1e7b1a30b3c24de72627c471810f7f7878d
[ROCm/clr commit: 5438b6362e]
German Andryeyev
Jaydeep Patel
kjayapra-amd
Satyanvesh Dittakavi
Jatin Chaudhary
taosang2
Saleel Kudchadker
Ioannis Assiouras
Vikram
Alex Xie
Anusha GodavarthySurya
jiabaxie