AI/rocm-systems
2
0 Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
3,050 Commits 3 Branches 0 Tags
5647e0ae3561208cceabc6e4f96ebc3bab94708e
Commit Graph
847 Commits
Author SHA1 Message Date
Alex Voicu 5647e0ae35 Missing handling nullary __global__ functions for mixed arity cases. 
[ROCm/clr commit: 1e64813716]
 2018-11-12 01:51:59 +00:00
Alex Voicu 0178ad536a Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_get_alignment_and_size_from_metadata 
[ROCm/clr commit: 01ec23538e]
 2018-11-11 22:37:09 +00:00
Yaxun Sam Liu 5d54fe864f Add more checks for fatbin 
[ROCm/clr commit: 80b382a37a]
 2018-11-08 16:46:01 -05:00
Alex Voicu c4acf6e8cc Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_get_alignment_and_size_from_metadata 
[ROCm/clr commit: 392e2ba1ba]
 2018-11-08 06:53:06 -08:00
Maneesh Gupta a3758b506d Merge pull request #730 from yxsamliu/dbg-fatbin 
Add more checks and debug outputs for fat binary

[ROCm/clr commit: e510be9395]
 2018-11-08 11:44:27 +05:30
Maneesh Gupta 387e6de9e6 Merge pull request #721 from fronteer/my-fix 
Make correct checking of the returned hipDeviceptr_t from read_global…

[ROCm/clr commit: 0bafc38135]
 2018-11-08 11:42:08 +05:30
Maneesh Gupta 46bf341c12 Merge pull request #742 from gargrahul/fix_module_global_tracking 
Fixes global symbols tracking in hip_module

[ROCm/clr commit: 8b1544c911]
 2018-11-08 11:41:35 +05:30
Alex Voicu e3c9869fbe Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_get_alignment_and_size_from_metadata 
[ROCm/clr commit: 1c525d7222]
 2018-11-01 23:17:27 +00:00
Siu Chi Chan e26a728173 Move the global arrays for hip malloc/free 
from a header into a source file such that
there's only an unique copy in an executable
and prevent wasting static memory on the host

Change-Id: Id5b62766f77809c8d7b47892cb7149c490dcbdb9


[ROCm/clr commit: cdd0109e70]
 2018-11-01 16:20:35 -04:00
Maneesh Gupta 7052a66935 Merge pull request #733 from ex-rzr/fix_array_alloc_size 
Fix allocation size of arrays with multiple and/or non-32-bit channels

[ROCm/clr commit: a75d4c0ecc]
 2018-10-31 14:08:58 +05:30
Maneesh Gupta 130edf2c43 Merge pull request #714 from aaronenyeshi/remove-static-hip-device 
Remove libhip_device.a static library

[ROCm/clr commit: c7f2198c6e]
 2018-10-31 14:08:29 +05:30
Maneesh Gupta 58aec2b859 Merge pull request #713 from scchan/master-hcc_early_finalize_support 
Support more than one device code blobs in a single .kernel section

[ROCm/clr commit: 8f667984af]
 2018-10-31 14:04:06 +05:30
Alex Voicu 58c5dedc16 If we've already seen a __global__ function we do not need to re-parse 
[ROCm/clr commit: 1621deee71]
 2018-10-30 23:34:27 +00:00
Rahul Garg 8856b89b21 Fixes global symbols tracking in hip_module 
[ROCm/clr commit: 592efa017f]
 2018-10-31 03:22:38 +05:30
Alex Voicu 8599d053b4 Handle the very confusing dual encoding of the symbol name. 
[ROCm/clr commit: 1fe33b4738]
 2018-10-30 01:55:09 +00:00
Anton Gorenko 600e5cdb58 Fix allocation size of arrays with multiple and/or non-32-bit channels 
hipMallocArray and hipMalloc3DArray must use sum of bits
of all components.


[ROCm/clr commit: 6e6297f3cd]
 2018-10-29 18:12:00 +06:00
Alex Voicu 44f0ee9cff Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_get_alignment_and_size_from_metadata 
[ROCm/clr commit: b68caf7653]
 2018-10-28 17:02:10 +00:00
Alex Voicu 1ec78f6d67 Rely on code object metadat for kernarg arguments alignof and sizeof. 
[ROCm/clr commit: 299b9d1382]
 2018-10-28 17:01:00 +00:00
Yaxun Sam Liu cc101b22ed Add HIP_DUMP_CODE_OBJECT 
[ROCm/clr commit: 0027a54d8b]
 2018-10-26 14:14:00 -04:00
Yaxun Sam Liu 4012f63d87 Adding checks and debug output for fat binary for hip-clang 
[ROCm/clr commit: da6c5c152e]
 2018-10-26 14:13:12 -04:00
Aaron Enye Shi 1c05b06bcd Remove libhip_device.a static library 
Move remaining function definitions from device_util.cpp to hip_runtime.h header. We can now remove the static library completely as device_util.cpp was the last file part of target hip::hip_device .


[ROCm/clr commit: ba43d1ce1b]
 2018-10-25 19:52:35 +00:00
Maneesh Gupta 127f7d0320 Make HIP functional again with HCC from ROCm 1.9.x 
Change-Id: I214acdfd0b79dcf783993e44fe31baee64fd4dc3


[ROCm/clr commit: b01ac26948]
 2018-10-24 10:41:56 +05:30
Qianfeng Zhang b12c16f72b Make correct checking of the returned hipDeviceptr_t from read_global_description() 
[ROCm/clr commit: 443698ce00]
 2018-10-23 21:13:11 +08:00
Siu Chi Chan 5d1a265efc Support more than one bundles in a single .kernel section 
When compiling with Early Finalization enabled in HCC,
the resulting .kernel section of the host object now may
contain more than one device code bundles.  This is to
teach the HIP runtime to correctly extract all the
bundles from the .kernel section.


[ROCm/clr commit: f39873c584]
 2018-10-18 17:00:27 -04:00
Alex Voicu c505c2be38 Re-sync with upstream. 
[ROCm/clr commit: 602280eb8b]
 2018-10-18 12:27:03 +01:00
Maneesh Gupta cf95358003 Merge pull request #692 from whchung/hip-reinit-take2 
HIP program state re-initialization logic (take 2)

[ROCm/clr commit: 9143ae6bdb]
 2018-10-18 12:06:41 +05:30
Maneesh Gupta bb3706b6f6 Merge pull request #703 from mangupta/stream_create_with_priority 
Implementation for stream priority

[ROCm/clr commit: 95be669f4a]
 2018-10-17 10:53:43 +05:30
Alex Voicu 50388a3bd1 Re-sync with upstream. 
[ROCm/clr commit: 9ec697c620]
 2018-10-10 11:43:49 +01:00
Maneesh Gupta c4ed6b18fb Implementation for stream priority 
- Requires ROCm 1.9.x or higher
- Requires HCC with PR#886 merged

Change-Id: Id7c95ea091ee610e80c9ad815f1cb989cba570ca


[ROCm/clr commit: 07ee1f07d8]
 2018-10-05 16:27:46 +05:30
Rahul Garg ed5c4bca1d Corrected the width calculation logic to accomodate multi channels 
[ROCm/clr commit: a78d68d639]
 2018-10-03 12:07:38 +05:30
Rahul Garg 8981438983 Fixed image width for linear resource type texture 
[ROCm/clr commit: 167265c279]
 2018-10-01 15:28:34 +05:30
Wen-Heng (Jack) Chung 73fc21f717 Keep the map which tracks GPU kernel symbols to grow monotonically 
[ROCm/clr commit: e257de95f3]
 2018-09-26 19:49:02 +00:00
Wen-Heng (Jack) Chung 1e956fcd70 Improve performance of re-initialization logic 
Keep track of shared libaries already discovered. Do not build HSA executables
for them.


[ROCm/clr commit: 060b3c0bf8]
 2018-09-26 19:48:56 +00:00
Wen-Heng (Jack) Chung 4e0f0b3e89 HIP program state re-initialization logic 
This commit is to support kernels dynamically loaded thru means such as
dlopen() after HIP runtime initializes.


[ROCm/clr commit: 319f007bf1]
 2018-09-26 19:48:47 +00:00
Rahul Garg e0efc8a09e Return hipSuccess when sizeBytes=0 in hipMemset 
[ROCm/clr commit: bd27310127]
 2018-09-26 12:47:36 +05:30
Alex Voicu f5d204450c Align with HC Next. 
[ROCm/clr commit: 26c26f63cb]
 2018-09-17 11:50:29 +03:00
Maneesh Gupta 20abaa15c4 Merge branch 'master' into support-malloc 
[ROCm/clr commit: b23609ad42]
 2018-09-17 10:17:25 +05:30
Yaxun Sam Liu cbb7e4af8b Add HIP_DB=fatbin for debugging fat binary issues 
[ROCm/clr commit: ce13c2061a]
 2018-08-17 11:53:45 -04:00
Maneesh Gupta bb9d95188a Merge pull request #621 from ROCm-Developer-Tools/disable_startup_loader 
Disable startup loader by default and guard with env var

[ROCm/clr commit: e10fa70173]
 2018-08-10 10:18:25 +05:30
sunway513 3621ec1971 resolve a segfault bug when env var not set; remove startup_kernel_loader class 
[ROCm/clr commit: 7c467220a8]
 2018-08-09 16:40:26 +00:00
sunway513 e1d72394da Add more check to ensure the startup loader only be enabled with the env var set to 1 
[ROCm/clr commit: 7d776ae665]
 2018-08-04 01:52:27 +00:00
sunway513 ef368c0f43 Add startup loader under HIP_STARTUP_LOADER env var, disable by default 
[ROCm/clr commit: 9af862d364]
 2018-08-04 01:48:06 +00:00
Wen-Heng (Jack) Chung d7b2bab183 Revert "HIP program state re-initialization logic" 
This reverts commit 865fe01b37.


[ROCm/clr commit: 1fe7f1be91]
 2018-08-03 17:03:04 -05:00
Wen-Heng (Jack) Chung fbf97dddb5 Revert "Improve performance of re-initialization logic" 
This reverts commit 1d1b402e5b.


[ROCm/clr commit: f6178c2159]
 2018-08-03 17:02:58 -05:00
Wen-Heng (Jack) Chung 1494fc442f Revert "Keep the map which tracks GPU kernel symbols to grow monotonically" 
This reverts commit ca62af218c.


[ROCm/clr commit: 329f3b7a5e]
 2018-08-03 17:02:50 -05:00
Sarunya Pumma 57c15b815d Remove device mapping from shareWithAll memory 
When shareWithAll memory (e.g., host memory) is allocated, set appId
in hc::AmPointerInfo to -1 to indicate that this memory is not mapped
to any device.  Peer checking in ihipStream_t::canSeeMemory is not
necessary if memory is shared with all devices.  Thus, it is skipped.

Note that earlier host memory is always mapped to device 0 and HIP
always performs peer checking for all kinds of hipMemcpy.  Since the
peer checking process requires context locking, hipMemcpy from/to host
memory always grabs device 0's context lock.  Therefore, if there is
another thread holding the context lock of device 0 (e.g.,
hipDeviceSynchronize on device 0), hipMemcpy will have to wait for the
lock until it can actually perform memcpy.  This can significantly
deteriorate execution performance.

Signed-off-by: Sarunya Pumma <sarunya.pumma@amd.com>


[ROCm/clr commit: 84aadb9274]
 2018-07-28 23:15:16 -07:00
Yaxun Sam Liu 9a4744566f Support malloc/free for hip-clang 
[ROCm/clr commit: d17dc21c9c]
 2018-07-27 16:24:51 -04:00
Rahul Garg 792b9ebe7c Null check before setting offset 
[ROCm/clr commit: f17c897751]
 2018-07-24 12:25:40 +05:30
Rahul Garg 99758b3064 Set offset in hipGetTextureAlignmentOffset 
[ROCm/clr commit: 0ec64e8b24]
 2018-07-24 10:11:26 +05:30
Alex Voicu 653c7c24f4 It is unclear what I was thinking when authoring the original code... 
[ROCm/clr commit: 633ce36ea7]
 2018-07-17 14:04:57 +01:00