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

766 Коммитов

Автор SHA1 Сообщение Дата
Jeffrey Poznanovic cbe87288c8 Adding hipblas include files
Change-Id: I73064d410acd8f655dc62eaeb6f4bdefc5381e35


[ROCm/hip commit: 48491a3978]
2016-08-12 11:59:25 +05:30
Aditya Atluri 25fa1336e6 Added fix for using HCC_VERSION
1. The variable is brought outside the conditional so that its scope is increased

Change-Id: I2d2689553e67930050fe5b3648739f0f72c3bbc8


[ROCm/hip commit: 3be747c41e]
2016-08-11 16:13:44 -05:00
Aditya Atluri fd564cc04c 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


[ROCm/hip commit: ef68f2f293]
2016-08-11 15:31:24 -05:00
Evgeny Mankov 4596f44fbe 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);


[ROCm/hip commit: 2aacb02358]
2016-08-11 22:29:55 +03:00
pensun 6322747871 add note in hip_faq regarding workaround that add keyword of static for all forceinline functions
Change-Id: Ia13ba59b1e54df8ead5a96a952084144431ec72a


[ROCm/hip commit: 4553e4e7f7]
2016-08-10 11:31:13 -05:00
Ben Sander 0221323977 Context update.
- Remove tls_deviceID.
- Add first passing test.

Change-Id: If3e2f254abf589028cfe4f9e6369745f04160de0


[ROCm/hip commit: 89164259ab]
2016-08-10 08:59:47 -05:00
Ben Sander c9c8c1323f Document workaround for parenthesis+macro+hipLaunchKernel
Change-Id: Ie04c99db92d6499ddde93028a96f9d8f72d3f992


[ROCm/hip commit: 1786b120ed]
2016-08-10 08:59:47 -05:00
Maneesh Gupta 8e648b1a33 Allow cmake to be run multiple times in directed tests
Change-Id: I9d68fdefd9f72895ad4bdb310fcf3c6e52dbbf02


[ROCm/hip commit: 2e9adefd71]
2016-08-10 09:49:07 +05:30
Maneesh Gupta c0398a8de6 Fix build issues due to refactoring changes
Change-Id: I0a709ff4864244ba1b49e1a25327e3901ed6e17f


[ROCm/hip commit: 76eeaf4fbb]
2016-08-09 22:49:32 +05:30
Rahul Garg e64ae4bda0 Changed StagingBuffer class to UnpinnedCopyEngine
Change-Id: I1e212bfc8030dcf225ecf78fd7b23fda9b1de92f


[ROCm/hip commit: 2ac93c340d]
2016-08-09 21:29:42 +05:30
Rahul Garg 99f4d4ebec Moved sync copy decision logic to staging buffer class
Change-Id: I5c398772375fcc1f174a7597eea1215ce7bf80b4


[ROCm/hip commit: 023b1ecf33]
2016-08-09 09:28:18 +05:30
Ben Sander db8c7e97bd 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/hip commit: 8f402132ba]
2016-08-08 17:49:02 -05:00
Ben Sander 28db2f79ba Code cleanup, use camelCase where appropriate.
Change-Id: I5a7ec50df8bbb3e7a3b313c0b12e2dd55ae4a09c


[ROCm/hip commit: ed0a2c02fe]
2016-08-08 14:54:38 -05:00
Ben Sander a671758882 Coding guidelines update
Change-Id: Ib8d8da4c3897d157aeb26eb2e99718d66fd260b1


[ROCm/hip commit: b1d8f9d00d]
2016-08-08 13:12:22 -05:00
Ben Sander e8757f8154 Move copy kernel templates into hip_memory.cpp
Change-Id: I862529f3fa8232372c6bacaa5d36f035bbdd32a1


[ROCm/hip commit: 2a798152d4]
2016-08-08 12:07:12 -05:00
Ben Sander 57df8a3967 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/hip commit: cfdacab32f]
2016-08-08 11:55:57 -05:00
Ben Sander 1f9d2201fe Add initial/partial coding guidelines
Change-Id: Ifd8cb3ad74b15d3ab2f38c3daa038a2808af6fa9


[ROCm/hip commit: 6aeb2dc8d6]
2016-08-08 11:55:41 -05:00
Ben Sander ee356ad0d4 Change Device->Ctx
Change ihipDevice_t -> ihipCtx_t (new)
Change ihipGetTlsDefaultDevice->ihipGetTlsDefaultCtx
Some other changes from device->ctx where appropriate.

Change-Id: I5c4ae93b2fd42c6303aa23d748eb166b7431925d


[ROCm/hip commit: 2dc3d3238b]
2016-08-07 21:47:12 -05:00
Ben Sander 308578d520 Remove ihipStream_r::_device_index
Replace with direct pointer to device.  Cleaner, and prep
for transition to contexts.

Change-Id: I0e550f34412923d46c541c0a14bb7d29c3fd4b11


[ROCm/hip commit: e7d7c5cbe8]
2016-08-07 20:47:06 -05:00
Evgeny Mankov d5d642d7bb 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);


[ROCm/hip commit: 114d5bfddf]
2016-08-05 21:35:58 +03:00
Rahul Garg 63c4dc3273 Region based apis to pool based api changes
Change-Id: If53019eebafe051ab4e811863995f78315297080


[ROCm/hip commit: fcb2fcce1e]
2016-08-05 15:05:57 +05:30
Maneesh Gupta ea3ad49a8c .hipconfig now lives in lib
Also moved reading .hipconfig to begining of hipcc

Change-Id: Ibb9d32bef64a79f189aa037c36814759c8cc8052


[ROCm/hip commit: bc394505cc]
2016-08-03 12:01:48 +05:30
scchan 41da8068f1 compile all hip objects with -fPIC so that they work in a shared library
Change-Id: I1f8355d4a81dbd5f408062c317c16a5efc668876


[ROCm/hip commit: 2f6ac8c36e]
2016-08-03 11:48:46 +05:30
Maneesh Gupta 8c98963134 hipcc reads .hipconfig
hipcc dies if HCC_VERSION is not same as the one used to build HIP

Change-Id: I180c0108812fe5cb6c5304477557c524a4a61f75


[ROCm/hip commit: 703a287bb0]
2016-08-03 11:32:08 +05:30
Maneesh Gupta 0229cddf67 Only create .hipconfig on install
Change-Id: I9c20440023401f5794b33a0165e0918372783a68


[ROCm/hip commit: c38d0466b8]
2016-08-03 11:25:59 +05:30
Maneesh Gupta 5448e4d280 Merge branch 'release_0.92.00' into amd-develop
[ROCm/hip commit: de39d01922]
2016-08-03 09:10:44 +05:30
Maneesh Gupta 960d1ec4cc Fix hipcc to work correctly when HIP_ATP_MARKER is not defined
Change-Id: I7cc525daccf896704e6ccf6d04ed395fda622031


[ROCm/hip commit: e818f2378d]
2016-08-03 09:08:40 +05:30
Maneesh Gupta 7bab8bfb27 hip_hcc rpm package does not depend on rocm-profiler
Change-Id: I291c046371e704ce395781d5de3f6430c0dcba7b


[ROCm/hip commit: b2035c98ac]
2016-08-03 09:07:58 +05:30
Aditya Atluri 8c104aaeeb Added configuration file to be used by hipcc
1. CMake will create .hip-config file in bin directory
Future Work: Need to make changes to hipcc to read the file

Change-Id: Ia7dc48d43787921d5af4ab07d7a5befbcf904465


[ROCm/hip commit: 9c45d9eaed]
2016-08-02 15:45:46 -05:00
Ben Sander 5ffc0a1e25 Remove faulty assert for kernelCnt==0
Change-Id: I8a925c95f48e857c0a31f44561499e90dc6df552


[ROCm/hip commit: f43d02027e]
2016-08-01 13:38:47 -05:00
Maneesh Gupta 7ea77fd1b6 Merge branch 'release_0.92.00' into amd-develop
Conflicts:
	RELEASE.md
	docs/markdown/hip_faq.md

Change-Id: Ifae1b64b6255a7872dfdc885bb8fb52f622464b7


[ROCm/hip commit: 03df23df2b]
2016-08-01 10:47:25 +05:30
Aditya Atluri 2f53a50718 Signal Fix: The signals in a stream are re-used
1. Before, the signal pool is increased depending on the usage
2. After, a static number of signals are allocated to the pool
Only these are used by hip in a stream
3. If the signals required are more than the pool size, the
stream has to wait to make sure all the signals are available
4. Once they are available, the stream can use them
5. Removed HIP_NUM_SIGNALS_PER_STREAM because of redundancy with HIP_STREAM_SIGNALS
6. Increased signal count from 2 to 32.
Future Work: Dynamically increase the pool size depending on the number of
streams allocated by the application. And, null stream should have more signals

Change-Id: I6be36e084f26bb04766fabf776c7210aee0f9e91


[ROCm/hip commit: 9062ebcf3a]
2016-07-28 23:01:35 -05:00
Ben Sander 3044711042 Remove dead enqueueBarrier function.
Change-Id: Ib18fe6bd96ce24dbeb342961ddb5721f7d03f2b2


[ROCm/hip commit: 666c227c7d]
2016-07-28 22:48:22 -05:00
Ben Sander 91368967bb 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/hip commit: 02dd7a7399]
2016-07-27 18:31:11 -05:00
Ben Sander 07c9fb8b38 Fix API string message for hipDeviceGetAttribute
Change-Id: I30f54627630c8ee835506be8c9921742bb68a43a


[ROCm/hip commit: 4cf0f86e99]
2016-07-27 16:18:14 -05:00
Aditya Atluri c2b4a3936c 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/hip commit: 1859c6e515]
2016-07-27 13:48:49 -05:00
Maneesh Gupta 1acf8919ee Update release notes for 0.92.00 release
Change-Id: I9ca588cd0d5d752dc6521e76ba943500eb55525f


[ROCm/hip commit: 059a8d51ba]
2016-07-27 20:30:04 +05:30
Aditya Atluri 6ff74b7780 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/hip commit: 0a31b47e2e]
2016-07-26 17:09:27 -05:00
Aditya Atluri eccab29360 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/hip commit: 53d7629a85]
2016-07-26 14:03:51 -05:00
Aditya Atluri 06abacce98 removed redundant signal destroy
Change-Id: Icf0cd76b2620d34c87cfb6c7a83049087c0a0bc4


[ROCm/hip commit: fa7933eb91]
2016-07-26 13:35:35 -05:00
Ben Sander 3027a1f269 Doc update for FAQ and future RELEASE notes
Change-Id: I7e7c32d9a19fdaea0a0e41f1d4fa4652e53640f2


[ROCm/hip commit: 99bc0f03e1]
2016-07-26 12:14:15 -05:00
Ben Sander 46cb4cdc7b Doc update for README.md - add more intro text, example
Change-Id: I99b8eaacd6460dfdbdbc8ddba3fe589647d877e7


[ROCm/hip commit: 6ea9d2eace]
2016-07-26 12:14:02 -05:00
Ben Sander b909cbb2a3 Make HCC ignore register keyword
- (previously would emit a warning)
- Also tweak documentation.

Change-Id: I0f4f00f82f8cc53d420112570f2d7675535e6aea


[ROCm/hip commit: 0a74e00ac8]
2016-07-26 12:13:48 -05:00
Aditya Atluri 69aed2c6ea 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/hip commit: 4bdf26a82e]
2016-07-26 09:22:59 -05:00
Maneesh Gupta adfaa47840 Documentation updates
Change-Id: Ia624d86915c4c96da0ac0242f767135f30ff73c6


[ROCm/hip commit: 90a3e4e1da]
2016-07-25 14:53:15 +05:30
Rahul Garg d506d58919 D2H and H2D unpinned memory transfer support
Change-Id: If6d6c970f435e5d917d5cc6cddc2ee2918cd1c37

Conflicts:
	src/hip_hcc.cpp


[ROCm/hip commit: 42a3ed544c]
2016-07-25 14:36:07 +05:30
Aditya Atluri 6f95ad0a8d Partial fix async after kernel launch signal issue
Change-Id: Ib48d6564379160035bded9493b93663fba361710


[ROCm/hip commit: c756bb3398]
2016-07-23 14:54:20 -05:00
pensun 2e5a6ab630 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/hip commit: f31668fee4]
2016-07-22 10:40:58 -05:00
Maneesh Gupta f02201e8c3 Replace calls to ihipInit with use of HIP_INIT_API macro
Change-Id: Iabf7df79f0238a8ddffea4607fe945df36642850


[ROCm/hip commit: 71d51170ef]
2016-07-22 15:46:55 +05:30
Maneesh Gupta 0fd69d97f4 hip_hcc package now depends on rocm-profiler
Change-Id: I80d1c6048cc18c47c2024efb90368b17139a09ad


[ROCm/hip commit: 1bf661223e]
2016-07-22 15:45:32 +05:30