Commit Graph

143 Commits

Author SHA1 Message Date
Yaxun (Sam) Liu f12b060a3b Add new kernel launching API for hip-clang 2019-09-26 20:15:24 -04:00
mhbliao 4691c1ef8c [HIP] Remove a circular including. (#1418) 2019-09-16 08:32:47 +00:00
Sarbojit2019 1ae43cbeba [HIP] Reclaiming hipLaunchKernel API (#1353)
* [HIP] Reclaiming hipLaunchKernel API

* Reclaiming hipLaunchKernel : Incorporated review comments

* Incorporated review comments

* Removed hipLaunchKernel Macro from nvcc path
2019-08-29 01:02:41 +00:00
Rahul Garg fbc9f7e20a Add hipMemcpy3DAsync (#1320)
* Add hipMemcpy3DAsync

* Fix CI build error

* Move back stream resolution to internal function

* Remove stream redefinition and check
2019-08-16 02:13:16 +00:00
Rahul Garg 569f35a258 Add hipMemcpyParam2DAsync (#1296)
* Add hipMemcpyParam2DAsync

* Add NVCC path changes

* Clean up

* Fix build issue

* Fix else use in both sync and async apis
2019-08-09 11:50:37 +00:00
Siu Chi Chan 3cf2d7d385 Compile HIP runtime with hidden visibility by default (#1303)
* add default visibility to most APIs in program_state

* remove unwanted C++ headers

* Add symbol visibility pragmas and compiler flags

* Add visibility attribute to APIs in channel_descriptor and hip_hcc

* remove unused headers

* simplify build flags with hcc

* add pragma visibility hidden to functional_grid_launch

* [CMake] add gfx908 back
2019-08-08 08:33:04 +00:00
Rahul Garg 8b6317d041 Add support for hipFuncGetAttribute (#1279)
* Add support for hipFunGetAttribute

* Support NVCC path

* Test using sample module_api_global

* Try fixing CI build failure due to hip_prof_gen scan

* Fix for CI build issue

* Resolve conflict

* Rebase and resolve conflicts with master

* Fix build error

* Fix NVCC path build error
2019-08-08 08:27:41 +00:00
Sarbojit2019 f23c1a1499 Enabled gcc for hip host code (#1214)
* Enabled gcc for hip host code

* Adding tests for hip code + (gcc & g++), without kernels

* Excluding nvcc platforms for gcc and g++ tests + Addressing review comments

* minor code clean-up

* Add rocm include path

* Added relative path for library

* Hiding non supported functions for gcc

* Incorporating review comments
2019-08-05 09:51:36 +00:00
Jeff Daily f337ae1edb 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
2019-08-05 09:51:02 +00:00
wkwchau 7b9801fe9a Added support of hipOccupancyMaxActiveBlocksPerMultiprocessor & hipOc… (#1240)
* Added support of hipOccupancyMaxActiveBlocksPerMultiprocessor & hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags APIs

* Taking into account of SGPR usage to determine the max active blocks in hipOccupancyMaxActiveBlocksPerMultiprocessor()
2019-08-01 08:58:48 +00:00
Evgeny Mankov f0832fd968 [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
2019-07-11 14:58:16 +03:00
Jatin Chaudhary 46553d719a Adding fix to initialize a dim3 variable from blockIdx. (#1212)
Adding gitignore for ctags
2019-07-11 09:16:03 +00:00
wkwchau 28c34ead70 Implement the hipOccupancyMaxPotentialBlockSize function (#1162)
* Implement the hipOccupancyMaxPotentialBlockSize function

* Replaced hipGetDeviceProperties() call by ihipGetDeviceProperties() in ihipOccupancyMaxPotentialBlockSize()

* Add test for hipOccupancyMaxPotentialBlockSize in Module API

* Added extern declaration for ihipGetDeviceProperties() to be accessed inside ihipOccupancyMaxPotentialBlockSize()

* fixed hipOccupancyMaxPotentialBlockSize test build issue

* Fix hipOccupancyMaxPotentialBlockSize dtest

* Add BUILD_CMD in hipOccupancyMaxPotentialBlockSize dtest

* Revert "Add BUILD_CMD in hipOccupancyMaxPotentialBlockSize dtest"

This reverts commit 0480ff56f1441fc515d2c26ce33783e303423938.

* Disable hipOccupancyMaxPotentialBlockSize dtest on NVCC

* move extern declaration of ihipGetDeviceProperties to hip_module.cpp

* Update the limiation of 32 wavefronts per CU and 800/512 SGPRs for VI/pre-VI chips to calculate the occupancy
2019-06-20 05:58:29 +05:30
Evgeny c6600ba26b prof layer includes refactoring 2019-06-11 20:13:29 -05:00
Maneesh Gupta 1d5d923d36 Merge pull request #1140 from scchan/program_state_stage_2-rebase-20190524
migrate more program_state logic from header into shared library (phase II)
2019-06-05 16:09:01 +05:30
Aryan Salmanpour d8e94fd5b5 Header change for new hip API hipExtLaunchMultiKernelMultiDevice 2019-05-30 18:04:05 -04:00
Maneesh Gupta 4af81134ba Header changes for cooperative groups
Change-Id: I5f3acca94275d74adc97adcb168aed9f74951189
2019-05-28 16:58:55 +05:30
Siu Chi Chan d40a4627c3 remove code_object_bundle.hpp, clean up the old Agent_global 2019-05-24 18:13:05 -04:00
Siu Chi Chan 6852be819f moving agent_globals_impl into hip_module 2019-05-24 16:43:38 -04:00
Evgeny Mankov 9cb3e9aa5e [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
2019-05-22 18:31:39 +03:00
Wen-Heng (Jack) Chung a4db991cbf Revert "HACK for SWDEV-173477" (#1004)
* Revert "HACK for SWDEV-173477"

This reverts commit cf7ad0f184.
2019-05-13 14:42:05 +05:30
Rahul Garg e1f3dc0c80 Add fine grained host memory lock support (#1095)
* Add fine grained host memory lock support

* Fix default flag check
2019-05-13 11:48:26 +05:30
Siu Chi Chan 05a1b696da migrate program_state logic from header into shared library (phase I) (#1077)
* Revert "Revert "Use COMgr to read Kernel Args Metadata (#1006)""

This reverts commit 235c6877c8.

* Revert "Use COMgr to read Kernel Args Metadata (#1006)"

This reverts commit 2c80975e9c.

* Revert "improve program state commentary"

This reverts commit f5e4fff6cc.

* Revert "load program state once per agent"

This reverts commit 2845b4c4b8.

* start moving function_names() into the hip shared lib

* start moving code_object_blobs to a new "state" object

* Consolidate various program state related static objects into a
single program_state object

* minor clean up

* move more stuffs from functional_grid_launch into program_state

* debug make_kernarg

* moving lookup for kernargs size_align into program_state

* clean up old code for kernarg size and alignment

* update hip_module to use newer api in program_state

* Create public member functions for program_state

* move most program state functions into shared library

* Pass the data buffer size to load_executable
Otherwise, it can't figure what the data size is
just from the char* (since the data is not really a string)

* turning free functions in program state into members of program_state_impl

* change the free function globals() into a member of program_state_impl

* replace the static mutex used for populating globals

* moving associate_code_object_symbols_with_host_allocation into
program_state_impl

* move load_code_object_and_freeze_executable into program_state_impl

* moving executables and functions_names into program_state_impl

* moving kernels() into program_state_impl

* moving functions() into program_state_impl

* move get_kernargs into program_state_impl

* moving kernel_descriptor into program_state_impl

* moving kernargs_size_align calculation into program_state_impl

* Changing the handle to program_state_impl to a pointer

* moving program_state_impl into a separate inline source file

* fixing/cleaning up some header file includes

* moving member function for kernargs_size_align into program_state.cpp

* moving Kernel_descriptor into program_state.inl

* add a new class to manage agent globals

* moving all agent globals processing functions into agent_globals_impl

* load program state once per agent

re-merging PR991 against other program state changes

* fix per-agent program state member initialization

* cache executables based on elf name, isa, and agent.

This avoids program state reloading executables after a shared library is dlopened.

re-merging PR1057 against other program state changes

* protect executables cache by a global mutex

* return ref to executables cache

* adapt PR#981 Make hipModuleGetGlobal be in HIP runtime
2019-05-12 19:24:03 +05:30
Maneesh Gupta 2f43f110d9 Implement hipExtGetLinkTypeAndHopCount for ROCm devices
Change-Id: Ie5bb4f640ac6d189c7fceeab22627a7494fd10bd
2019-05-06 15:54:31 +05:30
Maneesh Gupta 54cdeabe6e Merge pull request #1040 from eshcherb/roctracer-hip-frontend-190422
hip_prof_api.h include under __cplusplus
2019-04-24 16:50:27 +05:30
Rahul Garg 94769fc8dd Add hipMallocManaged default functional support (#1036)
* Add hipMallocManaged default functional support

* Fix build error

* Add dtest
2019-04-24 16:50:03 +05:30
Evgeny 165c42483b hip_prof_api.h include under __cplusplus 2019-04-22 21:14:18 -05:00
Mr-LiuSw 64bdf82265 add little changes in hip_runtime_api.h to work with c language (#1017)
* Update hip_runtime_api.h

when i try to use mpicc or gcc to compile a c language code which call some hip runtime api , error occured as
> /path/to/hcc_detail/hip_runtime_api.h:2268:33: error: unknown type name ‘hipFuncAttributes’; 
> hipFuncGetAttributes(hipFuncAttributes* attr, const void* func);
 
add ' struct ' for the first parameter of hipFuncGetAttributes will get ride of this problem.
2019-04-16 08:35:36 +05:30
Wen-Heng (Jack) Chung 04915cea2f Make hipModuleGetGlobal be in HIP runtime so it can be discovered at runtime (#981)
* Make hipModuleGetGlobal be in HIP runtime so it can be discovered at runtime

In HIP PR #929, quite a few HIP public APIs were made as inline functions with
hidden visibility. It was necessary to support applications with shared
libraries with GPU kernels launched via hipLaunchKernelGGL(), after HIP runtime
is initialized.

In empirical tests, the implementation has been proved to be a bit too
excessive, especially for hipModuleGetGlobal(). The function is used by another
type of client applications which relies on the existence of this function
within HIP runtime so global symbols from HSA code objects loaded dynamically
at runtime can be retrieved programmtically.

This commit moves hipModuleGetGlobal() back to src/hip_module.cpp, and makes it
visible and not inline, to fulfill requirements for applications
aforementioned. It does not change the behavior of applications depending on
hipLaunchKernelGGL().

* Add HIP_INIT_API into the implementation of hipModuleGetGlobal

Address review comments.

* Fix failing HIP unit tests
2019-03-29 03:45:04 +00:00
Jeff Daily f5e4fff6cc improve program state commentary
Disambiguate calling many varibles "agent".
More detail in exception message.
Create and discard map placeholders; no need to call std::vector::clear() on map value.
2019-03-27 21:40:27 +00:00
Jeff Daily 2845b4c4b8 load program state once per agent 2019-03-27 18:19:10 +00:00
Maneesh Gupta c20d233585 Merge pull request #970 from mangupta/swdev-172995
hipExtMallocWithFlags implementation
2019-03-25 07:46:53 +00:00
Wen-Heng (Jack) Chung cf7ad0f184 HACK for SWDEV-173477
For code objects with global symbols of length 0, ROCR runtime would
ignore them even though they exist in the symbol table. Therefore the
result from read_agent_globals() can't be trusted entirely.

As a workaround to tame applications which depend on the existence of
global symbols with length 0, always return hipSuccess here.

This behavior shall be reverted once ROCR runtime has been fixed to
address SWDEV-173477
2019-03-21 17:18:16 +00:00
Maneesh Gupta ce72890dcf Merge pull request #972 from yxsamliu/global
Add declaration of symbol related API for VDI
2019-03-20 05:12:21 +00:00
Siu Chi Chan fa9495841b reimplement HIP_INIT as hip_impl::hip_init(), add hip_init() to some of the inlined API (#966)
* reimplement HIP_INIT as a function, expose it as hip_impl::hip_init()
so that it could be called from hipLaunchKernelGGL and other inlined
HIP functions

* Don't call hip_init from ihipPreLaunchKernel
2019-03-20 05:11:15 +00:00
Yaxun Sam Liu fb3241a000 Add declaration of symbol related API for VDI 2019-03-19 11:11:49 -04:00
Maneesh Gupta e44de376f7 hipExtMallocWithFlags implementation
Change-Id: Iee9e119796472200b2933d5e23be60813f33bc75
2019-03-19 11:59:22 +05:30
Evgeny 2aa88a4505 tracing callback layer update 2019-03-14 22:43:52 -05:00
Siu Chi Chan 5044c9ba49 minor cleanup 2019-03-11 19:51:57 +00:00
Siu Chi Chan fa564a5345 Fix memory leak introduced by previous change to Agent_global.
Make Agent_global manage the lifetime of the name string
2019-03-11 19:51:32 +00:00
Aaron Enye Shi fcde9fe3df Fix Agent_global variables failing hipTestDeviceSymbol
Issue: Header uses std::vector<Agent_global> agent_globals which is created by hip_module.cpp
  - Move iterator fails to copy Agent_global from library source into header version
  - Due to different versions of std::string name in struct Agent_global
Fix: Change Agent_global to use char* name instead of std::string name
2019-03-11 19:51:25 +00:00
Aaron Enye Shi f8aca3159c Fix hash_for undefined reference in hipTestConstant test
Issue: mismatch undefined symbols in different user env
  - Binary expects modified return value std::string&
  - Fails to match libhip_hcc.so: return value is std::string& but doesn't match modified C++ env
Fix: Change return value to char*, create new key std::string in header from char*
2019-03-11 19:51:18 +00:00
Yaxun Sam Liu 2dc87b6019 Fix HIP/VDI build failure due to dlopen change 2019-03-07 14:45:45 -05:00
Alex Voicu ed48847237 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.
2019-03-06 17:31:44 +05:30
Wen-Heng (Jack) Chung 8b7baa0bd9 Address code review comments to use hipDeviceptr_t 2019-03-05 05:51:05 +00:00
Wen-Heng (Jack) Chung 392271f4db Add hipMemsetD32 and hipMemsetD32Async
Add 2 extra memset functions which fills memory with integer-typed data

Also change the parameters of ihipMemset to better explain the semantic
2019-03-04 17:00:33 +00:00
Laurent Morichetti 167bbffc4f Fix compilation issues with MSVC 2019-01-09 09:32:50 -08:00
Yaxun Sam Liu 57d1f6e6cb Use __dparm for default parameter 2018-12-31 13:07:19 -05:00
Yaxun Sam Liu 02deeadaa4 Add default arguments for hipConfigureCall 2018-12-18 15:48:21 -05:00
Maneesh Gupta 99bb89b756 Merge pull request #760 from eshcherb/roctracer-hip-frontend-181113
Roctracer hip frontend 181113
2018-11-23 11:08:25 +05:30