- 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
[ROCm/hip commit: 21e5c25225]
- Added memcpy and memset device functions
- Added test for memcpy and memset
Change-Id: Icd21a8dd964953b86d5e92889bf1664bee647219
[ROCm/hip commit: 8155f0e036]
- 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
[ROCm/hip commit: e51ce8fc09]
- 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
[ROCm/hip commit: f31668fee4]
add support for both cuda compatible implementation and hcc(faster)
implementation with test
Change-Id: I79a22344f458391d7dffac5f147619a542e97e4e
[ROCm/hip commit: 8264d5d6bd]
_ 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
[ROCm/hip commit: 9f29cc6989]
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
[ROCm/hip commit: bf4ba9e09e]
Add __host__ for some math functions to fix the following tests:
- hipSinglePrecisionMathHost
- hipDobulePrecisionMathHost
[ROCm/hip commit: 453d07daf4]
- Validate compile-time disables.
- Add README.md section explain how to install/use CodeXL tracing
- Add code docs on trace_helper.h
- fix color on hipLaunchKernel to green.
[ROCm/hip commit: 2569e15375]
- Control with HIP_DB=mask (env var). See src/hip_hcc.cpp for mask
values:
#define DB_API 0 /* 0x01 - shortcut to enable HIP_TRACE_API on single switch */
#define DB_SYNC 1 /* 0x02 - trace synchronization pieces */
#define DB_MEM 2 /* 0x04 - trace memory allocation / deallocation */
#define DB_COPY1 3 /* 0x08 - trace memory copy commands. . */
#define DB_SIGNAL 4 /* 0x10 - trace signal pool commands */
- Combine with HIP_TRACE to see debug with API trace.
- Use colors to distinguish different flows of debug.
- Add define COMPILE_DB_TRACE to allow removing all debug at compile-time
[ROCm/hip commit: aa03e1264c]
Ben Sander
Aditya Atluri
Evgeny Mankov
pensun
7SK
Maneesh Gupta
Jack Chung
Rahul Garg
Aditya Atluri
scchan