Commit Graph

6317 Commits

Author SHA1 Message Date
Wenkai Du e8e58e9ce5 Use NUMA distance for hop count calculation 2019-05-15 21:50:35 +00:00
Evgeny Mankov 70f01fad73 Merge pull request #1111 from emankov/master
[HIPIFY][tests] Add reverse engineered HIP sample "stream"
2019-05-15 20:18:51 +03:00
Evgeny Mankov 64eeeca6ce [HIPIFY][tests] Add reverse engineered HIP sample "stream"
+ Add additional checks for extern __shared__ due to [#1109]
2019-05-15 20:17:03 +03:00
Evgeny Mankov 84a42e1769 Merge pull request #1110 from emankov/master
[HIPIFY][fix][#1109] Do not preserve extern __shared__ for IncompleteArrayType
2019-05-15 20:09:00 +03:00
Evgeny Mankov fa3dda9107 [HIPIFY][fix][#1109] Do not preserve extern __shared__ for IncompleteArrayType
+ Update tests accordingly
2019-05-15 20:05:56 +03:00
foreman 15d1df9255 P4 to Git Change 1783301 by cpaquot@cpaquot-ocl-lc-lnx on 2019/05/15 11:57:57
SWDEV-189488 - [HIP] Caffe2 TensorTest.TensorSerializationMultiDevices fails
	1. Make sure to set attributes->device to current device for host malloc'd
	2. Return hipSuccess for hipDeviceCanAccessPeer

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#56 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_peer.cpp#4 edit
2019-05-15 12:22:50 -04:00
Konstantin Pyzhov c8f92bebf3 Upload pre-built kernel binary for hipModule test. 2019-05-15 07:19:40 -04:00
Jenkins efc3ffb9d2 Merge 'master' into 'amd-master'
Change-Id: I15e04f87004767e238af9aa1ab9567921594d07f
2019-05-15 05:45:15 -04:00
Evgeny Mankov 8def9412ff Merge pull request #1104 from emankov/master
[HIPIFY][tests] Add reverse engineered HIP sample Profiler
2019-05-14 16:59:36 +03:00
Evgeny Mankov d74d03aa74 [HIPIFY][tests] Add reverse engineered HIP sample Profiler
+ Add missing cuda_profiler_api.h to hip/hip_profile.h transformation.
NOTE: HIP Profiler API is under development. This is NOT WORKING example.
TODO: Find out a way to generate HIP_SCOPED_MARKER, HIP_BEGIN_MARKER, HIP_END_MARKER, declared in hip/hip_profile.h in particular place (signatures are to obtain).
2019-05-14 16:43:44 +03:00
Jenkins ab9e8a034e Merge 'master' into 'amd-master'
Change-Id: I476b12c4e051b47fb2e62dbea709d268d9a2f630
2019-05-14 02:12:10 -07:00
Evgeny Mankov aa32693e8f Merge pull request #1102 from emankov/master
[HIPIFY][tests] Add reverse engineered HIP sample hipEvent
2019-05-13 22:14:41 +03:00
Evgeny Mankov 3bc3b61fb4 [HIPIFY][tests] Add reverse engineered HIP sample hipEvent 2019-05-13 22:12:43 +03:00
Evgeny Mankov ee2823666b Merge pull request #1101 from emankov/master
[HIPIFY][tests] Add reverse engineered HIP sample MatrixTranspose
2019-05-13 19:39:16 +03:00
emankov 4b861bca39 [HIPIFY][tests] Add reverse engineered HIP sample MatrixTranspose 2019-05-13 19:37:18 +03:00
Jenkins 43c928aaa2 Merge 'master' into 'amd-master'
Change-Id: Ibca6232ef00edf0c11a5db5f6f10a5bcbb963625
2019-05-13 05:46:17 -04:00
Wen-Heng (Jack) Chung 9b9257f9b0 Revert "HACK for SWDEV-173477" (#1004)
* Revert "HACK for SWDEV-173477"

This reverts commit d941f19399.
2019-05-13 14:42:05 +05:30
Maneesh Gupta 693bd556d4 Merge pull request #1083 from gargrahul/fix_hip_impl_visible_agents
Maintain HIP_VISIBLE_DEVICES for kernel launch
2019-05-13 14:20:18 +05:30
Rahul Garg aeeab1b23f 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
Nick Curtis 5257b54a39 Markdown fixes & Whitespace cleanup for samples (#1096)
* Fix multiline code blocks in README's

* Whitespace cleanup
2019-05-12 19:27:44 +05:30
Maneesh Gupta 8c4b161a45 Merge pull request #1094 from mangupta/hit_improvements
[dtests] Add new tests to directed tests
2019-05-12 19:25:21 +05:30
Siu Chi Chan f5eb91d53d 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 a3d118eaa8.

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

This reverts commit 8a548bf40b.

* Revert "improve program state commentary"

This reverts commit 7aada87cbd.

* Revert "load program state once per agent"

This reverts commit c9117de8eb.

* 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
foreman 680e1a6552 P4 to Git Change 1781540 by kpyzhov@hip-vdi-pal-win on 2019/05/10 17:10:30
SWDEV-188234 - Corrected reading binary module file in hipModuleLoad().

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#26 edit
2019-05-10 19:19:11 -04:00
foreman 0949beeeb2 P4 to Git Change 1781384 by michliao@hliao-dev-00-hip.rocm-workspace on 2019/05/10 14:00:04
SWDEV-189216 - Export `hipExtModuleLaunchKernel`

	RB: http://ocltc.amd.com/reviews/r/17332/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#16 edit
2019-05-10 15:50:37 -04:00
Jenkins 08fd279dd9 Merge 'master' into 'amd-master'
Change-Id: I390e88de292d69072474b3ed1e5537d2477e9a0f
2019-05-10 11:30:25 -04:00
foreman 1d60371d01 P4 to Git Change 1780704 by yaxunl@yaxunl-lc10 on 2019/05/09 13:26:13
SWDEV-145570 - Fix device name mismatch for gfx906.

	For now hip-clang can only emits gfx906 ISA with conservative configurations, i.e. with ecc on and xnack on, therefore it is always gfx906. It is still under discussion how to encode the target id for xnack off or ecc off.

	Therefore, the reasonable solution for now is just allow code object marked as gfx906 to be loaded on any device name that starts with gfx906. We will have more detailed control once hip-clang is able to emit code object for gfx906 with ecc off or xnack off.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#30 edit
2019-05-09 14:19:54 -04:00
Maneesh Gupta 5b607e14a6 Merge pull request #1084 from mhbliao/hliao/master/api_ext
[hip] Add API `hipExtModuleLaunchKernel` in HIP runtime
2019-05-09 18:26:31 +05:30
Maneesh Gupta 88abfde2f8 [dtests] Fix hipModule test for nvcc path
Change-Id: If918b87b848a825242e06b0d552a7be188a1c4b6
2019-05-09 18:17:19 +05:30
Maneesh Gupta eb637766b9 [dtests] Add complex_loading_behavior test
Change-Id: Iadf135cb727a1a3761abef20336d652b159c7dcd
2019-05-09 18:03:42 +05:30
Maneesh Gupta 79843f3b12 [dtests] Add hipModule test to unit tests
Change-Id: I1dac38f8580265e2e9c82d88e4f070a2ff87f60b
2019-05-09 11:36:46 +05:30
Maneesh Gupta 49a2d785d0 [hit] Add support for BUILD_CMD 2019-05-09 11:36:26 +05:30
Maneesh Gupta 622ea32964 [hit] Remove CUSTOM_CMD
Change-Id: Ia156fe6aab9cfcc11284823ea5131e33eaf962bc
2019-05-09 09:59:18 +05:30
Maneesh Gupta 9f2d1453fb [hit] Rename RUN -> TEST & RUN_NAMED -> TEST_NAMED
Change-Id: I75e24f15129973cee15fc9dac65d678bd2172074
2019-05-09 09:59:18 +05:30
Evgeny Mankov bfb0524e13 Merge pull request #1090 from emankov/master
[HIPIFY][python] Initial support of hipify-python generation from hipify-clang
2019-05-08 19:12:08 +03:00
Evgeny Mankov 6b370e7743 [HIPIFY][python] Initial support of hipify-python generation from hipify-clang
+ Only a generation of transformation map of CUDA entities is implemented.
+ 2 hipify-clang options are added: -python, -o-python-map-dir.
+ Explicitly set -roc option for cuda_to_hip_mappings.py generation.
+ Generated file already might be used by pytorch team.
2019-05-08 19:08:55 +03:00
Evgeny Mankov 86fe1d3f1a Merge pull request #1089 from emankov/master
[HIPIFY][perl] Support of hipify-perl generation from hipify-clang: n…
2019-05-08 15:59:45 +03:00
Evgeny Mankov 9ddc316fa7 [HIPIFY][perl] Support of hipify-perl generation from hipify-clang: next steps
+ Generate transformation map sorted by entity type.
+ Add a generation of supported header files.
2019-05-08 15:25:06 +03:00
Maneesh Gupta 450c7ab295 Merge pull request #1088 from ROCm-Developer-Tools/mangupta-patch-1
[ci] Enable tests on ROCm 2.4
2019-05-08 12:44:02 +05:30
Maneesh Gupta 07bd3ecfad [ci] Enable tests on ROCm 2.4 2019-05-08 12:07:33 +05:30
Jenkins 2b8431f59b Merge 'master' into 'amd-master'
Change-Id: I0b8ccb710d733f41d62227750daf4adc38f401da
2019-05-07 12:15:02 -07:00
Evgeny Mankov 26d4677091 Merge pull request #1085 from emankov/master
[HIPIFY][perl] Initial support of hipify-perl generation from hipify-clang
2019-05-07 17:30:39 +03:00
Evgeny Mankov 5a3d33a338 [HIPIFY][perl] Initial support of hipify-perl generation from hipify-clang
+ Only a generation of transformation map of CUDA entities supported by HIP is implemented.
+ 3 hipify-clang options are added: -perl, -o-perl-map, -o-perl-map-dir.
+ OptionsParser mode is changed from OneOrMore to Optional to support hipify-perl generation without actual hipification.
+ Add explicit control of source files specification absence in case of no perl generation.
2019-05-07 17:27:34 +03:00
Maneesh Gupta 7264f6b64e Merge pull request #1082 from gargrahul/fix_hipmemcpy_symbol_nvcc
Fix symbol address issue on NVCC path
2019-05-07 16:17:01 +05:30
Maneesh Gupta d5abe65668 Merge pull request #1081 from mangupta/swdev-181624
Implement hipExtGetLinkTypeAndHopCount for ROCm devices
2019-05-07 16:15:41 +05:30
Maneesh Gupta f931152280 Merge pull request #1075 from mhbliao/hliao/master/test_fix2
[test] Add device variant of `std::declval`.
2019-05-07 16:15:01 +05:30
Maneesh Gupta 1d4941e487 Merge pull request #1074 from mhbliao/hliao/master/test_fix
[test] Use explicit cast for address space cast.
2019-05-07 16:09:15 +05:30
Maneesh Gupta 98ab402fcb Merge pull request #1073 from kpyzhov/multi-thread-device-test
hipMultiThreadDevice test: Reduced maximum number of created HIP stre…
2019-05-07 16:08:37 +05:30
Maneesh Gupta fea21dc6d5 Merge pull request #1072 from kpyzhov/master
Refined hipSetDevice test.
2019-05-07 16:07:36 +05:30
Maneesh Gupta 730763c817 Merge pull request #1069 from mhbliao/hliao/master/test_cleanup
[test] Remove unused common routines.
2019-05-07 16:02:57 +05:30
Maneesh Gupta d82d6b499e Merge pull request #1068 from mhbliao/hliao/master/dev_vec_func
[devfunc] Add necessary `__device__` and `__host__` attributes.
2019-05-07 16:01:48 +05:30