Wykres commitów

53 Commity

Autor SHA1 Wiadomość Data
Jack Chung bf4ba9e09e Squashed commit of the following:
commit 9548493fa754b3bf5c31cbdc2211db1e73e8c07c
Author: Jack Chung <whchung@gmail.com>
Date:   Mon May 23 11:57:23 2016 +0800

    Rename hipExternShared test to hipDynamicShared

    Change-Id: I180d9d539420fb69cfc121eceaa7db9da03483b2

commit 827081f8244a38f010789d556db0c4ff7b6422d8
Author: Jack Chung <whchung@gmail.com>
Date:   Mon May 23 11:56:27 2016 +0800

    Rename HIP_DECLARE_EXTERN_SHARED to HIP_DYNAMIC_SHARED

    Change-Id: I22362d179812ac547e0f11ba4e2bb999050e08ae

commit 4c277228ed41af187739610fa17eab1fb144c947
Author: Jack Chung <whchung@gmail.com>
Date:   Thu May 19 17:49:52 2016 +0800

    Adopt new interface to get dynamic LDS in hc.hpp

    Change-Id: I47b433b714633a4c97df87c40a0b1d3386429a00

commit 5a36117d777064113a528dc47b42e8c8413baa97
Author: Jack Chung <whchung@gmail.com>
Date:   Thu May 19 11:29:24 2016 +0800

    Add test patterns for regular expression to match "extern __shared__"

    These test patterns should better be saved as an individual test case, but I'm
    not familiar with HIP test structures so I leave them as comments in hipify as
    of now.

    Change-Id: I7fee89c89b9e73de2133357a226ec0c769733531

commit 1b26284168c7f5339f63338fd0149bed5d994656
Author: Jack Chung <whchung@gmail.com>
Date:   Thu May 19 11:25:23 2016 +0800

    Add one HIP unit test to use HIP_DECLARE_EXTERN_SHARED

    Change-Id: I4d9907815920693a74ea9d575fe26e7c67636109

commit 77b816ee5972b13d829d5bbcf06fbfd07acea2af
Author: Jack Chung <whchung@gmail.com>
Date:   Wed May 18 19:18:59 2016 +0800

    Adopt HIP_ prefix for DECLARE_EXTERN_SHARED macro

    Change-Id: I555ded16b449b67d2e20904013d86fe1ded6a2be

commit ef0997939c3578a9ae11621bf21c0416f04d2622
Author: Jack Chung <whchung@gmail.com>
Date:   Wed May 18 17:42:04 2016 +0800

    Modify hipify to support converting extern __shared__ to DECLARE_EXTERN_SHARED macro

    Added regular expression to search & replace extern __shared__ declarations to
    DECLARE_EXTERN_SHARED macro.

    Limitation:
    - Won't work if "extern __shared__" is declared at global scope

    Sample Usages:
    extern __shared__ double foo[];
    extern __shared__ unsigned int foo[];
    extern volatile __shared__ double foo[];
    extern volatile __shared__ unsigned int sdata[];
    extern __shared__ volatile unsigned int sdata[];
    extern __shared__ T s[];
    extern __shared__ T::type s[];
    extern __shared__ blah<T>::type s[];
    extern __shared__ typename mapper<Float>::type s_data[];
    extern __attribute__((used)) __shared__ typename mapper<Float>::type s_data[];

    Change-Id: I2be0b7039adeddb789f5a2b067d403a43fdc3e26

commit 93ff268724493aedfacdcd5a5aa9a100f4ebaed0
Author: Jack Chung <whchung@gmail.com>
Date:   Wed May 18 15:13:09 2016 +0800

    Introduce DECLARE_EXTERN_SHARED macro to encapsulate "extern __shared__" decls

    Change-Id: I93b2d37c763195b0ca9fd0afee78605a1e3272db

commit cff9c95412de343cc6405158b5acc4f1029267ff
Author: Jack Chung <whchung@gmail.com>
Date:   Wed May 18 12:53:54 2016 +0800

    Add __get_dynamic_groupbaseptr() to point to dynamic LDS

    Change-Id: I97b548d8a691488057617c551a8f331cad7afc77

Change-Id: I84e7875b76fa1f59e860e19c93bd4209cdd1fd2c
2016-06-05 06:20:44 -04:00
Rahul Garg 381dc93b96 Support for Atomic inc and dec in HIP
Change-Id: I783e4917cece5cc379894f0d293382315fbfa8b0
2016-05-12 11:10:48 +05:30
Jack Chung 453d07daf4 Fix math functions so they can be used on host codes as well
Add __host__ for some math functions to fix the following tests:

- hipSinglePrecisionMathHost
- hipDobulePrecisionMathHost
2016-05-06 08:09:52 +08:00
Ben Sander db7877320f remove kalmar_defines.h 2016-04-28 17:40:17 -05:00
Ben Sander 558975c046 remove hacks to disable math or LDG functions. 2016-04-28 14:14:22 -05:00
Ben Sander 9026146951 Properly pair #if/#endif 2016-04-28 13:43:32 -05:00
Ben Sander efaa278cbf merge fix 2016-04-28 13:24:40 -05:00
Ben Sander bae6462e3b hipcc uses -stdlib=stdlibc++ by default.
hipcc switches to -stdlib=libc++ if user requests it.
hack to disable LDG.
2016-04-28 13:17:49 -05:00
Jack Chung 6756b8f0bc Fix compilation error in hipSinglePrecisionMathHost test 2016-04-27 17:46:27 -05:00
Aditya Atluri cac8110a4f changed to guard from hc.hpp 2016-04-27 17:46:27 -05:00
Maneesh Gupta ffdf6ab23b Merge branch 'release_0.84.00' into privatestaging
Conflicts:
	include/hcc_detail/hip_runtime.h
	src/hip_hcc.cpp
2016-04-22 10:55:58 +05:30
Aditya Atluri 805b268ad4 added support for __ldg 2016-04-20 12:25:40 -05:00
pensun 883a3fe831 update API trace information for hipLaunchKernel 2016-04-20 09:57:55 -05:00
pensun 2ddade96d7 Merge branch 'release_0.84.00' of https://github.com/AMDComputeLibraries/HIP-privatestaging into release_0.84.00 2016-04-20 09:48:52 -05:00
pensun 7f37fc4ec7 update hipLaunchKernel API trace information 2016-04-20 09:48:35 -05:00
Jack Chung 07efbcf0d0 Fix compilation error in hipSinglePrecisionMathHost test 2016-04-19 23:17:09 +08:00
Maneesh Gupta bc4f388b69 Prefix HIP includes with hip/ 2016-04-19 15:02:12 +05:30
streamhsa 76f2670cef Add __ldg 2016-04-16 21:12:09 +08:00
Aditya Atluri a0aef11da7 changed to guard from hc.hpp 2016-04-15 11:21:45 -05:00
Aditya Atluri 4dfe27e8b5 Removed pragma once 2016-04-12 09:22:03 -05:00
Aditya Atluri d5feabfa1c added C guard to hip_runtime.h 2016-04-11 13:01:02 -05:00
Aditya Atluri 26985b746b added device math functions 2016-03-31 12:29:19 -05:00
Aditya Atluri 72c72d8748 moved cpp specific device code from headers to source 2016-03-31 04:45:00 -05:00
Aditya Atluri 496da0aabd C guarding the apis 2016-03-30 06:08:50 -05:00
Aditya Atluri da1d9caf6b Moved flag to hip_runtime.h 2016-03-30 03:32:39 -05:00
Aditya Atluri 61cd2bb399 Calling intrinsics from global kernel 2016-03-29 11:17:55 -05:00
Aditya Atluri 794007c3e8 Split files based on functionality and changed cmake file 2016-03-24 04:57:30 -05:00
Aditya Atluri 7920fd9a47 moved variables to headers 2016-03-23 12:59:52 -05:00
scchan 0eead76d4e add compiler version guard to the grid launch fix 2016-03-23 11:24:37 -05:00
scchan 56a1863c74 fix builtins mapping for IDs and Dims 2016-03-23 11:08:28 -05:00
Ben Sander 2569e15375 Improve trace API
- Validate compile-time disables.
- Add README.md section explain how to install/use CodeXL tracing
- Add code docs on trace_helper.h
- fix color on hipLaunchKernel to green.
2016-03-23 02:57:52 -05:00
Ben Sander aa03e1264c Enhance HIP trace debug functions.
- Control with HIP_DB=mask (env var).  See src/hip_hcc.cpp for mask
  values:
    #define DB_API    0 /* 0x01 - shortcut to enable HIP_TRACE_API on single switch */
    #define DB_SYNC   1 /* 0x02 - trace synchronization pieces */
    #define DB_MEM    2 /* 0x04 - trace memory allocation / deallocation */
    #define DB_COPY1  3 /* 0x08 - trace memory copy commands. . */
    #define DB_SIGNAL 4 /* 0x10 - trace signal pool commands */
- Combine with HIP_TRACE to see debug with API trace.
- Use colors to distinguish different flows of debug.
- Add define COMPILE_DB_TRACE to allow removing all debug at compile-time
2016-03-06 23:50:52 -06:00
pensun 11ca71bd76 resolve conflicts of doc_update 2016-02-27 15:08:45 -06:00
Ben Sander 4adab7b7ef Merge branch 'memtracker' into privatestaging
Conflicts:
	src/hip_hcc.cpp
2016-02-25 19:38:46 -06:00
gargrahul a2fbf06129 Update for shared atomics support 2016-02-22 16:21:52 +05:30
Ben Sander d33d806a5b Track last command to a stream.
Passing simple tests.
2016-02-20 11:02:07 -06:00
Ben Sander db3a63360b Move warpSize to header, have shuffles use default warpsize. 2016-02-15 05:41:09 -06:00
gargrahul 8c40a4ace4 Removed atomicInc and atomicDec support from HIP 2016-02-10 04:29:55 +05:30
Peng Sun 1fb48b0714 Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into doc_update 2016-02-09 15:08:39 -06:00
Maneesh Gupta 3291e0ec96 Move HIP_DEVICE_COMPILE defines to hip_common.h 2016-02-09 10:57:20 +05:30
Ben Sander 9e2c3c8df3 minor doc touchup 2016-02-08 22:11:11 -06:00
Peng Sun fb3b11774b fix merging conflicts 2016-02-08 15:35:49 -06:00
Peng Sun 3d5608ea84 Additional typo and extra space fix 2016-02-03 09:42:16 -06:00
scchan 265c42500f add inline attribute to shfl functions 2016-02-02 12:53:17 -06:00
streamhsa 974d491902 Adjusted the value of __any as per CUDA -sandeep 2016-02-02 15:25:42 +05:30
streamhsa 23904df99b ADDED Support for __ffs() and __ffsll() having signed input -sandeep 2016-02-02 15:05:46 +05:30
scchan 04f3e3e598 adding shfl, shfl_up, shfl_down, shfl_xor intrinsics 2016-02-01 23:55:31 -06:00
Maneesh Gupta 861cba6f75 Add double and integer intrinsics to test 2016-02-01 16:00:45 +05:30
Maneesh Gupta d2c6125a7c Add few more single precision intrinsics to hcc_detail/hip_runtime.h 2016-02-01 14:29:50 +05:30
Maneesh Gupta 3b19fd578d Restrict using namespace hc::precise_math to device only 2016-02-01 14:26:50 +05:30