Commit Graph

238 Commits

Author SHA1 Message Date
Rahul Garg 84a789c85c Added further hipCtxXXX Apis
Change-Id: I286d962a06cee656c1c652b3f6b45078587fbc41


[ROCm/clr commit: 96de030c2d]
2016-08-17 16:28:22 +05:30
pensun 4353cefd6e add occupancy support for NV path; fix hipPeekAtLastError on HCC path
Change-Id: I26b0e1875c19d7c636ffcc18f1738926572ded81


[ROCm/clr commit: e70cfd3582]
2016-08-16 16:25:03 -05:00
Aditya Atluri 5513a08601 Added kernel compilation driver apis
1. Added 2 new driver apis, hipModuleLoad, hipModuleGetFunction

Change-Id: If464a7fad178121e3da791c7ac9e17ebc01a9cd0
Issues: When a sample written with them shows Aborted (core dumped) when exiting


[ROCm/clr commit: b8597512ca]
2016-08-16 14:36:25 -05:00
Evgeny Mankov e88b7756b7 #define HIP_DYNAMIC_SHARED_ATTRIBUTE is added
[ROCm/clr commit: 1031f422c2]
2016-08-16 17:58:57 +03:00
Rahul Garg a23ea89cae Implementation of hipCtxGetDevice
Change-Id: I067572e486323c3aad6f744a2c0c4997c8696af6


[ROCm/clr commit: 5eb889f73c]
2016-08-13 01:17:46 +05:30
Rahul Garg b55d509790 First implementation of hipCtxXXX functions
Change-Id: I4609cbe6bd90a1fff8655bff4fdd773864397aba


[ROCm/clr commit: f734f12411]
2016-08-13 00:09:08 +05:30
Jeffrey Poznanovic cea99c239b Adding hipblas include files
Change-Id: I73064d410acd8f655dc62eaeb6f4bdefc5381e35


[ROCm/clr commit: dd680ad86a]
2016-08-12 11:59:25 +05:30
Ben Sander a4bfccf628 Context update.
- Remove tls_deviceID.
- Add first passing test.

Change-Id: If3e2f254abf589028cfe4f9e6369745f04160de0


[ROCm/clr commit: a9bcee3d77]
2016-08-10 08:59:47 -05:00
Rahul Garg dbf9fd8312 Changed StagingBuffer class to UnpinnedCopyEngine
Change-Id: I1e212bfc8030dcf225ecf78fd7b23fda9b1de92f


[ROCm/clr commit: d08b1239ee]
2016-08-09 21:29:42 +05:30
Rahul Garg cb771c60d3 Moved sync copy decision logic to staging buffer class
Change-Id: I5c398772375fcc1f174a7597eea1215ce7bf80b4


[ROCm/clr commit: 4044cb412c]
2016-08-09 09:28:18 +05:30
Ben Sander 0c8ca4b37d Add initial context implementation.
APIs: hipInit, hipCtxCreate.
Track TLS default ctx.  Set deviceID now changes the ctx.
Add first context test.

Change-Id: If1cb9989b5a04a36147e25e84904336c7b6f3d88


[ROCm/clr commit: f19f2248bf]
2016-08-08 17:49:02 -05:00
Ben Sander c85f34bf8a Code cleanup, use camelCase where appropriate.
Change-Id: I5a7ec50df8bbb3e7a3b313c0b12e2dd55ae4a09c


[ROCm/clr commit: ac173c3ae8]
2016-08-08 14:54:38 -05:00
Ben Sander e6dd8281d0 Move copy kernel templates into hip_memory.cpp
Change-Id: I862529f3fa8232372c6bacaa5d36f035bbdd32a1


[ROCm/clr commit: 6dbe554894]
2016-08-08 12:07:12 -05:00
Ben Sander 3ee6e7f8ab Split ihipCtx_t into ihipCtx_t and ihipDevice_t .
Major change to existing code base.
    Ctx holds streams, enables peers, and flags.
    Device holds accelerator, hsa-agent, device props.

Add hipCtx_t.

Add peer APIs that accept hipCtx_t (in addition to deviceId)

Compiles and passes directed tests.

Change-Id: Iddab1eb9edbf90caad2ef5959c6b811d658197f1


[ROCm/clr commit: d09b19bb6c]
2016-08-08 11:55:57 -05:00
Ben Sander e5300e2b40 Change Device->Ctx
Change ihipDevice_t -> ihipCtx_t (new)
Change ihipGetTlsDefaultDevice->ihipGetTlsDefaultCtx
Some other changes from device->ctx where appropriate.

Change-Id: I5c4ae93b2fd42c6303aa23d748eb166b7431925d


[ROCm/clr commit: 0d16565061]
2016-08-07 21:47:12 -05:00
Ben Sander aa43c9e006 Remove ihipStream_r::_device_index
Replace with direct pointer to device.  Cleaner, and prep
for transition to contexts.

Change-Id: I0e550f34412923d46c541c0a14bb7d29c3fd4b11


[ROCm/clr commit: 3c604b6430]
2016-08-07 20:47:06 -05:00
Rahul Garg 7b7f9dab90 Region based apis to pool based api changes
Change-Id: If53019eebafe051ab4e811863995f78315297080


[ROCm/clr commit: 52712e2134]
2016-08-05 15:05:57 +05:30
Ben Sander a853d10446 Cleanup sync code.
Remove dead depFutures, enqueueBarrier call.
Rename some parms to reflect usage.
Add comments to better explain tricky parts of sync code.

Change-Id: I763296421d9c2b3b58fc8cef5f010b12ab49553c


[ROCm/clr commit: ef61aae878]
2016-07-27 18:31:11 -05:00
Aditya Atluri 96fe085832 Signal Fix: Added signal limit to allocSignal
1. Did not change the logic in allocSignal
2. Added guard to wait on signal limit

Change-Id: I78f29097e6a584b3c3d78319dac19869067bd1fe


[ROCm/clr commit: 1b2a24d0b8]
2016-07-27 13:48:49 -05:00
Aditya Atluri b99410bd07 Signal Fix: Moved kernel count to critical stream
1. Added environment variable HIP_NUM_KERNELS_INFLIGHT
2. Moved kernelcount variable inside stream critical section

Change-Id: I51d24d0a2a109467209170de117a6d02ba4e308e


[ROCm/clr commit: 7be196de48]
2016-07-26 17:09:27 -05:00
Aditya Atluri d8ea207f61 Signal Fix: Changed global signal count to per stream signal count
1. The number of kernels that can use signals are increased to 128
2. The kernel count is now specific to the stream

Change-Id: Ie6d1aa3f437aad8f08c3333fe48bd3f46e551e60


[ROCm/clr commit: 2e754d27dc]
2016-07-26 14:03:51 -05:00
Aditya Atluri e207805f81 Added re-fix for memcpy kernel sync
1. The patch uses HIP signal pools to sync between copy and kernel commands
2. The hsa_signal_create is removed
3. Left the redundant enqueueBarrier method just in case

Change-Id: I3dff3e8ee57fff3cd49bec802ff735ed128e5ca1


[ROCm/clr commit: 0232e6bbb4]
2016-07-26 09:22:59 -05:00
Rahul Garg 7a70966cb1 D2H and H2D unpinned memory transfer support
Change-Id: If6d6c970f435e5d917d5cc6cddc2ee2918cd1c37

Conflicts:
	src/hip_hcc.cpp


[ROCm/clr commit: d11d65d401]
2016-07-25 14:36:07 +05:30
Aditya Atluri 776ea7a1f0 Partial fix async after kernel launch signal issue
Change-Id: Ib48d6564379160035bded9493b93663fba361710


[ROCm/clr commit: 1704006bed]
2016-07-23 14:54:20 -05:00
pensun 2c945dba23 Add empty stubs for threadfence family routines, changes include:
- 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/clr commit: 6db08e5135]
2016-07-22 10:40:58 -05:00
Maneesh Gupta 64d0ffb404 Replace calls to ihipInit with use of HIP_INIT_API macro
Change-Id: Iabf7df79f0238a8ddffea4607fe945df36642850


[ROCm/clr commit: b485470819]
2016-07-22 15:46:55 +05:30
Maneesh Gupta 68b1a199c4 Fix using ATP markers
Change-Id: If2d04f80b580237426c569737551e2001a8cd35a


[ROCm/clr commit: dffed956fb]
2016-07-21 16:02:51 +05:30
Maneesh Gupta 997cb27be8 Merge branch 'hiparray' into amd-develop
Change-Id: I63ca7b1db7b593ac5cfb3fd7cd5d08d6e4075a4c


[ROCm/clr commit: 7d5cffdc17]
2016-07-21 12:29:56 +05:30
Maneesh Gupta d03904578c Merge branch 'amd-master' into amd-develop
[ROCm/clr commit: 5947731054]
2016-07-05 21:40:22 +05:30
Aditya Atluri c17209f530 added more nvcc event functions
Change-Id: I79ee20ef444d4c1ab6ada3c0d56730ce754ab6b6


[ROCm/clr commit: adad6ef71f]
2016-06-30 21:03:19 -05:00
Maneesh Gupta 4e84678f27 Merge branch 'amd-develop' into amd-master
Change-Id: I04f85b207e15e66c1a546675dc0937726ee08362


[ROCm/clr commit: 2577b6158f]
2016-06-30 18:36:07 +05:30
Aditya Atluri cf3186cd63 moved half support to a source file
Change-Id: I7c09b41877e22c1b743dea25a585e5307427dafd


[ROCm/clr commit: 38720f8a4e]
2016-06-30 18:23:29 +05:30
Aditya Atluri 29c5bf181a added fp16 software support
Change-Id: Ic0fdd9f8248a66911169fc00d3af71f50b36e233


[ROCm/clr commit: 4a6ebce880]
2016-06-30 18:23:29 +05:30
7SK 308a0f7c31 NVCC_COMPAT
add support for both cuda compatible implementation and hcc(faster)
implementation with test

Change-Id: I79a22344f458391d7dffac5f147619a542e97e4e

[ROCm/clr commit: 54034e5048]
2016-06-28 09:36:06 +05:30
Ben Sander 5edb4185a2 Use GRID_LAUNCH_VERSION define provided by HCC grid_launch.h
Change-Id: I5a7b31f13c4055cdca3a8854728d59db7dc5a7bb


[ROCm/clr commit: 5fecff5a59]
2016-06-21 23:52:25 -05:00
Ben Sander 050430789c Grid-launch updates to 2.0 and cleanup of old.
_ 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/clr commit: e27b5cc927]
2016-06-21 23:24:38 -05:00
Ben Sander 5e376bf121 Switch to hc_* coordinate builtins (replace amp_)
Change-Id: I0a8871f0c9f047eb45a7391fd032100af2bbd4e0


[ROCm/clr commit: fefc3f3f1e]
2016-06-21 23:22:38 -05:00
Maneesh Gupta 22edc8805a Merge branch 'amd-master' into amd-develop
[ROCm/clr commit: d3a0de99a1]
2016-06-21 15:14:04 +05:30
Maneesh Gupta 72feba608b Merge "Add compiler annotations to make_ functions" into amd-master
[ROCm/clr commit: 47b26ff5f0]
2016-06-21 05:34:18 -04:00
Aditya Atluri adace999e7 able to pass non-dim launch parm to kernel launch
Change-Id: I0411849a27efcba597a1a9aa08be179635e04988


[ROCm/clr commit: ffcfc95360]
2016-06-18 11:28:20 -05:00
Ben Sander a03c51dc9e Clean up old work-week and USE_* refs
Change-Id: I929c979fa085f8e5205194cbccca46e9b5516aa9


[ROCm/clr commit: 44d117ba63]
2016-06-17 15:18:57 -05:00
Aditya Atluri 1e390cd4af added tests for host math functions
Change-Id: I66a5c574a27190e32054586f07ecf20e1ff71292


[ROCm/clr commit: ba262ea855]
2016-06-17 15:05:33 -05:00
Maneesh Gupta 75ecf38bb7 hipLaunchKernel supports empty kernel with no arguments
Change-Id: Iaa46ced7fe0e3d72734d02250dfb48b98bef615d


[ROCm/clr commit: 275c5516b6]
2016-06-17 14:11:46 +05:30
Ben Sander e317f07cb3 Add NVCC implementation for hipPeeakAtLastError
[ROCm/clr commit: 7492351a3c]
2016-06-16 09:38:34 -05:00
Ben Sander 62c5658e93 remove old KERNELBEGIN/KERNELEND macros
[ROCm/clr commit: dbffc7bba2]
2016-06-16 09:38:07 -05:00
Maneesh Gupta 475fd56684 Add compiler annotations to make_ functions
Change-Id: Ic16cfe402dc5cb387b1b63ad311ce8b57ea07361


[ROCm/clr commit: bdfd943b4a]
2016-06-16 16:18:35 +05:30
Ben Sander c46b2e7450 NVCC improvements.
- 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


[ROCm/clr commit: 6a2a140f34]
2016-06-16 06:13:51 +05:30
Aditya Atluri d01b51eab9 added more host functions and tests
Change-Id: I9904e65e14c5479ba33d836c5c0b763cb5af71e3


[ROCm/clr commit: c4e667cf90]
2016-06-15 11:45:19 -05:00
Aditya Atluri 657a1222d7 added erfinv software implementation
Change-Id: Ib1a5584f6c81ab3afa70f7bcbfd7780e156454e3


[ROCm/clr commit: 720fa16355]
2016-06-14 00:09:41 -04:00
Aditya Atluri 570eba42ae added more device functions
Change-Id: I191919060b393772ee442cc19d83479217c5a4ce


[ROCm/clr commit: ae96fe4d12]
2016-06-13 11:55:12 -05:00