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
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
- 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
Add support of HIP_FORCE_DEV_KERNARG under PAL.
Fix persistent memory detection for a resource view.
Change-Id: Ifb7db2db14e0c2205a9661cfa53887ec61ab26a4
- 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
Move context allocation into Device::init() method to simplify the logic and handle
HIP_VISIBLE_DEVICES properly
Change-Id: I0fc6f37c7ae39bedbdad0290295d6794c66d6c54
- Support graph with different types of nodes with single
branch when DEBUG_CLR_GRAPH_PACKET_CAPTURE flag is enabled
Change-Id: I149a8629769cd0d5849ffefb04f1352668a685b6
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
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
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
Pass LUID infromation from PAL to HIP device properties.
NodeMask is calculated based on the order of devices in PAL
Change-Id: Ic53707b68b16f392e5284a0105bf1f3941b46653
If a WGL context wasn't attachhed to OGL, then OGL fails to return
the entry points for AMD extension.
Change-Id: I2bfa5f4f4d038785282e1c9af455e4593e75d3ff
- 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
- 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
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
Some pytorch tests use a tracer plugin and rely on profiling information
to be reported right after hipDeviceSynchronize()
Change-Id: Ib021a1e7b1a30b3c24de72627c471810f7f7878d
Support hipExternalMemoryGetMappedMipmappedArray().
Add ImageExternalBuffer to differiate ImageBuffer.
Currently we only support tiling_optimal mode as
vulkan driver doesn't provide tiling information.
Change-Id: I7e3524cdde53e4df9f728894bcebf4bd3f58d4d9
Add trap handler code into runtime and compile/load during
device initialization. The current interface for trap handler in
PAL is obsolete and the new interface will be provided later.
Change-Id: I1fa702c5d1f2e6731f781369c980d546cf422328
This reverts commit f3dc04a50d.
Reason for revert: There are currently some outstanding issues with the COMPILE_SOURCE_WITH_DEVICE_LIBS Comgr action (https://ontrack-internal.amd.com/browse/SWDEV-386072). Once these LLVM issues have been resovled, we can safely re-apply this patch
Change-Id: I8501967af8496ea50d6e4a97399e45db51bbed1e
- Acquire and Release scopes for AQL need to be balanced. These were
missing on the AMD_OPT_FLUSH=0 path.
Change-Id: Ibf4132eb96800f155d7b664359c790d68a353e60
This is related to SWDEV-410182, but it's not enough to fix it.
Functions from device-libs are precompiled into llvm-ir in a "target agnostic" way
(in reality, it's not 100% target agnostic, which brings us many headaches).
When linking builtins (like device-libs) from the command line, we use the flag
-mlink-builtin-bitcode. The difference between regular linking of bitcode and
this flag is that the later propagates target-specific attributes. If this
attributes are not propagated, we can end up with incosistent target attributes.
Comgr provides the action AMD_COMGR_ACTION_COMPILE_SOURCE_WITH_DEVICE_LIBS_TO_BC
for this exact reason. The old action is currently deprecated and this one should
be used.
Change-Id: I518415214debdf4fedf0b1d81456d6e9fb8a3d19
Vikram
Ioannis Assiouras
Saleel Kudchadker
German
Alex Xie
Anusha GodavarthySurya
taosang2
kjayapra-amd
Jatin Chaudhary
jiabaxie
Ajay
Jaydeep Patel
Julia Jiang
Jacob Lambert
Sourabh Betigeri
Satyanvesh Dittakavi
Juan Manuel Martinez