Граф коммитов

265 Коммитов

Автор SHA1 Сообщение Дата
Aditya Atluri df6f400a8e added doxygen comments for hipModuleLaunchKernel
Change-Id: I8a52d2e62f4b7eea8e05d779b9fda49f0ac45130


[ROCm/clr commit: 6bfbf17fb6]
2016-08-31 10:18:46 -05:00
Rahul Garg cf294e9b60 CUresult to hipError mapping updates
Change-Id: I602a70acda798a47ebbbef84b438b21c399821c3


[ROCm/clr commit: e0c5ad6ff6]
2016-08-30 15:55:51 +05:30
Aditya Atluri 50a2e2db2a Added doxygen comments for module apis
Change-Id: I1825249bf91efe7d058f9026a82ec47855759c98


[ROCm/clr commit: d87bfd425c]
2016-08-29 21:24:19 -05:00
Ben Sander 78cfa5f6d6 Refactor trace code for hipLaunchKernel.
- 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/clr commit: fa9b95838a]
2016-08-29 18:37:57 -05:00
Aditya Atluri 9af7591223 added hipInit for cuInit nvcc path
Change-Id: I594d08d936ac5d06f16e42c63062ac5776b65a0c


[ROCm/clr commit: 5d3c72a7ff]
2016-08-29 15:20:09 -05:00
Aditya Atluri 41a363a6b7 Changed module api to work with nvcc and hipcc path
- Added cuda and cudart linking for nvcc path in hipcc
- Added hipMemcpyHtoD and hipMemcpyDtoH for nvcc path
- Changed hipDeviceptr to hipDeviceptr_t
- Fixed hipMemcpy*to* API signatues

Change-Id: I6ef076b507f92502efda475c83dcdcdd462afc37


[ROCm/clr commit: fb0c432095]
2016-08-29 15:05:12 -05:00
Aditya Atluri 25fe3f0777 Fixed nvcc path for module apis
Change-Id: I43c7a6bc226f4f270c37f3c4af86b3b3efd0b175


[ROCm/clr commit: 645975a05b]
2016-08-28 22:50:44 -05:00
Aditya Atluri 9fe6924154 Cleaned up module api
- Moved HIP_INIT_API from internal to all public apis
- renamed hipLaunchModuleKernel to hipModuleLaunchKernel
- Changed tests according to the changed api

Change-Id: I822ff63c7c5b7dad340ece49456baf9d89428e9f


[ROCm/clr commit: 71991ed8f8]
2016-08-28 16:48:57 -05:00
Ben Sander 40ca19440c Remove deprecated message from hipLaunchModuleKernel
Change-Id: I87675453ae4363e3340a9d1491bb00543fa8c6e0


[ROCm/clr commit: f7ee14e637]
2016-08-28 16:34:34 -05:00
Aditya Atluri 6a596c353a added memset and memcpy device functions
- Added memcpy and memset device functions
- Added test for memcpy and memset

Change-Id: Icd21a8dd964953b86d5e92889bf1664bee647219


[ROCm/clr commit: 5e7c396bcd]
2016-08-27 11:13:56 -05:00
Aditya Atluri fa734c5a6a Added explicit memory copy direction apis
- Fixed stale printf in context api
- Added 4 sync memcpy apis
  1. hipMemcpyHtoD
  2. hipMemcpyDtoH
  3. hipMemcpyDtoD
  4. hipMemcpyHtoH
- Added test for added apis

Change-Id: I4a9c382445b62631f8d0bcbb9a670322288b72b1


[ROCm/clr commit: f722a132a4]
2016-08-26 13:11:01 -05:00
Aditya Atluri 33a38a0e26 Added NVCC support and name changes
- Added NVCC support for module APIs
- Changed hipFunction and hipModule data types to hipFunction_t and hipModule_t
- Created new intenal ihipModuleGetFunction as it is used twice
- Changed test to match with the new data types

Change-Id: I300a1c7fd40ed7065b1b8b9de97e3a06b96ed729


[ROCm/clr commit: 74a8fb132d]
2016-08-26 10:32:01 -05:00
Rahul Garg c43c365811 Added logic to update primary ctx when ctx stack is empty, updated hipCtxDestroy and ctxGetCurrent functions
Change-Id: Ia0a8943c121bc1279788a1cfa9be59af614b04a6


[ROCm/clr commit: 452108b0e6]
2016-08-26 19:03:23 +05:30
Rahul Garg 5830aeb77e Resolved errors due to hipCtxXXX APIs
Change-Id: Iffac0095c4352864eca622ea318d2291571b5153


[ROCm/clr commit: 77991293ba]
2016-08-26 15:32:49 +05:30
Rahul Garg 93e5faf1ce NVCC path support for hipCtxXXX APIs
Change-Id: Ic7dbfbdaee9d00c0de1363c50758e5e29a96a8b2


[ROCm/clr commit: da80bf6326]
2016-08-26 14:10:36 +05:30
Rahul Garg 8e09eacd0b Addition of hipCtxEnablePeerAccess and hipCtxDisablePeerAccess functions
Change-Id: I381c8cbbde17eae7d9bb5d4cb1596cebf4bda039


[ROCm/clr commit: a1d401336b]
2016-08-26 13:51:33 +05:30
Aditya Atluri 82deaeb581 Changed how hipEvent_t is typedefed internall
- Mapped hipEvent_t directly to ihipEvent_t* instead of a handle

Change-Id: I5a8bcca0ef962932e0738c03eb1fc914d23022ae


[ROCm/clr commit: 25be4fb532]
2016-08-25 14:34:41 -05:00
Aditya Atluri d769b71efc Added hipModuleGetGlobal and hipModuleLoadData
Change-Id: Iaec873f7d86b72911b6ad32e067a4dfe3d552fe6


[ROCm/clr commit: f22a3189a3]
2016-08-25 14:16:53 -05:00
Aditya Atluri 069d132f41 changed internal structure of hipFunction and hipModule
Change-Id: Ifa343782e29d7e056efc47e56253311013005093


[ROCm/clr commit: af4dc556eb]
2016-08-24 09:47:11 -05:00
Aditya Atluri 9c36b77d39 Module test correction and hipModuleUnload API
- Corrected the hipModule.cpp test to minimal code
- Added hipModuleUnload API
- Added hipModuleUnload API test

Change-Id: I9c40337043d7972a570b795e1bfc104bd2c4d8aa


[ROCm/clr commit: 768287f2a2]
2016-08-23 14:19:15 -05:00
Aditya Atluri 52305186ef Added stream synchronisation for hipLaunchModuleKernel
- The module kernel launch is now in sync with commands in its stream
- Moved launch kernel inside ihipStream

Change-Id: Ic00cfcf4882bf81b6203c36881a52575ea68b529


[ROCm/clr commit: c0fbe7891c]
2016-08-22 14:17:55 -05:00
Aditya Atluri b9f564ae19 Added nvcc path for hipComplex APIs
- Changed from inline to static inline for hipComplex AMD APIs
- Added NVCC path for hipComplex APIs mapped to cuComplex APIs

Change-Id: I809cf3a11b5b1c8bbc7a57c5fbcc3dc6745ccb95


[ROCm/clr commit: 98c740c51f]
2016-08-22 10:29:46 -05:00
Rahul Garg f1f04949ac Added support for hipCtxSynchronize and hipCtxGetFlags,modified hipDeviceSynchronize
Change-Id: If7bac667a262fa8c0cb3dc93e97f2534855acd07


[ROCm/clr commit: b0ba622ed5]
2016-08-22 16:15:27 +05:30
Aditya Atluri b2301b8912 Added more complex apis and copyright
- New header which redirects to CUDA/HIP path added for hipComplex.h
- Added more complex device api including fma
- Added copyright to new files

Change-Id: Iff0dece4c438e97d0ae33efa4312975d465a6464


[ROCm/clr commit: a362489b27]
2016-08-19 23:02:04 -05:00
Aditya Atluri d5cc978648 Added support for complex device functions
- Added complex number arithmetic operation for float and double datatypes
- TODO: make them host functions and support half
- Added new function which is not in CUDA, hipCsqabs which is square of absolute value

Change-Id: Ib96e194ad45dc64fcba29eb19ad0376542e0591d


[ROCm/clr commit: 4d278259b9]
2016-08-19 21:48:23 -05:00
Aditya Atluri 9571f355ac Added support for executable and symbols for data structures
- symbol handle is added to hipFunction
- executable handle is added to hipModule
- This way, the APIs doesn't need to track the values

Change-Id: I7cf05329cf79fe946319d7746bd9f5503268fda4


[ROCm/clr commit: 9eaabb507c]
2016-08-19 08:49:34 -05:00
Aditya Atluri a0740e524c Added hipLaunchModuleKernel and new error codes
- 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/clr commit: 792811fd52]
2016-08-18 11:26:55 -05:00
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