Creation of ReferenceCountedObject will increase reference count by 1.
Clear the commands from Node after capture so that they wont be reference later.
Change-Id: I1cc4085939cf65218ec2aa2e25ab6d737f7cacd3
Capture AQL packets during GraphInstantiation and enqueue AQL packets during graph launch.
Added support to capture single graph memset node.
Capture support for memset node is currently disabled.
Memset capture will be enabled when capture for multiple packets are supported..
Change-Id: I14dfbc41731025cc3a548a730558915def3fa384
Since we made the members public, we can optimize some operations which
do not require redundant conversions to half_raw types.
Change-Id: I31555ef18e695d8e24b89f0418187fa4e932a38a
On some platforms user can ask for extended shared memory for a
particular kernel in some cases. This feature does not exist on HIP at
the moment. So we are setting it to sharedMemPerBlock which is the
maximum user can expect for their kernels.
Change-Id: I81005cf0d1c9fb941e77d34fb8385241ffe5bdd0
For refactoring of childGraph to have its own graphExec,
kernelArgs needs to be separated from the graphExec object.
All the childNodes part of graph should share same kernelArg pool.
Otherwise we endup creating multiple device kernel arg memory chucks
for single graphExec.
Change-Id: I4029a46ebc1fa112d87df64ab1fecbf288fabe5e
In case when the tile size is greater than the number of active threads,
the coalesced group size should be equal to the number of active threads.
Change-Id: I1d41322f2428a07862a590cb5d34b01243383b7c
Remove the redundant copies inside sub folders. This was useful when
these projects were independent but now since they are merged they
should have one single .clang-format file.
Change-Id: I60510d7b78b129c761e84f13403492bd0c5d941a
This change modifies the readback mechanism to use a pointer to volatile
instead of a volatile pointer. This ensures that the compiler does not
optimize away the read operation.
Change-Id: I79ff925d615aa8cc4f950e8ff4b7e608fcb179a4
- Introduce a lock when checking isUserObjectValid. We need a lock
here as one can remove the userObject T2, leading to buffer overflow
when checking ranges in T1.
Change-Id: I058144b8cc463c90ab6bf5cf96bf937897742917
This reverts commit d240b03969.
Reason for revert: <rocm-llvm package name change not required for static builds>
Change-Id: Ib2214a74162e5b015b096dc286151ecbd3ca0a80
If graph has multiple branches, End command is enqueued on launch stream which
makes sure all the internal parallel streams are finsihed.
When node is removed from the graph, indegree and outdegree are not getting update correctly for parent, child nodes and
resulting in endNode not having deps on parallel commands. Resulting in graph sync issues.
Change-Id: I33cc2f21220e1c017d88099b29b542e05b683f73
Resolved an issue where a freed virtual buffer was incorrectly
added to the global mapping causing an assertion error during
teardown process.
Change-Id: I4801157a28603ce9be1ca0131982b700ff884f7a
Changed find_package call to prioritize the package that is
found under the rocm installation over other system locations
Change-Id: Ice93c94bbb9cdebd467d3e88bb2e4bfb7a1e76d9
If the graph has kernels that does device side allocation, during packet capture, heap is
allocated because heap pointer has to be added to the AQL packet, and initialized during
graph launch.
Handle race with wait when 2 kernels with device heap are enqueued on multiple streams.
Change-Id: I45933b77fcaf7bc8fdf1bc906462e32b5d8d3688
The hip libraries depends on multiple standard shared libraries.
Corresponding packages are added to the dependency list.
Removed hipcc from hip-runtime package dependency list
Rearranged the package dependencies as well
The base (docker) images usually comes with the standard packages installed.
So even without the standard packages in dependency list, the installation will go through.
But its good to have all the required packages in the dependency list.
Change-Id: Iacab7993dffaa00efd1344d3eafb55f6317d1509
We more clearly define what happens in the case of amdgpu-arch
failure, and instruct developers on how to workaround the failure
Change-Id: Ifff569a7d688b6545ca457116e534a1830fa03c0
Adding a safety check prevents an invalid memory access
if timestamps and kernelNames vectors are of different size.
The patch also moves the addKernelNames for the accumulate command
into dispatchAqlPacket function.
Change-Id: Iea0927e1253800403a1ae3f3d72de1e7d96476c3
cmake config files were part of hip runtime package rather than dev/devel package
Corrected the same
Change-Id: I5e52658f28c551c830294d3b1525907f1b7cbc50
- Update the intra socket weight for partitions within single socket as
it is changed to 13 by the driver.
- Use the PCIe function to distinguish the partitions of the same device
such as TPX mode in gfx942.
Change-Id: I8e64023d44e37c2dbb105cbb343441a48021ba7b
The warpSize variable is set to the value of the __AMDGCN_WAVEFRONT_SIZE macro,
which is a meaningless default in host code.
The resolution for SWDEV-449015 will introduce diagnostics for uses of this
macro in host code, which includes the current definition of the warpSize
variable. With the __device__ specifier, the definition of the warpSize
variable will not cause these diagnostics.
This change does not stop the variable from being used in host code since clang
intentionally does not diagnose uses of __device__ constexpr variables in host
code.
Change-Id: I0317217affe94fdf2dfd9ad0f134e68f5173245f
The installation of libraries should use the target signature rather than programs.
The TARGETS form will provide the right permission for the libraries
Change-Id: Iaaaeff960a8d85df92f5025e905082642d2225d9
Anusha GodavarthySurya
Jatin Chaudhary
kjayapra-amd
Marko Arandjelovic
Julia Jiang
sdashmiz
Ioannis Assiouras
Jaydeep Patel
Saleel Kudchadker
Ranjith Ramakrishnan
Rahul Manocha
taosang2
Jacob Lambert
Branislav Brzak
Satyanvesh Dittakavi
Fabian Ritter