Support gfx9-4-generic target to cover mi3XX.
Support features sramecc and xnack in generic target.
Improve some code formats.
Add more log on compiler.
Change-Id: I6b3c6af55c60cffd43ce6f17b75998f751b75713
Some libs use __HIP_NO_HALF_OPERATORS__ and __HIP_NO_HALF_CONVERSIONS__
which results in operators being hidden and can cause errors.
Change-Id: I83c194d7d727cba30b46d7c296f7d396549f5fca
On Windows, hipHostRegister may add a single object in the MemObjMap
that maps to memory that is allocated on different devices.
This change ensures that the offset that is returned from
getMemoryObject() is computed relative to the memory that is allocated
on the current device.
Change-Id: I5fd3af200bf6f4926fdeaea12dcb9d0154d3a843
- Header files inside rocclr/utils when included from hipamd or opencl should be included as #include "rocclr/utils/xxx.h" instead of "utils/xxx.h"
Change-Id: Ic0760c33b9d091f5620dec67e5482c9698d22093
Some functions were __device__ only, but should be __host__ and
__device__, changed them to __HOST_DEVICE__.
Some functions were __HOST_DEVICE__ but were using ockl functions,
changed them to __device__ only.
Change-Id: Ife9e7abe60415bda68f5f9a101e6e7c39ad51064
The hipGraph will use VMM by default when allocating memory.
However, the handle of Phy mem has been added to Memobj by default.
Since the Memobj will track the whole address range from handle to
handle + size, this needs the system to reserve the whole address
range. If the system range have not reserved by the system, then it
will have the potential issue that clr finds the Memobj incorrectly.
This patch removes the handle from the Memobj to fix this potential
issue.
Change-Id: I2da38e6b2d11d0d48e1afe66c46899500c290624
This change replaces some asserts, that were only available in debug
mode, with standard error handling.
Signed-off-by: Sebastian Luzynski <Sebastian.Luzynski@amd.com>
Change-Id: I112f9e56f921abd72daf0d11e4ecdcb7b1a9f9e6
=> If null stream is not created during sync skip nullstrm creation
=> Do cpu wait on blocking & null stream if it exists
Change-Id: I90d6ced6a2dd1782ba58f3fed4e3608fc0efa55a
hipMallocAsync/hipFreeAsync APIs should return error stating
operation is not supported, if a stream is actively capturing
and is different from the passed stream
Change-Id: I2a1b8260c5eb22d99a936ac529d6788a83f81a17
1) for case where kernelParam_.func = nullptr, the validation
fails in setParam call and memory is not alloced for kernelParams
2) destructor path segfaults trying to free the kernelParams memory.
3) copy of params is done after function validation is successful.
Change-Id: I6338e0c89f259632e4115f0508e2f240bc207fd9
This reverts commit efce2f77c4.
Reason for revert: Even though this change is valid, this would break backward compatibility.
Change-Id: I9c7cab83198c8d5c8485b11194099162e3e7a874
1) For Dynamic CO variables, free the device pointer in
DynCO destructor instead of DeviceVar destructor.
2) For Static CO Remove Fatbinary,
only call hipFree for valid device Vars instead of all devices.
Change-Id: I84291f5371b2c05d1d0bcdb4f9c6bd122e7c9b21
Modified hipFuncSetAttribute to handle pointers to dynamic functions
returned by hipModuleGetFunction.
Change-Id: I54b98f9d31a79630dd7edcd363fad81f1d89219b
This reverts commit 68e5aeb93d.
Reason for revert: Breaking change that will be merged in at a later date
Change-Id: Idd300492cc08a57c50decc22df287ddcc5463c88
Anusha GodavarthySurya
Istvan Kiss
Sourabh Betigeri
Tao Sang
Saleel Kudchadker
Jaydeep Patel
Maneesh Gupta
Jatin Chaudhary
Todd tiantuo Li
Branislav Brzak
Ioannis Assiouras
Marko Arandjelovic
Shane Xiao
German Andryeyev
Sameer Sahasrabuddhe
Vladana Stojiljkovic
Sebastian Luzynski
Marko Arandjelovic
Satyanvesh Dittakavi
Aidan Belton-Schure
Rahul Manocha
kjayapra-amd
Julia Jiang