Graf commitů

236 Commity

Autor SHA1 Zpráva Datum
Alex Voicu a5ad7bc67b Address post-staging issues in #1809 (#1894)
Fixes SWDEV-223910 and SWDEV-223663

[ROCm/hip commit: d830dad3be]
2020-02-27 16:21:12 +05:30
Alex Voicu 745a66c60e Tweak synchronous memcpy implementation (#1809)
The existing one can have issues on certain systems, therefore this limits use of direct memcpy via largeBAR to sizes where it is unequivocally better.

Also addresses SWDEV-220030 and SWDEV-222237.

[ROCm/hip commit: 9b4f39e1d8]
2020-02-18 20:50:27 +05:30
Rahul Garg f1746197c7 Fix hipMemcpy3D (#1798)
Fixes #1790 and #1791. hipMemcpy3D still requires further refactoring for different input and output combinations.

[ROCm/hip commit: 8c5e5e435b]
2020-02-17 19:35:35 +05:30
ansurya ac41b1d0e0 Fix texture tests, always pass row pitch to HSA API for linear layout images (#1785)
=> New ROCr calculates pitch as per HSA specification and addrlib is used to check whether HW can support that configuration. Hence few texture tests are failing with HSA_EXT_STATUS_ERROR_IMAGE_PITCH_UNSUPPORTED.

=> Determine pitch for linear images and always pass rowpitch to HSA API's.

[ROCm/hip commit: df20c17f12]
2020-01-16 08:54:30 +05:30
satyanveshd 33b5e15e9a hipMemcpy[To/From]Symbol(Async) fixes (#1774)
[ROCm/hip commit: 6b5ea15dfe]
2020-01-07 08:11:53 +05:30
Rahul Garg b6a3c3a95e Fix nested designator warnings (#1768)
[ROCm/hip commit: 785a4440c8]
2020-01-06 12:33:34 +05:30
Evgeny Mankov 6ac4d18753 Merge pull request #1759 from emankov/master
[HIP] Unify hipError_t (Step 2)

[ROCm/hip commit: 0dadb23327]
2019-12-30 19:21:09 +03:00
ansurya 8f0daf4019 Fix texture 3D & 2D layered with N components (#1746)
SWDEV-151670: Issue with 3D texture with 4 components
SWDEV-151671: Issue with 2D layered texture with 4 components

Fixed memcpy when memory is allocated with driver API's.

Github issues: #1755
Fixed 3D default case when array type is not set during memory allocation.

[ROCm/hip commit: 391e99cd12]
2019-12-30 12:41:42 +05:30
Evgeny Mankov abef353b5b [HIP] Clean-up deprecated HIP error codes
hipErrorMemoryAllocation -> hipErrorOutOfMemory
hipErrorInitializationError -> hipErrorNotInitialized
hipErrorMapBufferObjectFailed -> hipErrorMapFailed
hipErrorInvalidResourceHandle -> hipErrorInvalidHandle


[ROCm/hip commit: 4921678b6c]
2019-12-23 17:01:35 +03:00
Alex Voicu 1f5ecc0f6a Fix late-coming issues. (#1724)
Implementation for hipMemcpyWithStream.


[ROCm/hip commit: 75a11330aa]
2019-12-23 19:11:24 +05:30
Rahul Garg ebdd811632 Revert - Changes related to hipMemcpyWithStream (#1718)
Reverting #1673, #1697 and #1707.
Support for hipMemcpyWithStream and memcpy optimizations, will be brought in again once issues seen with these are resolved independently.

[ROCm/hip commit: e53fc316f1]
2019-12-06 09:51:53 +05:30
Rahul Garg fa87abd58a Fix hipMemcpyWithStream (#1707)
* Fix hipMemcpyWithStream
* Add a dtest for the same.


[ROCm/hip commit: 892d2a562e]
2019-12-04 11:56:23 +05:30
Jeff Daily f1544b621d fix hipMempcy precondition, test for zero size first (#1697)
[ROCm/hip commit: 299bbe2325]
2019-11-25 19:03:03 -08:00
Alex Voicu 1baa84eae4 General sync memcpy improvements. Add hipMemcpyWithStream (#1673)
* General sync memcpy improvements. Add `hipMemcpyWithStream`

* Update hip_memory.cpp


[ROCm/hip commit: 5a1f823739]
2019-11-20 21:36:37 +05:30
Jeff Daily d4fe8ff822 General hipMemset improvements (#1495)
* hipMemset et al can use HSA API directly for synchronous cases

* lock and flush stream in hipMemset, hold lock until complete

* move hipMemset async check to front of conditional

* use hsa_amd_memory_fill for additional sync memset cases

code cleanup/review for all memset calls

* Fix inversion of execution mutating value.

* ihipMemsetSync fall back to kernel if HSA memset fails

* Never fallback, never surrender.

* Allow NULL stream.

* Optimise memset kernel. Remove deadwood.

* Update hip_memory.cpp

* Clean up stream logic in sync memset

* Revert "Clean up stream logic in sync memset"

This reverts commit 6117dedf673367f44cc704192573a117a3d92477.


[ROCm/hip commit: e31e0ca12e]
2019-11-07 13:19:54 +05:30
ansurya 92dcba11ac Fixed texture 2D mapping for pitched arrays & 3D Texture read (#1415)
Texture 2D image mapping for pitched arrays:
github issue: Texture Object's Buffer seems to be Misaligned #886
JIRA ticket: SWDEV-199313

SWDEV-151670 : Fixed issue with 3D texture with 4 components
SWDEV-151671 : Issue with 2D layered texture with 4 components

[ROCm/hip commit: e07926ce0f]
2019-11-07 13:17:46 +05:30
Rahul Garg 75cf902cdb Add stream
[ROCm/hip commit: 85d70086cb]
2019-10-31 12:15:56 -04:00
Rahul Garg 73ca647852 Fix HIP init calls in hipMemcpy2DFromArray
[ROCm/hip commit: efe6fa86dc]
2019-10-31 12:15:56 -04:00
Rahul Garg c315da2028 Add hipMemcpy2DfromArray (#1510)
Adds hipMemcpy2DFromArray and hipMemcpy2DFromArrayAsync equivalent to cudaMemcpy2DFromArray and cudaMemcpy2DFromArrayAsync.

[ROCm/hip commit: 356765a223]
2019-10-25 15:43:33 +05:30
Aryan Salmanpour 9ab561dd66 [hip] add support for implicit kernel argument for multi-grid sync (#1456)
* [hip] add support for implicit kernel argument for multi-grid sync

* modified code for calculating the prev_sum

* change the impCoopArg type to size_t

* add memory clean up

* launch init_gws and main kernels into two separate loops


[ROCm/hip commit: 359dc79101]
2019-10-24 17:43:30 +05:30
Vladislav Sytchenko 370d039d24 Update the declarations of hipMemsetD8, hipMemsetD8Async, hipMemsetD16, hipMemsetD16Async. These functions are type aware and take in as their third argument the number of elements in the buffer, not the buffer size. Change the name of this argument from sizeBytes to count to align with the above description.
[ROCm/hip commit: 0b52c1d9d8]
2019-10-15 14:18:42 -04:00
ansurya 66e94b1b1f Added new Memory API's (#1399)
Added new memory API's hipMemAllocPitch, hipMemAllocHost, hipMemsetD16, hipMemsetD16Async, hipMemsetD8Async
Modified to support all scenarios hipMemcpyParam2DAsync, hipMemcpyParam2D.


[ROCm/hip commit: ba9c6e13e4]
2019-10-04 13:36:31 +05:30
satyanveshd bfb64c43a4 Reimplement hipMemGetInfo (#1447)
Addresses SWDEV-136570. hipMemGetInfo changed to compute free memory based on information from kfd instead of relying on hc::am_tracker.


[ROCm/hip commit: 3d661e4706]
2019-10-01 12:40:36 +05:30
satyanveshd fe1fae3ec2 Fix checks in hipMemcpy[D/H]to[D/H]() (#1397)
* SWDEV-202043 hipMemcpyDtoD() issue


[ROCm/hip commit: cfbd0e77fe]
2019-09-16 08:32:38 +00:00
Rahul Garg 60adfab243 Fix typo and colon in comment
[ROCm/hip commit: f0a1f95864]
2019-09-13 03:09:01 +05:30
Rahul Garg e79f2d0ff8 [HACK] Temporary fix for hipFree for hipManagedMalloc
[ROCm/hip commit: fe47b2185c]
2019-09-13 02:10:21 +05:30
Rahul Garg 83d9c2769b Revert "Using HSA API for hipMemsetAsync (#1346)" (#1381)
This reverts commit 7dca0455e9.

[ROCm/hip commit: 6545521d6c]
2019-09-03 05:13:46 +00:00
Rahul Garg 4aa011eec6 Fix memcpy with IPC slowness (#1321)
* Fix memcpy with IPC slowness

* Make early erroneous returns

* Real Clean up

* Real Clean up++


[ROCm/hip commit: 71559200c0]
2019-08-23 09:19:18 +00:00
Jatin Chaudhary 7dca0455e9 Using HSA API for hipMemsetAsync (#1346)
[ROCm/hip commit: ac62d7a5c0]
2019-08-21 10:00:10 +00:00
Rahul Garg a984acf245 Add hipMemcpy3DAsync (#1320)
* Add hipMemcpy3DAsync

* Fix CI build error

* Move back stream resolution to internal function

* Remove stream redefinition and check


[ROCm/hip commit: 2405621f62]
2019-08-16 02:13:16 +00:00
Rahul Garg d42844182c Add hipMemcpyParam2DAsync (#1296)
* Add hipMemcpyParam2DAsync

* Add NVCC path changes

* Clean up

* Fix build issue

* Fix else use in both sync and async apis


[ROCm/hip commit: 45b73e0961]
2019-08-09 11:50:37 +00:00
Jeff Daily 9b44993343 consolidate thread local storage (#915)
* all thread local access now through single struct

* clean up old commented-out code, more use of GET_TLS()

* fewer calls to GET_TLS by passing tls as a funtion argument

* revert unnecessary change to printf

* fix failing tests due to TLS change

* fix merge conflicts in ihipOccupancyMaxActiveBlocksPerMultiprocessor


[ROCm/hip commit: 1eb3dbf065]
2019-08-05 09:51:02 +00:00
Rahul Garg 0009dc1067 Change hipErrorUnknown to hipErrorInvalidValue
[ROCm/hip commit: 483aab031f]
2019-07-31 00:28:30 +05:30
Evgeny Mankov d11063e64c [HIP] Fix segfault on uninitialized struct members in hipArrayCreate and hipArray3DCreate
[ROCm/hip commit: 09162d9a53]
2019-07-12 16:38:26 +03:00
Evgeny Mankov 96801f7b3a [HIP][HIPIFY] Split HIP_ARRAY_DESCRIPTOR struct to HIP_ARRAY_DESCRIPTOR and HIP_ARRAY3D_DESCRIPTOR
[Reason] To be compatible with CUDA [#1133]

Update HIP code, hipify-clang, tests and docs

[TODO] Add support of the corresponding functions on nvcc fallback path


[ROCm/hip commit: c7117df91b]
2019-07-11 14:58:16 +03:00
Jatin Chaudhary c7f8ffe41e Adding bounds check before hipMemset (#1190)
* Adding bounds check in ihipMemset

* Adding ihipMemPtrGetInfo to hipMemPtrGetInfo


[ROCm/hip commit: 5ed16432f8]
2019-07-08 11:00:38 +00:00
Anusha Godavarthy Surya f1d6b56fc4 Added missing NULL checks and corrected API return values as per validation
[ROCm/hip commit: 3d5f6be1c7]
2019-06-27 00:19:05 +05:30
Evgeny Mankov cd309b6638 [HIP][HIPIFY] Make hipMemcpyParam2D coherent with cuMemcpy2D
+ Makes hip_Memcpy2D struct compatible with CUDA_MEMCPY2D struct
+ Add hipMemcpyParam2D support in nvcc fallback path
+ Update hipify-clang, tests and docs accordingly


[ROCm/hip commit: 8f059b0ee9]
2019-05-22 18:31:39 +03:00
Rahul Garg c4567ad01a Add fine grained host memory lock support (#1095)
* Add fine grained host memory lock support

* Fix default flag check


[ROCm/hip commit: aeeab1b23f]
2019-05-13 11:48:26 +05:30
Rahul Garg c01236f679 Add hipMallocManaged default functional support (#1036)
* Add hipMallocManaged default functional support

* Fix build error

* Add dtest


[ROCm/hip commit: 2bc2c46d4d]
2019-04-24 16:50:03 +05:30
Jeff Daily a0172ca884 In hipFree, synchronize owner of memory (#1018)
* In hipFree, if memory is associated with a device, synchronize that device's streams.

This changes the behavior from synchronizing the currently set TLS device.

* All devices sync in hipFree for _appId=-1 case.

* Revert "All devices sync in hipFree for _appId=-1 case."

This reverts commit 1efb34d6a8426661e45bc5f763422a1147aeac10.

* add HIP_SYNC_FREE env var


[ROCm/hip commit: 2b3037a6ea]
2019-04-16 08:35:55 +05:30
Rahul Garg 0eaa29ad06 Handle D2D in memcpy2D
[ROCm/hip commit: 0c55db8552]
2019-03-28 02:21:45 +05:30
Rahul Garg d98d5ca12a Let hipHostMalloc always share/map pinned host ptr
[ROCm/hip commit: f0af073793]
2019-03-26 10:19:13 +05:30
Rahul Garg c6ef785464 Avoid double mapping of devices to hostMalloc buffer
[ROCm/hip commit: 5e917d70f3]
2019-03-25 23:07:05 +05:30
Maneesh Gupta 82fd86e63f Merge pull request #970 from mangupta/swdev-172995
hipExtMallocWithFlags implementation

[ROCm/hip commit: 30b5c02ec4]
2019-03-25 07:46:53 +00:00
Maneesh Gupta 67819c0395 hipExtMallocWithFlags needs hcc workweek 19115 or higher
[ROCm/hip commit: cab119c8b2]
2019-03-25 11:41:20 +05:30
Maneesh Gupta 9ac6005d35 hipExtMallocWithFlags implementation
Change-Id: Iee9e119796472200b2933d5e23be60813f33bc75


[ROCm/hip commit: 73ec5d54b5]
2019-03-19 11:59:22 +05:30
Rahul Garg a3fb908a0a Add 2D fallback to use copy kernel
[ROCm/hip commit: 918d7e3a40]
2019-03-14 13:03:06 +05:30
Alex Voicu 0c16497abd dlopen() fixes (#929)
* Initial attempt to switch over to internally linked state.

* Add missing CMake update.

* hipLaunchKernelGGLImpl must be inline as well. Ensure internal linkage.

* Ensure global retrieval uses internally linked state.

* Hide HC in the implementation. Minimise ADL woes.

* Strange software exists, and must be catered to.

* Use a less spammy mechanism for ensuring internal linkage / non-export.

* Remove leftover internal detail.


[ROCm/hip commit: ea0fcf3e61]
2019-03-06 17:31:44 +05:30
Wen-Heng (Jack) Chung da589e38ed Address code review comments to use hipDeviceptr_t
[ROCm/hip commit: 5cbd28f29b]
2019-03-05 05:51:05 +00:00