AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
2,442 Commits 3 Branches 0 Tags
e8dd9297ae516b41ed2e5a6fbff758aa48a6fe71
Commit Graph

84 Commits

Author SHA1 Message Date
Deven Desai 30f0c19cf1 Fixing a copy/paste error in my previous checkin 2018-04-27 12:59:51 +00:00
Deven Desai 76e808304c Checkin to add unit tests for fixes in my previous commit 
This change adds unit tests for

1. __hisinf
2. __hisnan
2. abs(long long)
3. lgamma(double)
 2018-04-25 17:53:21 +00:00
Maneesh Gupta 9e47fccc89 Apply .clangformat to all repo source files 
Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344
 2018-03-12 11:29:03 +05:30
Phaneendr-kumar Lanka f69762b300 [nvccWarnings] Fix -Wno-deprecated-declarations in hip_anyall and hip_ballot 2017-12-20 12:05:21 +05:30
Phaneendr-kumar Lanka 0ac125e3db [nvccWarnings] Fix warnings seen with dtests on nvcc path 2017-12-14 14:10:37 +05:30
Ben Sander 721d862089 Temporarily disable a couple tests pending some HCC work 2017-12-01 21:46:28 +00:00
Alex Voicu 4966518846 Revert "Revert adoption of CUDA indexing in general - this can only work with later versions of the compiler, just like module based dispatch, and thus must be guarded against usage in earlier (e.g. 1.6) versions." 
This reverts commit 4792475
 2017-11-29 21:49:10 +00:00
Alex Voicu 2557000b56 Revert "Revert adoption of CUDA indexing in general - this can only work with later versions of the compiler, just like module based dispatch, and thus must be guarded against usage in earlier (e.g. 1.6) versions." 
This reverts commit d2fd1f5
 2017-11-29 21:36:29 +00:00
Alex Voicu 0755f1fc26 Modify the set component of the memcpy test (unclear why there is a memset component to begin with). 2017-11-21 17:52:01 +00:00
Alex Voicu 30d90dab38 Re-sync with upstream. 2017-11-20 15:34:50 +00:00
Maneesh Gupta 4477d3d314 Merge pull request #266 from gargrahul/fix_half2_gfx900 
Fixed half2 issue on gfx900
 2017-11-20 07:28:41 +05:30
Alex Voicu 0cc921f103 This implements the trivial change needed to move back from the hip{Something}_{x, y, z} macros to the natural CUDA syntax of Something.{x, y, z}. This is contained in lines 384-404 in hip_runtime.h. All of the other changes have to do with changing unit tests to use this syntax. The macros are retained for backwards compatibility. 2017-11-19 01:54:12 +00:00
Rahul Garg cfc9f87465 Fixed test case for GFX900 2017-11-16 09:34:52 +05:30
Rahul Garg 20947f80fb Fixed half2 issue on gfx900 2017-11-15 18:52:59 +05:30
Alex Voicu 58a18eaf0c This corrects some interesting choices that were present in the HIP 
unit tests such as e.g. de-allocating memory allocated with new[] using
free. All of these were identified via cppcheck.
 2017-11-01 22:51:22 +00:00
Alex Voicu 28f87f7d2e This switches HIP from its currently convoluted macro + pfe based dispatch mechanism to a more natural one partially based on the existing module API. The basic idea is that HCC will always correctly emit __global__ functions: as empty-bodied stubs, on host, and as kernels, on device. It then becomes trivial to obtain the mangled name on host, at dispatch, from the function's address, and then to use the mangled name to retrieve the kernel. This should address all problems stemming from serialisation, dubious mismatches due to the manufactured functor, macro-isms et al. It also immediately enables support for generalised globals as a consequence of that being available in the module API. Finally, it will make debug much easier, since the actual names of the __global__ functions will automatically be used in traces etc. One detail is that due to how dispatch works now (hipLaunchKernel and hipLaunchKernelGGL are themselves variadic function templates which deduce the function type of the callee), in certain cases it may be necesssary to insert explicit casts to ensure that the variadic argument list selects a viable overload - this can be observed in some unit tests. Eventually we may be able to remove this limitation, but for now it does not appear terribly onerous. The code is not extremely HIPpie, nor is it fully optimised, but rather is intended as a starting point for the HIP team to make its own. 2017-11-01 15:09:59 +00:00
Maneesh Gupta 88ff8d9da4 [dtests] Fix hip_threadfence_system test's BUILD command 
Change-Id: I037c523369e4bddd7a386ffe117fe97a178900fd
 2017-08-11 14:49:17 +05:30
Siu Chi Chan 8c2cda8f9e add C++11 compilation flags and minor bug fixes 2017-08-10 17:58:36 +00:00
Maneesh Gupta 0d537f9966 [dtests] Disable hip_threadfence_system.tst on nvcc 
The test needs some more modifications before it is ready for building
on nvcc path. Disabling it for now so that the pull request can be
merged.

Change-Id: I19a894fcda7b2159f86a4e4e95a409c5655d3760
 2017-08-10 12:05:58 +05:30
Siu Chi Chan b6e22b5c93 implement __threadfence_system 2017-08-02 08:50:18 +00:00
Maneesh Gupta 1bd1c6bea8 Renable frexp(f) device math function 
Change-Id: I53c022b8ddf38cd17ddb42eba457b9020db66395
 2017-07-20 14:41:30 +05:30
Rahul Garg f2d6da3e34 Enable memcpytosymbol family functions test 2017-07-18 12:37:37 +05:30
Wen-Heng (Jack) Chung 0237aeab06 Remove explicit address spaces declarations 2017-07-07 12:46:29 -05:00
Maneesh Gupta a50f5ca0ac hipDeviceMemcpy test: make it functional on nvcc path 
Change-Id: Id10c79b48747ed701adbd0a233c53cd60cfa743b
 2017-06-07 15:24:44 +05:30
Maneesh Gupta 2145e94049 Disable rcbrtf, scalblnf, scalbnf in single precision device test 
Change-Id: I8a250a64a0cb05132d022a11d9766ced9cdf11a7
 2017-05-31 10:16:19 +05:30
Maneesh Gupta da19087ae2 Disable rcbrt, scalbln and scalbn double precision device test 
Change-Id: I46bd895701c46d3592b553090cafba99e41a2e2d
 2017-05-31 10:15:41 +05:30
Aditya Atluri 42739c37ef added __host__ for complex functions and corrected memset and memcpy test 
Change-Id: I9ffefb7a0025aa111a54d20d2766982df15532e7
 2017-04-06 09:29:44 -05:00
Sun, Peng 38c4331861 fix hipVectorTypesDevice direct test with GGL enabled 
Change-Id: I7a63b87348f08f094cd709e87397d9e0fc24e4c2
 2017-03-30 17:14:55 -05:00
Rahul Garg c33c58d687 Update hipTestDeviceSymbol sample 
Change-Id: If5ba99c60cd30c4491ca3a4856764224163d3ddf
 2017-03-24 10:39:11 +05:30
Ben Sander f4256cb57c Add USE_PROMOTE_FREE_HCC for smooth transition to new HCC caps. 
ADDRESS_SPACE_1 defines
 2017-03-17 11:04:39 -05:00
Ben Sander c9f64bbd2d Add simple device-side assert macro 
Currently swallows asserts but will compile.
 2017-03-17 10:25:56 -05:00
Maneesh Gupta 8536af7b71 Disable broken tests on hcc path 
Change-Id: Id6234da576566faa32d5fdf42dca6d6267596823
 2017-03-15 12:03:44 +05:30
Wen-Heng (Jack) Chung 17af837120 Revert "Changes to HIP to cope with Promote-free HCC" 
This reverts commit 0005dd5f66.

Change-Id: I20a9bab3883ad09913b320210344d37599cb8fcd
 2017-03-14 22:59:27 +08:00
Wen-Heng (Jack) Chung 0005dd5f66 Changes to HIP to cope with Promote-free HCC 
Squashed commit of the following:

commit c111b5bd10d7c2a5b0b1ad8b07f6e81185b47b39
Author: Wen-Heng (Jack) Chung <whchung@gmail.com>
Date:   Sat Mar 4 17:06:46 2017 +0800

    Use __device__ for all variables and functions to be used in kernel path

    Abolish __device and adopt [[hc]] in HIP implementation, so __device__ can be
    used on all HIP applications, no matter they are variables or functions.

    Change-Id: I20ca25857ce3bc3e42a5ebf65cafea2c8492f4c7

commit 30c0e4e4701bbf6bd9a7182e0320a71ff73d3a83
Author: Wen-Heng (Jack) Chung <whchung@gmail.com>
Date:   Thu Mar 2 12:14:11 2017 +0800

    XXX FIXME get around LDS spills caused in Promote-free HCC

    hipDynamicShared2 uses all 64KB of LDS for computation. But in Promote-free HCC
    there are cases where LDS spills would occur, which would make the test case to
    hang.

    In this workaround commit we reduce the size of dynamic LDS used to get around
    this known issue, and will revert this commit when LDS spills are resolved in
    HCC.

    Change-Id: If648b36200a4f9143951a8129192bcb7ed0bef5e

commit e803173be2d73e2f132a7ff7f61e7a20b4083d34
Author: Wen-Heng (Jack) Chung <whchung@gmail.com>
Date:   Wed Mar 1 21:41:41 2017 +0800

    Fix math functions which take pointer arguments

    Change-Id: I332c997e640edbc44824691e2a9434c6b3dadefa

commit de590c469e213c42090ff83dbd060f25bb1d6047
Author: Wen-Heng (Jack) Chung <whchung@gmail.com>
Date:   Wed Mar 1 18:38:54 2017 +0800

    Changes to cope with Promote-free HCC

    - abolish usage of address_space GNU attribute
    - use __device in file-scope global variables which would be accessed by GPU kernels
    - temporarily disable some math functions which take pointer arguments

    Change-Id: I730311dee848e20e763e35cd3980317fce0dce0d

Change-Id: I1f6b970b53b9401eeaaab08f04a7b9fed0fb8cf0
 2017-03-08 01:32:59 +08:00
Maneesh Gupta 578dc68b0a Disable some tests which are broken on nvcc path 
Change-Id: I6f8df7687ff1798dc17f6c9b8a7f4cd029ce45d8
 2017-02-27 13:19:06 +05:30
Aditya Atluri 7ac5017cb9 Added initial support for hipMemcpyFromSymbol. But not working! 
Change-Id: I48d8c7de4ec9f85c6c942be995fb488a3931f5d7
 2017-02-23 11:29:06 -06:00
Aditya Atluri 3d348b2d81 added runtime api hipMemcpyFromSymbolAsync 
Change-Id: Ibaf925faf0ba464dd0ed6c5ea74c224c2ce38889
 2017-02-22 19:16:35 -06:00
Aditya Atluri 2e3f3c7d26 Enable symbol tests 
Change-Id: I6bd036bf00c8051c8ff728ee60562c4ebd222160
 2017-02-22 13:42:03 -06:00
scchan a6ac4e7097 calls isfinite,isinf,isnan from the std namespace on the host 
Change-Id: Ica2370075b89713eecfd96102e2f4e0ab9961ce4
 2017-02-14 11:52:09 -05:00
Maneesh Gupta 6659eb8d16 Disable failing directed tests 
- hipTestDeviceSymbol
- hipTestConstant
- hipTestMallocKernel

Change-Id: Ibfe9fc0b8a59882f1de64b42e18777a7bd56ee97
 2017-02-09 14:48:22 +05:30
Aditya Atluri c6c90f6f5a added architecture specific macros 
1. Added __HIP_ARCH_GFXNUM__
2. Usage, -D__HIP_ARCH_GFX803__=1

Change-Id: I68b3a85d62cfab3a45d2b7a70cb3518ab2565236
 2017-02-08 19:45:32 -06:00
Aditya Atluri 400d0d4f78 fixed hipcc for new compiler flags 
Change-Id: I49ec059be20ff26b7482c84d91ab7a43826c6a8d
 2017-02-08 14:06:01 -06:00
Aditya Atluri 26b5e57cfd include arch specific ir on fallback path 
Change-Id: Ib04996aae2c21eb73ef2a9f6305915e0caccd704
 2017-02-08 12:19:06 -06:00
Aditya Atluri f009f3533d changed __global__ attribute 
1. Moved around tests and added them to HIT

Change-Id: I5d75280c42a5af852670ebabc7305ee56721ec7b
 2017-02-03 10:53:36 -06:00
Aditya Atluri 0d4e6ae60a fixed symbol memcpy issue 
Change-Id: I89d7401be51d194bcbf771020ba66e3d3b6a18f8
 2017-02-01 17:54:59 -06:00
Aditya Atluri a5d017e406 removed host math functions from math_functions.h 
Change-Id: I90d8784e2d6b58c6fade9f0fa12c0db3ee417d3e
 2017-01-27 17:38:43 -06:00
Aditya Atluri eee520def9 changed device code tests to work not to work as one 
Change-Id: I0eec1eab19dda3b703bc3a0d778a6bbb2802a412
 2017-01-27 09:20:14 -06:00
Aditya Atluri fd2e6ac2f0 added more test coverage for vector data types 
Change-Id: I9f57a8b597bd2ee4b265eadfd0859531497a6ada
 2017-01-20 13:52:02 -06:00
Aditya Atluri 02190736e3 fixed compilation issues for vector types and math functions 
1. Added math_functions.h to hip_runtime.h
2. Changed operator overloading classifier static to static inline
3. Added vector types test for gpu
4. Seperated __host__ and __device__ for math functions in headers

Change-Id: I499862fad5d7b10da686da9011d7ecefe523f8e2
 2017-01-20 09:49:11 -06:00
Aditya Atluri 41a46effef fixed compilation issues 
1. Fixed compilation issues for tests
2. Added missing intrinsics + math functions
3. Disabled some device functions as they are causing linking error with HCC

Change-Id: I79d52c4c7a539cc8ef40580247ad97ffcb975f09
 2017-01-18 11:53:47 -06:00