AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
2,709 Commits 3 Branches 0 Tags
a0bb8820c543de4a4c2c50cb5bba0083a009d00d
Commit Graph

165 Commits

Author SHA1 Message Date
Rahul Garg a796eedde9 Revert "Use memcpy kernel for all pinned memory cases in hipMemcpy2DAsync" 
[ROCm/clr commit: f554e48db3]
 2018-07-02 14:32:11 +05:30
Rahul Garg 1c2acf1467 TEMP- fix memcpy2dAsync for trsm issue 
[ROCm/clr commit: 007e2a4b5f]
 2018-06-15 16:08:29 +05:30
Rahul Garg a41cb78337 Fix stream resolution in memcpy2dasync 
[ROCm/clr commit: 2ae3be9773]
 2018-06-14 11:58:56 +05:30
Rahul Garg dfe8629088 Fix retrieved locked ptr offset 
[ROCm/clr commit: 68554e155b]
 2018-06-13 23:10:05 +05:30
Maneesh Gupta e34247ca19 Merge pull request #497 from gargrahul/fix_memcpy3d_fastpath 
Fix hipMemcpy3D for fast path

[ROCm/clr commit: 53037472ff]
 2018-06-06 14:44:02 +05:30
Rahul Garg 914a0aaf9c Fix hipMemcpy3D for fast path 
[ROCm/clr commit: 163d4a5b03]
 2018-06-05 18:54:33 +05:30
Rahul Garg 3a6e8aadbd Fix depth value for 3D allocations 
[ROCm/clr commit: fa6ce7a724]
 2018-06-04 18:00:22 +05:30
Rahul Garg 84cb355df6 Fix memcpy2D for malloc+ hostRegister 
[ROCm/clr commit: a3609eaf61]
 2018-05-31 13:14:27 +05:30
Rahul Garg 0ab5c9d6f1 Use 64x4 grid dims 
[ROCm/clr commit: 27e0566c3a]
 2018-05-24 23:51:52 +05:30
Rahul Garg 439655d213 Clean up and fix remaining bytes copy 
[ROCm/clr commit: aed2653857]
 2018-05-24 23:30:27 +05:30
Rahul Garg 4fc64ff4c5 Fix memcpy2d kernel dims 
[ROCm/clr commit: 53a7c61e9b]
 2018-05-24 17:00:12 +05:30
Rahul Garg 4e52a634a1 Correct remaining bytes in copy 2d kernel 
[ROCm/clr commit: 96b4618d26]
 2018-05-24 08:27:24 +05:30
Rahul Garg 432467d88b Optimize memcpy2D kernel use 
[ROCm/clr commit: 2c3d1498d4]
 2018-05-23 14:43:47 +05:30
Maneesh Gupta 335dad2feb Merge pull request #464 from gargrahul/fix_memcpy2d_pinned_mem_case 
Fixed memcpy2D for pinned memory case using 2D kernel

[ROCm/clr commit: 7042fe6067]
 2018-05-22 10:42:28 +05:30
Rahul Garg 54497d1e2b Fixed memcpy2D for pinned memory case using 2D kernel 
[ROCm/clr commit: 40fb44dbe6]
 2018-05-21 22:14:45 +05:30
Maneesh Gupta efe6b5e4a0 hipMemcpy returns success if sizeBytes is 0. 
Fixes SWDEV-153754 & SWDEV-154178.

[ROCm/clr commit: 66d05e6fc3]
 2018-05-21 15:38:44 +05:30
Rahul Garg 8575112052 Fix for memcpy2DAsync for pinned host memory case 
[ROCm/clr commit: 4f5bdb071c]
 2018-05-18 21:09:50 +05:30
Maneesh Gupta 106f3f591e Merge pull request #433 from gargrahul/add_hipmemset3d 
Added hipMemset3D

[ROCm/clr commit: 1c93e11cdf]
 2018-05-18 14:54:15 +05:30
Rahul Garg 84089036b0 Fixed hipMemcpy2D to handle 1D memcpy case 
[ROCm/clr commit: 8413fb51e1]
 2018-05-16 11:07:10 +05:30
Rahul Garg a7be5be448 Added hipMemset3D 
[ROCm/clr commit: 78568435da]
 2018-05-07 10:24:30 +05:30
Lakhan Singh 6278830c1a Null checks added for hipmallocpitch and hipmemcpy apis 
[ROCm/clr commit: 12d8a47c0c]
 2018-05-03 09:27:50 +05:30
Rahul Garg 6f05c37cde Fix texture 3D for HIP/NVCC 
[ROCm/clr commit: 1d76c48e3d]
 2018-05-02 11:56:37 +05:30
Lakhan Singh 06dd83f7f1 SWDEV-141024 
[ROCm/clr commit: 74faa61d52]
 2018-04-20 17:40:00 +05:30
Rahul Garg 39a19cc9e8 Added hipMemset2DAsync support 
[ROCm/clr commit: fcc0866681]
 2018-04-17 18:27:27 +05:30
Rahul Garg cf6322da63 Correct missed ihipMemsetCopyDataType change 
[ROCm/clr commit: 2ee80397f2]
 2018-04-12 10:27:19 +05:30
Rahul Garg 312b764721 Changed ihipMemsetCopyDataType to ihipMemsetDataType 
[ROCm/clr commit: 2658963f5b]
 2018-04-12 09:29:22 +05:30
Rahul Garg aff8d83cdf Fix hipMemset stream resolution 
[ROCm/clr commit: 88073a17c1]
 2018-04-11 19:01:53 +05:30
Rahul Garg a3bc29976b hipMemset refactoring 
[ROCm/clr commit: 40846f6f8e]
 2018-04-11 15:58:48 +05:30
Maneesh Gupta cdb7f8f354 hipMemcpyAsync returns success when trying to copy 0 bytes 
Change-Id: I4c0ee7ccc7563e2df657b50356cdd7fec9a1ef15


[ROCm/clr commit: 445de66560]
 2018-04-09 12:39:44 +05:30
Maneesh Gupta 46ddefedee Apply .clangformat to all repo source files 
Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344


[ROCm/clr commit: 9e47fccc89]
 2018-03-12 11:29:03 +05:30
Alex Voicu c4d704b2ad Change directory name to match HIP lowercase style. 
[ROCm/clr commit: 696f0595a3]
 2018-02-22 13:15:10 +00:00
Maneesh Gupta 2dc84aeaf9 Merge pull request #321 from gargrahul/hipMemcpyArray_Functions 
Added support for hipMemcpy Array functions-

[ROCm/clr commit: 647d1ba310]
 2018-02-12 10:36:38 +05:30
Rahul Garg f15543d5cf Fixed host allocated globals address lookup for host usage 
Fixed texture driver APIs failure


[ROCm/clr commit: b8c23f979b]
 2018-01-30 18:06:31 +05:30
Rahul Garg ec4cc27986 Added support for - 
- hipMemcpyFromArray
- hipMemcpyAtoH
- hipMemcpyHtoA


[ROCm/clr commit: ca5bcb5af4]
 2018-01-16 11:44:19 +05:30
Rahul Garg 467b908368 Added support for 
- 3D texture driver APIs
- hipMalloc3D
- hipMemcpy3D for destination other than array


[ROCm/clr commit: 105df94cd0]
 2017-12-05 14:11:13 +05:30
Ben Sander cabf4d817c Fix some cppcheck style issues. 
[ROCm/clr commit: 5933ca7300]
 2017-12-01 20:45:34 +00:00
Ben Sander b1aee08dc1 Fix warning from default cppchek. 
[ROCm/clr commit: 743c54a012]
 2017-12-01 20:45:33 +00:00
Alex Voicu 76336d46b4 Fix legacy mode detection of the address of an agent allocated variable. In this mode, there exist two executables per each code object, one created by HCC and one created by HIP. Since we dispatch through HCC in legacy mode, we should obtain the address for an agent allocated variable from the latter's executable. Also add two omitted validity checks, whose absence could lead to segfaults when the current process had no .kernel section and / or when an invalid or empty blob was extracted from the latter. 
[ROCm/clr commit: 33bb425013]
 2017-11-30 03:29:04 +00:00
Alex Voicu a27fa25d76 Revert "Revert adoption of CUDA indexing in general - this can only work with later versions of the compiler, just like module based dispatch, and thus must be guarded against usage in earlier (e.g. 1.6) versions." 
This reverts commit 1cd6416


[ROCm/clr commit: 4966518846]
 2017-11-29 21:49:10 +00:00
Alex Voicu a1d5f0ab0c Revert "Revert adoption of CUDA indexing in general - this can only work with later versions of the compiler, just like module based dispatch, and thus must be guarded against usage in earlier (e.g. 1.6) versions." 
This reverts commit d2fd1f5


[ROCm/clr commit: 2557000b56]
 2017-11-29 21:36:29 +00:00
Alex Voicu abf92dd6e7 Change memset kernel to use memcpy instead of placement new. Simplify indexers. 
[ROCm/clr commit: d5c8de3f41]
 2017-11-28 19:45:47 +00:00
Alex Voicu 579a3187da Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe 
# Conflicts:
#	src/hip_module.cpp


[ROCm/clr commit: d37a5a6008]
 2017-11-28 17:29:11 +00:00
Rahul Garg 212afdf56b Fixed review comments 
[ROCm/clr commit: 38029f2849]
 2017-11-21 21:19:06 +05:30
Rahul Garg dc2c56c43d Changed function hipMemcpy_2D to hipMemcpyParam2D 
[ROCm/clr commit: 24307fe5c4]
 2017-11-21 12:36:24 +05:30
Alex Voicu 1366b2d5ae Refactor the __device__ versions of memset and memcpy to be less awkward i.e. not return nullptr as opposed to the destination pointer (it can only be assumed it was done for maximum confusion) and actually unroll as they claim to. Change all of the {to, from}Symbol functions to use hipModuleGetGlobal, as opposed to hc::accelerator::get_symbol_address which is no longer valid with module based dispatch. 
[ROCm/clr commit: f8c1c1b38e]
 2017-11-21 02:40:34 +00:00
Rahul Garg bb8c31c198 Texture driver APIs support 
[ROCm/clr commit: 1851c153f6]
 2017-11-09 22:10:55 +05:30
Alex Voicu 142213a52b Clean up trailing whitespace so as to reduce noise in #246. 
[ROCm/clr commit: 0ad1308efd]
 2017-11-08 00:08:55 +00:00
Alex Voicu 21d1d3f1ed Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe 
# Conflicts:
#	src/hip_module.cpp


[ROCm/clr commit: 17753cbd92]
 2017-11-03 10:53:39 +00:00
Alex Voicu 840895525a Correctly deal with functions from shared objects, wherein the program visible VA == so_base_va + st_value(function_symbol). Remove quaint usage of pfe for hipMemset (which is actually fill_n). 
[ROCm/clr commit: 4693c5e56c]
 2017-11-01 22:33:13 +00:00
Alex Voicu b738633bcb This switches HIP from its currently convoluted macro + pfe based dispatch mechanism to a more natural one partially based on the existing module API. The basic idea is that HCC will always correctly emit __global__ functions: as empty-bodied stubs, on host, and as kernels, on device. It then becomes trivial to obtain the mangled name on host, at dispatch, from the function's address, and then to use the mangled name to retrieve the kernel. This should address all problems stemming from serialisation, dubious mismatches due to the manufactured functor, macro-isms et al. It also immediately enables support for generalised globals as a consequence of that being available in the module API. Finally, it will make debug much easier, since the actual names of the __global__ functions will automatically be used in traces etc. One detail is that due to how dispatch works now (hipLaunchKernel and hipLaunchKernelGGL are themselves variadic function templates which deduce the function type of the callee), in certain cases it may be necesssary to insert explicit casts to ensure that the variadic argument list selects a viable overload - this can be observed in some unit tests. Eventually we may be able to remove this limitation, but for now it does not appear terribly onerous. The code is not extremely HIPpie, nor is it fully optimised, but rather is intended as a starting point for the HIP team to make its own. 
[ROCm/clr commit: 28f87f7d2e]
 2017-11-01 15:09:59 +00:00