Commit Graph

798 Commits

Author SHA1 Message Date
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
Ben Sander 694c0bad62 Add initial/partial coding guidelines
Change-Id: Ifd8cb3ad74b15d3ab2f38c3daa038a2808af6fa9
2016-08-08 11:55:41 -05:00
Ben Sander 0d16565061 Change Device->Ctx
Change ihipDevice_t -> ihipCtx_t (new)
Change ihipGetTlsDefaultDevice->ihipGetTlsDefaultCtx
Some other changes from device->ctx where appropriate.

Change-Id: I5c4ae93b2fd42c6303aa23d748eb166b7431925d
2016-08-07 21:47:12 -05:00
Ben Sander 3c604b6430 Remove ihipStream_r::_device_index
Replace with direct pointer to device.  Cleaner, and prep
for transition to contexts.

Change-Id: I0e550f34412923d46c541c0a14bb7d29c3fd4b11
2016-08-07 20:47:06 -05:00
Evgeny Mankov b7ac63e202 clang-hipify: Transformation of declarations with external linkage and shared attribute for IncompleteArrayType (aka C array[]) only.
Example:
extern __shared__ uint sRadix1[]; =>  HIP_DYNAMIC_SHARED(unsigned int, sRadix1);
2016-08-05 21:35:58 +03:00
Rahul Garg 52712e2134 Region based apis to pool based api changes
Change-Id: If53019eebafe051ab4e811863995f78315297080
2016-08-05 15:05:57 +05:30
Maneesh Gupta 6a2e92f7c8 .hipconfig now lives in lib
Also moved reading .hipconfig to begining of hipcc

Change-Id: Ibb9d32bef64a79f189aa037c36814759c8cc8052
2016-08-03 12:01:48 +05:30
scchan 992195b395 compile all hip objects with -fPIC so that they work in a shared library
Change-Id: I1f8355d4a81dbd5f408062c317c16a5efc668876
2016-08-03 11:48:46 +05:30
Maneesh Gupta 1381681335 hipcc reads .hipconfig
hipcc dies if HCC_VERSION is not same as the one used to build HIP

Change-Id: I180c0108812fe5cb6c5304477557c524a4a61f75
2016-08-03 11:32:08 +05:30
Maneesh Gupta b732fd74b4 Only create .hipconfig on install
Change-Id: I9c20440023401f5794b33a0165e0918372783a68
2016-08-03 11:25:59 +05:30
Maneesh Gupta 25c65c5fc4 Merge branch 'release_0.92.00' into amd-develop 2016-08-03 09:10:44 +05:30