- Fix calls to HIP_INIT_API to pass all function arguments.
- Change ihipFunction to follow coding convention:
- leading underscore for member fields,
- camelCase for member fields.
- move kernel print function inside ihipPreLaunchKernel.
- add HIP_TRACE_API_COLOR, control color of messages.
- add ihipLogStatus wrapper to hipDeviceSynchronize()
Change-Id: I20bbb644da213f821404648945197254e3648fc9
- move _av into stream critical section. ( HCC accelerator_view is not
thread-safe but HIP steram is. )
- Refactored many places in code that need to acquire critical section.
some were previously thread races, ie enqueueing marker.
-remove support for GRID_LAUNCH_VERSION < 20
-Enable USE_AV_COPY based on HCC work-week.
- Review hipModule docs, some calrity/editing.
Change-Id: I3ce7c25ece048c3504f55ecd4683e506bb1fc8b6
- Use standard print functions for streams.
- Add HIP_INIT macro, for cases where we want to initialize HIP but not
log an API (ihipPreKernelLaunch).
Change-Id: If43cf8a363d918bcd3722a2e6a965d4cfa2e03e7
- hipLaunchModuleKernel maps to cuLaunchKernel
- Whole lot of new error codes added for the use of driver api
- KernelParams arguments is not yet supported
- hipLaunchModuleKernel is a synchronous api (will change eventually)
- All the commands in a stream will wait on host when hipLaunchModuleKernel is called on it
Change-Id: Ib4a4fae1db06fbb3a81d5a5575b026aa821264ed
- stubs and documentation in include/hcc_details/hip_runtime.h
- stubs with "no-op" in src/hip_memory.cpp
- document update in hip_kernel_language.md, add suggestions to
disable L1 and L2 caches when using the threadfence routines.
Change-Id: Ic0753170f802003055bca9d7476d7f48817b98b7
_ Use fields from GRID_LAUNCH_20 structure
(See USE_GRID_LAUNCH_20 define, currently set to 0)
"1" will require HCC support.
- Remove old DISABLE_GRID_LAUNCH support.
Change-Id: I584ce648d217251789a6283cf27feb24cb7dc8d1
1. Added copyright for device float test
2. Added device double functions support
3. Added device double functions test
4. Corrected device function signatures in headers
Change-Id: I13c8829682c925992f5cad84062bc9f702fe4048
commit 9548493fa754b3bf5c31cbdc2211db1e73e8c07c
Author: Jack Chung <whchung@gmail.com>
Date: Mon May 23 11:57:23 2016 +0800
Rename hipExternShared test to hipDynamicShared
Change-Id: I180d9d539420fb69cfc121eceaa7db9da03483b2
commit 827081f8244a38f010789d556db0c4ff7b6422d8
Author: Jack Chung <whchung@gmail.com>
Date: Mon May 23 11:56:27 2016 +0800
Rename HIP_DECLARE_EXTERN_SHARED to HIP_DYNAMIC_SHARED
Change-Id: I22362d179812ac547e0f11ba4e2bb999050e08ae
commit 4c277228ed41af187739610fa17eab1fb144c947
Author: Jack Chung <whchung@gmail.com>
Date: Thu May 19 17:49:52 2016 +0800
Adopt new interface to get dynamic LDS in hc.hpp
Change-Id: I47b433b714633a4c97df87c40a0b1d3386429a00
commit 5a36117d777064113a528dc47b42e8c8413baa97
Author: Jack Chung <whchung@gmail.com>
Date: Thu May 19 11:29:24 2016 +0800
Add test patterns for regular expression to match "extern __shared__"
These test patterns should better be saved as an individual test case, but I'm
not familiar with HIP test structures so I leave them as comments in hipify as
of now.
Change-Id: I7fee89c89b9e73de2133357a226ec0c769733531
commit 1b26284168c7f5339f63338fd0149bed5d994656
Author: Jack Chung <whchung@gmail.com>
Date: Thu May 19 11:25:23 2016 +0800
Add one HIP unit test to use HIP_DECLARE_EXTERN_SHARED
Change-Id: I4d9907815920693a74ea9d575fe26e7c67636109
commit 77b816ee5972b13d829d5bbcf06fbfd07acea2af
Author: Jack Chung <whchung@gmail.com>
Date: Wed May 18 19:18:59 2016 +0800
Adopt HIP_ prefix for DECLARE_EXTERN_SHARED macro
Change-Id: I555ded16b449b67d2e20904013d86fe1ded6a2be
commit ef0997939c3578a9ae11621bf21c0416f04d2622
Author: Jack Chung <whchung@gmail.com>
Date: Wed May 18 17:42:04 2016 +0800
Modify hipify to support converting extern __shared__ to DECLARE_EXTERN_SHARED macro
Added regular expression to search & replace extern __shared__ declarations to
DECLARE_EXTERN_SHARED macro.
Limitation:
- Won't work if "extern __shared__" is declared at global scope
Sample Usages:
extern __shared__ double foo[];
extern __shared__ unsigned int foo[];
extern volatile __shared__ double foo[];
extern volatile __shared__ unsigned int sdata[];
extern __shared__ volatile unsigned int sdata[];
extern __shared__ T s[];
extern __shared__ T::type s[];
extern __shared__ blah<T>::type s[];
extern __shared__ typename mapper<Float>::type s_data[];
extern __attribute__((used)) __shared__ typename mapper<Float>::type s_data[];
Change-Id: I2be0b7039adeddb789f5a2b067d403a43fdc3e26
commit 93ff268724493aedfacdcd5a5aa9a100f4ebaed0
Author: Jack Chung <whchung@gmail.com>
Date: Wed May 18 15:13:09 2016 +0800
Introduce DECLARE_EXTERN_SHARED macro to encapsulate "extern __shared__" decls
Change-Id: I93b2d37c763195b0ca9fd0afee78605a1e3272db
commit cff9c95412de343cc6405158b5acc4f1029267ff
Author: Jack Chung <whchung@gmail.com>
Date: Wed May 18 12:53:54 2016 +0800
Add __get_dynamic_groupbaseptr() to point to dynamic LDS
Change-Id: I97b548d8a691488057617c551a8f331cad7afc77
Change-Id: I84e7875b76fa1f59e860e19c93bd4209cdd1fd2c
Maneesh Gupta
Ben Sander
pensun
Aditya Atluri
Evgeny Mankov
7SK
Jack Chung
Rahul Garg
Aditya Atluri