Commit gráf

808 Commit-ok

Szerző SHA1 Üzenet Dátum
Aditya Atluri 25be4fb532 Changed how hipEvent_t is typedefed internall
- Mapped hipEvent_t directly to ihipEvent_t* instead of a handle

Change-Id: I5a8bcca0ef962932e0738c03eb1fc914d23022ae
2016-08-25 14:34:41 -05:00
Aditya Atluri f22a3189a3 Added hipModuleGetGlobal and hipModuleLoadData
Change-Id: Iaec873f7d86b72911b6ad32e067a4dfe3d552fe6
2016-08-25 14:16:53 -05:00
Evgeny Mankov 5cca5b3dca clang-hipify: code refactoring - API (Driver/Runtime/Blas) distinguishing is added. 2016-08-25 19:36:37 +03:00
Evgeny Mankov 207eb02736 clang-hipify: code refactoring and performance improvement 2016-08-24 18:51:36 +03:00
Aditya Atluri af4dc556eb changed internal structure of hipFunction and hipModule
Change-Id: Ifa343782e29d7e056efc47e56253311013005093
2016-08-24 09:47:11 -05:00
Aditya Atluri 768287f2a2 Module test correction and hipModuleUnload API
- Corrected the hipModule.cpp test to minimal code
- Added hipModuleUnload API
- Added hipModuleUnload API test

Change-Id: I9c40337043d7972a570b795e1bfc104bd2c4d8aa
2016-08-23 14:19:15 -05:00
Aditya Atluri 2f426aaaea Added module api test with gcn binary
Change-Id: I61c3ecc2b34168d10f1a7b15d668630eb2c69c8c
2016-08-23 13:50:19 -05:00
Aditya Atluri c0fbe7891c 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
2016-08-22 14:17:55 -05:00
Aditya Atluri 98c740c51f 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
2016-08-22 10:29:46 -05:00
Rahul Garg b0ba622ed5 Added support for hipCtxSynchronize and hipCtxGetFlags,modified hipDeviceSynchronize
Change-Id: If7bac667a262fa8c0cb3dc93e97f2534855acd07
2016-08-22 16:15:27 +05:30
Rahul Garg 778bb2e78a Added initial draft for performance optimizations, started with unpinned memory transfers
Change-Id: Icbce2aec347d015bc66cc0c08f6193057bf36b4c
2016-08-22 11:37:37 +05:30
Aditya Atluri a362489b27 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
2016-08-19 23:02:04 -05:00
Aditya Atluri 4d278259b9 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
2016-08-19 21:48:23 -05:00
Aditya Atluri 9eaabb507c 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
2016-08-19 08:49:34 -05:00
Maneesh Gupta 8f19a51521 Make it easier to switch between HIP library formats
Change-Id: Id16406795a81f4bb64dbcb76b9b8763ffe59aac6
2016-08-19 13:07:22 +05:30
Maneesh Gupta 78576ab1a8 CMakeLists: use macro for appending to config file
Change-Id: I71ca3769b58b146f95368c2a2f6615c1eb47e121
2016-08-19 12:17:00 +05:30
Evgeny Mankov 00b9782187 clang-hipify: After translating any symbol forcibly include the hip header file in case it wasn’t.
Fixes https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/issues/34
2016-08-18 20:59:51 +03:00
Aditya Atluri 792811fd52 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
2016-08-18 11:26:55 -05:00
Maneesh Gupta 49d8faef54 Enable hipDynamicShared directed test
Change-Id: I31e7e83ecb3e15fb25b63d6bb6fa9291484c9ef5
2016-08-18 14:15:50 +05:30
Maneesh Gupta eb90dddbc4 Fix version related bug in CMakeLists.txt
Change-Id: I31c567575185a4e85f5f24d3f105f7cb1beed425
2016-08-18 13:56:25 +05:30
Maneesh Gupta 7b8f091f0b Rename 2_Advanced to 7_Advanced
Change-Id: I51e5fa7f4c1dbf467f2d7182ec69d12d5fe548d0
2016-08-18 12:40:30 +05:30
Maneesh Gupta aa23a003cf Update directed tests with all supported math functions
Change-Id: Id32a94313288e78bf2957bd19efb30877e20221d
2016-08-18 12:38:25 +05:30
Maneesh Gupta 260c0dc4ab Fix normcdf signature
Change-Id: I36b225cfe03db687f295aeea8a006d535bc14231
2016-08-17 22:07:06 +05:30
Aditya Atluri b6ff7df81e Added copyright to hip_module.cpp file
Change-Id: Ifc5f1e251d5c52a5b59f372b3fada938dbecb34a
2016-08-17 10:36:28 -05:00
Rahul Garg 96de030c2d Added further hipCtxXXX Apis
Change-Id: I286d962a06cee656c1c652b3f6b45078587fbc41
2016-08-17 16:28:22 +05:30
Aditya Atluri a3cb0f37c4 corrected issues from hipModule API
Change-Id: I9d07884db20df5632f5a69b1a89a0e6ca531712b
2016-08-16 16:49:42 -05:00
pensun e70cfd3582 add occupancy support for NV path; fix hipPeekAtLastError on HCC path
Change-Id: I26b0e1875c19d7c636ffcc18f1738926572ded81
2016-08-16 16:25:03 -05:00
Aditya Atluri b8597512ca 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
2016-08-16 14:36:25 -05:00
Evgeny Mankov 1031f422c2 #define HIP_DYNAMIC_SHARED_ATTRIBUTE is added 2016-08-16 17:58:57 +03:00
Maneesh Gupta 85784149fa Several improvements to hipcc, hipconfig and CMakeLists.txt
- moved hip base version information back to hipconfig
 - fixed bug in hip patch version generation
 - renamed .hipconfig to .buildInfo
 - HCC_HOME is logged to .buildInfo only when HIP_DEVELOPER environment
 - variable is defined
 - hipcc and hipconfig require perl 5.10.1 or above
 - replaced unless defined with logic-defined or
 - added ROCM_TARGET for carrizo
 - moved config parsing to a subroutine
 - config parsing does not set VALID key anymore
 - hipcc honors HCC_HOME environment variable above buildInfo
 - hip_base package now bundles .version in bin directory

Change-Id: Iaeea6d7529671220b02e07337946aaee0af90310
2016-08-14 16:22:25 +05:30
Rahul Garg 5eb889f73c Implementation of hipCtxGetDevice
Change-Id: I067572e486323c3aad6f744a2c0c4997c8696af6
2016-08-13 01:17:46 +05:30
Rahul Garg f734f12411 First implementation of hipCtxXXX functions
Change-Id: I4609cbe6bd90a1fff8655bff4fdd773864397aba
2016-08-13 00:09:08 +05:30
Maneesh Gupta 346ae7bc42 Updates to HIP version string format
HIP patch version is now a monotonically increasing number

Change-Id: Ia6e35794b702bbd2018c502564d0a65997819687
2016-08-12 23:21:37 +05:30
Maneesh Gupta 7c03a0604a Add simple hipblas saxpy sample
Change-Id: I67ae83e1e5397d5191a3c644aba068f06ff97830
2016-08-12 13:50:22 +05:30
Jeffrey Poznanovic dd680ad86a Adding hipblas include files
Change-Id: I73064d410acd8f655dc62eaeb6f4bdefc5381e35
2016-08-12 11:59:25 +05:30
Aditya Atluri 5a47758033 Added fix for using HCC_VERSION
1. The variable is brought outside the conditional so that its scope is increased

Change-Id: I2d2689553e67930050fe5b3648739f0f72c3bbc8
2016-08-11 16:13:44 -05:00
Aditya Atluri df1dd53453 Change hipcc to take HCC_HOME from hipconfig by default
1. Current implementation checks both env var and value in hipconfig and reports error
2. New implementation gives value in hipconfig with highest priority
3. If hipconfig is not present, fall back to env variables.

To Devs: No need to switch between environment variables for different HCC + different HIP.

Change-Id: I6cdf37e1429d7f07be3a68c7e5ba1533d832962b
2016-08-11 15:31:24 -05:00
Evgeny Mankov aba73d6673 clang-hipify: Add support for nested macro expansion and translation.
Fixes bug “HIPIFY: nested macro is not hipified”
https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/issues/33

Example:
#include "cuda_runtime.h"
#define MY_MACRO(func, flags) (func, flags)
...
cudaEvent_t *event = NULL;
MY_MACRO(cudaEventCreateWithFlags(event, cudaEventDisableTiming), NULL);

where cudaEventDisableTiming is a defined numeric literal and thus a nested MACRO:
#define cudaEventDisableTiming 0x02 /**< Event will not record timing data */

After hipifying now:
MY_MACRO(hipEventCreateWithFlags(event, cudaEventDisableTiming), NULL);

Should be:
MY_MACRO(hipEventCreateWithFlags(event, hipEventDisableTiming), NULL);
2016-08-11 22:29:55 +03:00
pensun 3b268c54f8 add note in hip_faq regarding workaround that add keyword of static for all forceinline functions
Change-Id: Ia13ba59b1e54df8ead5a96a952084144431ec72a
2016-08-10 11:31:13 -05:00
Ben Sander a9bcee3d77 Context update.
- Remove tls_deviceID.
- Add first passing test.

Change-Id: If3e2f254abf589028cfe4f9e6369745f04160de0
2016-08-10 08:59:47 -05:00
Ben Sander e23cd0dd3c Document workaround for parenthesis+macro+hipLaunchKernel
Change-Id: Ie04c99db92d6499ddde93028a96f9d8f72d3f992
2016-08-10 08:59:47 -05:00
Maneesh Gupta 5fe7159baf Allow cmake to be run multiple times in directed tests
Change-Id: I9d68fdefd9f72895ad4bdb310fcf3c6e52dbbf02
2016-08-10 09:49:07 +05:30
Maneesh Gupta fcafb975a6 Fix build issues due to refactoring changes
Change-Id: I0a709ff4864244ba1b49e1a25327e3901ed6e17f
2016-08-09 22:49:32 +05:30
Rahul Garg d08b1239ee Changed StagingBuffer class to UnpinnedCopyEngine
Change-Id: I1e212bfc8030dcf225ecf78fd7b23fda9b1de92f
2016-08-09 21:29:42 +05:30
Rahul Garg 4044cb412c Moved sync copy decision logic to staging buffer class
Change-Id: I5c398772375fcc1f174a7597eea1215ce7bf80b4
2016-08-09 09:28:18 +05:30
Ben Sander f19f2248bf Add initial context implementation.
APIs: hipInit, hipCtxCreate.
Track TLS default ctx.  Set deviceID now changes the ctx.
Add first context test.

Change-Id: If1cb9989b5a04a36147e25e84904336c7b6f3d88
2016-08-08 17:49:02 -05:00
Ben Sander ac173c3ae8 Code cleanup, use camelCase where appropriate.
Change-Id: I5a7ec50df8bbb3e7a3b313c0b12e2dd55ae4a09c
2016-08-08 14:54:38 -05:00
Ben Sander 7c5a611df4 Coding guidelines update
Change-Id: Ib8d8da4c3897d157aeb26eb2e99718d66fd260b1
2016-08-08 13:12:22 -05:00
Ben Sander 6dbe554894 Move copy kernel templates into hip_memory.cpp
Change-Id: I862529f3fa8232372c6bacaa5d36f035bbdd32a1
2016-08-08 12:07:12 -05:00
Ben Sander d09b19bb6c 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
2016-08-08 11:55:57 -05:00