AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad
2.162 Commits 3 Ramas 0 Etiquetas
b3fb8ea92f08eded8bdef383b68375fed3fe74e3
Gráfico de commits

2162 Commits

This Branch
This Branch
Todas las Ramas
Autor SHA1 Mensaje Fecha
Maneesh Gupta b3fb8ea92f Simplify square sample's Makefile 
Change-Id: I44349a880a3c57ca0e833d67d9c380b706655b1e
 2017-12-05 11:54:50 +05:30
Ben Sander 657e65733f Remove coordinate conversions (threadIdx.x -> hipThreadIdx_x) 
Since these are now supported directly.
 2017-12-02 07:44:54 -06:00
Ben Sander b9fa704521 Update square sample for recent HIP ease-of-use improvements 2017-12-02 07:44:27 -06:00
Ben Sander 03d1263a51 Update RELEASE.md and bugs to reflect HIP infra improvements. 2017-12-01 16:09:24 -06:00
Ben Sander 7d2675520a Merge pull request #267 from AlexVlx/feature_sanitise_conversions 
Fix curious bitwise conversions.
 2017-12-01 15:51:08 -06:00
Ben Sander dac57ba854 Merge pull request #246 from AlexVlx/feature_use_module_based_dispatch_instead_of_pfe 
Use natural dispatch syntax
 2017-12-01 14:21:25 -06:00
Alex Voicu 7c0b9a005b 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. 2017-11-30 03:29:04 +00:00
Alex Voicu 20fc68c9a1 Add missing space (the final frontier). 2017-11-29 21:50:43 +00:00
Alex Voicu 32e11e7dc6 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
 2017-11-29 21:49:10 +00:00
Alex Voicu 7acb1e6ff6 Use a much simpler guard for version 1.6, which allows for direct CUDA indexing to be used. 2017-11-29 21:47:04 +00:00
Alex Voicu 61453e15cc Merge branch 'feature_use_module_based_dispatch_instead_of_pfe' of https://github.com/AlexVlx/HIP into feature_use_module_based_dispatch_instead_of_pfe 2017-11-29 21:45:56 +00:00
Alex Voicu fbaf729f88 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
 2017-11-29 21:36:29 +00:00
Ben Sander a0cec19d6e Merge pull request #275 from mangupta/bump_version_1.5 
Bumped HIP version to 1.5 and updated summary in HIP release notes
 2017-11-29 15:06:02 -06:00
Alex Voicu d2fd1f5544 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. 2017-11-29 21:01:28 +00:00
Maneesh Gupta f4187b0f16 Bumped HIP version to 1.5 and updated summary in HIP release notes 
Change-Id: I4b612da88236674fac2c585ab6891bdd495511b0
 2017-11-29 13:56:17 +05:30
Alex Voicu b881cf713c Fix compiler version check. 2017-11-29 03:05:53 +00:00
Alex Voicu 3ed8897a5a Add missing file. 2017-11-29 02:16:44 +00:00
Alex Voicu faa546d194 Fix oversight in selection mechanism which led to erroneous code to be compiled for the grid_launch_GGL component. 2017-11-29 01:37:52 +00:00
Alex Voicu 89e9399427 Choose whether or not to use functional grid_launch based on the version of HCC used to compile. 2017-11-29 00:17:44 +00:00
Alex Voicu 5aeb5dcd6f Remove leftover agent allocated globals. 2017-11-28 19:56:04 +00:00
Alex Voicu 6e4ca3fbb4 Change memset kernel to use memcpy instead of placement new. Simplify indexers. 2017-11-28 19:45:47 +00:00
Alex Voicu 02c2bfc7ef Re-sync with upstream and re-factor platform global management for texture references. 2017-11-28 19:15:29 +00:00
Alex Voicu dc67ca3feb Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe 
# Conflicts:
#	src/hip_module.cpp
 2017-11-28 17:29:11 +00:00
Maneesh Gupta 265c3b224e Fix float2int rounding functions 
Change-Id: I67943859a6344c5eec0eaa23418c9b802ef72468
 2017-11-28 17:23:43 +00:00
Alex Voicu a6ccaf3d57 This actually (tries) to do the right thing all the way, by using memcpy for bitcasting, and not rely on undefined behaviour of a different flavour as a substitute for the original undefined behaviour. Note that the compiler will (should) optimise down to the same emitted code, since this is a pattern it understands. 2017-11-28 17:23:06 +00:00
Alex Voicu a401ce6e5d This fixes some outright quaint choices made when implementing HIP's bitwise conversion functions, by using simple reinterpret_casts, as is idiomatic. These functions are supposed to be re-entrant, correct and efficient. Sadly, they were neither: they hid a massive race condition against a value stored in global memory, which means that they were also unreasonably slow if they ever managed to be correct, and relied on union based type punning which is in a grey area of the standard. It is difficult to ascertain what may have been the reason for coming up with this quirky solution. 2017-11-28 17:23:06 +00:00
Chris Kitching 9d47a4d980 Add hipify mappings for all CUDA headers that have HIP equivalents 
I'm particularly running into issues with `device_types.h` in real
CUDA code...
 2017-11-28 17:23:06 +00:00
Ben Sander 0da0426f94 Merge pull request #256 from gargrahul/texture_driver_api_support 
Texture driver APIs support
 2017-11-27 13:52:39 -06:00
Rahul Garg a139b533bb Added note in Release notes for texture driver APIs 2017-11-27 22:06:50 +05:30
Evgeny Mankov aa05b3d84e Merge pull request #262 from ChrisKitching/frontendaction 
[HIPIFY] Mostly fix preprocessor-or-template induced issues
 2017-11-27 17:30:11 +03:00
Alex Voicu 907d41df77 Re-sync with upstream. 2017-11-24 13:04:12 +00:00
Rahul Garg 04bc5a1d1f Porting guides update for texture APIs usage 2017-11-24 12:00:55 +05:30
Ben Sander 814548407d Merge pull request #273 from mangupta/swdev-129574 
Fix float2int rounding functions
 2017-11-23 12:04:36 -06:00
Maneesh Gupta 4c96882366 Fix float2int rounding functions 
Change-Id: I67943859a6344c5eec0eaa23418c9b802ef72468
 2017-11-23 09:57:24 +05:30
Alex Voicu 08f252e4bf Remove leftover comment. 2017-11-22 19:37:03 +00:00
Evgeny Mankov 4e92a034d0 Merge pull request #263 from ChrisKitching/headers 
[HIPIFY] Add hipify mappings for all CUDA headers that have HIP equivalents
 2017-11-22 21:24:21 +03:00
Alex Voicu 4131b47134 Modify the set component of the memcpy test (unclear why there is a memset component to begin with). 2017-11-21 17:52:01 +00:00
Rahul Garg 56862b1c35 Fixed review comments 2017-11-21 21:19:06 +05:30
Alex Voicu 5e16ee0d1f This corrects how addresses are formed for symbols which reside in shared objects. For this case, the .value component of an ELF symbol holds the offset from the base VA where the shared object was loaded. Thus, to correctly obtain the VA of the object refered by the symbol, we must add the offset to the VA where the shared object is loaded. We were already doing this correctly for symbols denoting functions, but we were incorrect for those denoting objects. 2017-11-21 13:15:13 +00:00
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
Maneesh Gupta db378fbc9e Merge pull request #266 from gargrahul/fix_half2_gfx900 
Fixed half2 issue on gfx900
 2017-11-20 07:28:41 +05:30
Maneesh Gupta 1174534e85 Merge pull request #265 from phani544/nvccTests 
[nvccTests]Enabled inline_asm_vadd on nvcc
 2017-11-20 07:28:29 +05:30
Ben Sander 59956a57ca Fix test on cuda 2017-11-19 15:31:02 -06:00
Ben Sander 5a7a28ad29 Merge branch 'feature_natural_indexing' of https://github.com/AlexVlx/HIP 2017-11-19 15:25:17 -06:00