AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność
2 375 Commity 3 Gałęzie 0 Tagi
d98eaf592309300a710cbf00f4e179f0f0d59c7d
Wykres commitów

706 Commity

Autor SHA1 Wiadomość Data
Rahul Garg 9866fa250d Changed function hipMemcpy_2D to hipMemcpyParam2D 2017-11-21 12:36:24 +05:30
Alex Voicu 9d088d2283 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. 2017-11-21 02:40:34 +00:00
Alex Voicu 1824fb7698 Clean-up some remaining noise in program_state.cpp. 2017-11-20 22:41:46 +00:00
Alex Voicu 7d5a45ac1a Correct ill-formed merge in earlier commit and adjust for differences with the new CUDA natural indexing mechanism. 2017-11-20 16:33:52 +00:00
Alex Voicu c5f2b22d0d Re-sync with upstream. 2017-11-20 15:34:50 +00:00
Ben Sander e8ede28ec4 Merge pull request #264 from pzins/missing_end_marker 
Fix missing MARKER_END
 2017-11-20 06:08:01 -06:00
Rahul Garg f97c5f9a64 -Moved coGlobals in hipModule class (takes care of multi module case) 
-Used mutex scope for updating coGlobals
 2017-11-20 16:23:18 +05:30
Rahul Garg c7d60a7a75 Update hipModuleGetTexRef API 2017-11-19 22:10:46 +05:30
Alex Voicu cffd0e14eb This implements the trivial change needed to move back from the hip{Something}_{x, y, z} macros to the natural CUDA syntax of Something.{x, y, z}. This is contained in lines 384-404 in hip_runtime.h. All of the other changes have to do with changing unit tests to use this syntax. The macros are retained for backwards compatibility. 2017-11-19 01:54:12 +00:00
Rahul Garg ae1eb7a03a Removed redundant desc variable 2017-11-15 18:28:27 +05:30
Rahul Garg 4b19c2aa0c -Fixed texture driver API sample 
-Added hipTexRefSetAddress and hipTexRefSetAddress2D APIs
 2017-11-15 18:23:28 +05:30
Rahul Garg 63680edd30 Texture code reorganized 2017-11-14 11:09:35 +05:30
Pierre 6baaed8e48 Fix missing MARKER_END 
Logging status of hipCtxSynchronize was missing
Test if hip profiling is active for MARKER_END in ihipPostLaunchKernel
Add MARKER_END after the completion of a kernel launched through
the "grid launch"
 2017-11-13 16:13:19 -05:00
Alex Voicu f7726cd416 Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe 2017-11-09 23:43:07 +00:00
Rahul Garg ef09c4918d Texture driver APIs support 2017-11-09 22:10:55 +05:30
Maneesh Gupta 31bcb59f62 Merge pull request #250 from AlexVlx/feature_add_agent_global_support 
Support for agent globals
 2017-11-09 07:52:09 +05:30
Alex Voicu 3d248927e4 Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe 
# Conflicts:
#	tests/src/runtimeApi/stream/hipStreamSync2.cpp
 2017-11-08 10:26:30 +00:00
Alex Voicu d8e323d4b5 Clean up trailing whitespace so as to reduce noise in #246. 2017-11-08 00:08:55 +00:00
Alex Voicu adaf6b8dff Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe 2017-11-07 00:01:22 +00:00
Ben Sander f278f67d2d Check for null event in hipEventElapsedTime 2017-11-06 23:49:31 +00:00
Ben Sander 16708dd2e0 hipStreamWaitEvent returns success if event created but not recorded 2017-11-06 23:49:31 +00:00
Ben Sander 4a2e6f8955 Make hipEvent_t thread safe. 
Support re-recording of same event by different threads.

- Add criticalData structure to hipEvent_t, similar to mechanism used
  for streams, contexts, device.  Events are always locked
  after streams to avoid deadlock.
- ihipEvent_t::locked_copyCrit can be used to copy critical state
  including marker.  The critical state in the event can then
  be re-recorded.
- refactor hipEventElapsedTime.  Remmove stale debug code, native signal
  refs.
 2017-11-06 23:49:25 +00:00
Maneesh Gupta 1131c9e41a Merge pull request #251 from ROCm-Developer-Tools/fix_event_state 
Set event state AFTER it is recorded.
 2017-11-06 07:28:11 +05:30
Maneesh Gupta a62d5aa875 Merge pull request #249 from bensander/warn_event 
Add HIP_DB=warn + message if sync on dangerous event.
 2017-11-06 07:25:40 +05:30
Ben Sander 4c3b65a5cd Set event state AFTER it is recorded. 2017-11-05 10:33:18 -06:00
Alex Voicu bb1176001f Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe 
# Conflicts:
#	src/hip_module.cpp
 2017-11-03 10:53:39 +00:00
Alex Voicu 328c18b886 This introduces correct support for agent global variables, and implements hipModuleGetGlobal as an actual equivalent for cuModuleGetGlobal. 2017-11-03 01:44:48 +00:00
Alex Voicu 2cacda91bb 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). 2017-11-01 22:33:13 +00:00
Ben Sander 09d866a639 Merge pull request #237 from bensander/use_ctxptr_for_p2p 
Use ctxptr for p2p
 2017-11-01 18:55:25 +01:00
Ben Sander 70c25bdf8e Add HIP_DB=warn + message if sync on dangerous event. 2017-11-01 10:44:34 -07:00
Ben Sander 86f62accfd Merge pull request #245 from scchan/centos_fixes 
various fixes for centos/rhel
 2017-11-01 18:10:29 +01:00
Alex Voicu c2482d1255 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. 2017-11-01 15:09:59 +00:00
Siu Chi Chan 99d32a195f Centos/RHEL - remove usage of constexpr since libc++ doesn't enable ctor for constexpr pair in C++11 2017-10-31 18:16:12 +00:00
Ben Sander 7e908bdec8 Add ns-level timer for HIP API routines 
Refactor some miuses of ihipLogStatus, these should only be in top-level
HIP APIs and should be paired with HIP_API_INIT calls.
 2017-10-30 20:20:51 +00:00
Ben Sander 2e8ec71e40 Merge pull request #222 from bensander/fix_device_prop 
Fix device prop
 2017-10-30 17:58:48 +01:00
Ben Sander d610f16c47 Check for null copyEngine before looking at peers. 2017-10-30 16:58:03 +00:00
Siu Chi Chan a9789ddcda Merge remote-tracking branch 'origin/master' into HEAD 2017-10-27 01:18:28 -04:00
Ben Sander f288f24e95 Merge pull request #198 from AlexVlx/feature_support_globals_for_module_api 
Feature support globals for module api
 2017-10-27 01:53:34 +02:00
Ben Sander 7d30f32332 Fix bug with peer-to-peer combined with context API 
- Store context inside the tracker rather than using int deviceID that
  was always mapped to primary context
- IsPeerWatcher now based on device IDs rather than specific peers.
 2017-10-26 19:44:22 +00:00
Aditya Atluri 5d646d0fe3 Enhance debug for copy pointers 
- show more pointer tracking fields
- show pointer info before and after "tailoring'
 2017-10-26 19:44:22 +00:00
Siu Chi Chan d91a4f5bd6 add HC_FEATURE_PRINTF around the printf buffer definition 2017-10-25 12:00:02 -04:00
Siu Chi Chan 1ddee10c2f printf support for module API 2017-10-24 00:55:41 -04:00
Siu Chi Chan 5b9ce032d6 replace __hcc_workweek__ with HC_FEATURE_PRINTF flag 2017-10-23 18:30:08 -04:00
Maneesh Gupta b792f9f507 Make elfio headers private 
Change-Id: I3ba174bb46e84a75380207d93a0da6fe3703689e
 2017-10-23 10:24:36 +05:30
Ben Sander dd24983571 Remove printf 2017-10-20 13:24:04 -07:00
Ben Sander acf89b43d4 Update device properties. 
- clear properties to defined initial state.
- enable some property flags which are now supported.
 2017-10-20 15:52:13 +00:00
Ben Sander c9f906c2ce Modify device properties to use pool API. 
- Also better error code checking
 2017-10-20 14:49:29 +00:00
Siu Chi Chan ccef1cbd6e hipDeviceReset(): make sure to reinitialize the printf buffer in hcc RT 2017-10-18 16:26:13 -04:00
Alex Voicu 70786a5563 This fixes incorrect usage of the reader object, which created arcane 
mismatches when one reader accessed another's section.
 2017-10-09 15:46:38 +01:00
Alex Voicu ffffe052b6 This adds cursory support for globals to the HIP module loading API. The 
style is purposefully alien so as to signal that HIP experts should turn
it into HIP worthy code as soon as possible.
 2017-10-09 13:27:11 +01:00