- Complete translation tables for cudaError <-> hipError_t.
- Remove some odd errors that were not correctly translated or not used.
- Add HIPCHECK_API to test infrastructure. Used for negative testing
an API ; if a mismatch occurs it shows the expected return error
code. Can also print a warning rather than error.
- Enable hipMemoryAllocate on NV system, and review error coded.
- Add hipErrorName to nvcc.
Change-Id: I680427dcf32a5796d5913cf9e7f3b4c6f6b91599
Conflicts:
tests/src/CMakeLists.txt
Bug fixes and improved docs for hipFree and hipHostFree.
- Passing NULL pointer initialized runtime and return hipSuccess
(not an error like before).
- add negative test for this. (hipMemoryAllocate, improved)
- Match NVCC errors for invalid pointers, add to test.
- Update hipFree and hipHostFree docs.
- hipGetDevicePointer always set *devicePointer=NULL, even for
invalid flags.
- Gate shared memory usage on specific HCC work-week.
Change-Id: I533b4fd3280a3d6cdbf05eb768976f0c7506c012
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
Note hipHostMalloc (not hipHostAlloc or hipMallocHost).
- the hipHost* is used for all HIP APIs dealing with Host memory.
(including hipHostMalloc, hipHostFree, hipHostRegister,
hipHostUnregister, hipHostGetFlags, hipHostGetDevicePointer).
- hipMallocHost is consistent with "hipMalloc" for allocating device
memory. Enumerations hipHostMalloc* also used as optional
flags parm to hipHostMalloc.
These will print compiler warnings if used, so we can weed them out
before removing.
Also add a default flags args for hipHostAlloc, in the C++ functioin
headers. So you can replace hipMallocHost(&ptr, size( with hipHostAlloc(&ptr, size)
On HIP path property obtaining done through hsa_iterate_agents and counting the devices of HSA_DEVICE_TYPE_GPU type.
P.S.
On multi-boards systems it might be problems with detection what board a GPU plugged into (not tested).
Device property MaxSharedMemoryPerMultiprocessor set equal to totalGlobalMem (HIP path).
Reason: MaxSharedMemoryPerMultiprocessor should be as the same as group memory size. Group memory will not be paged out, so, the physical memory size = total shared memory size = group region size. NVCC path remains untouched: CUDA's device property MaxSharedMemoryPerMultiprocessor is reported.
hipify is updated as well.
For HCC path concurrentKernels is set to true since all ROCR hardware supports this feature.
For NVCC path concurrentKernels is obtained from CUDA's device property cudaDeviceProp::concurrentKernels.
Aditya Atluri
Maneesh Gupta
Ben Sander
Jack Chung
pensun
Aditya Atluri
Aditya Atluri
Aditya Atluri
Evgeny Mankov
Sam Kolton