Commit Graph

354 Commits

Author SHA1 Message Date
Aditya Atluri 3577c19d66 v2: Fixed hipModule memory management
1. Changed test to assert for same hipFunction values
2. Added better memory management for hipModule

Change-Id: I10d7aef13c215a2211e262f3c79017f26a17d9a7


[ROCm/clr commit: 554fe191e1]
2017-02-10 13:32:13 -06:00
Aditya Atluri 5f8799a0a2 fixed shared memory usage implementation for hipDynamicShared2.cpp test
Change-Id: I34c72cb905f78de0f37e94174382e8be5c532028


[ROCm/clr commit: 307a21c946]
2017-02-09 18:17:42 -06:00
Aditya Atluri c5bcf5c096 added new dynamic shared mem test
Change-Id: Ic2a12fc9bc5b67e85f1e6c6008f10c7c66388377


[ROCm/clr commit: d6b5c02862]
2017-02-09 18:07:48 -06:00
Aditya Atluri 6c76483ac8 fixed hipFunction memory management
Change-Id: I7ebb323419bcd220ebd6466a8eb38e7bfdb1520a


[ROCm/clr commit: 6110fcdfb6]
2017-02-09 17:22:55 -06:00
Maneesh Gupta 90a03f8684 Disable failing directed tests
- hipTestDeviceSymbol
- hipTestConstant
- hipTestMallocKernel

Change-Id: Ibfe9fc0b8a59882f1de64b42e18777a7bd56ee97


[ROCm/clr commit: 6659eb8d16]
2017-02-09 14:48:22 +05:30
Aditya Atluri ba45e02007 added architecture specific macros
1. Added __HIP_ARCH_GFXNUM__
2. Usage, -D__HIP_ARCH_GFX803__=1

Change-Id: I68b3a85d62cfab3a45d2b7a70cb3518ab2565236


[ROCm/clr commit: c6c90f6f5a]
2017-02-08 19:45:32 -06:00
Aditya Atluri 76a58521f5 fixed hipcc for new compiler flags
Change-Id: I49ec059be20ff26b7482c84d91ab7a43826c6a8d


[ROCm/clr commit: 400d0d4f78]
2017-02-08 14:06:01 -06:00
Aditya Atluri 248b42d3e3 include arch specific ir on fallback path
Change-Id: Ib04996aae2c21eb73ef2a9f6305915e0caccd704


[ROCm/clr commit: 26b5e57cfd]
2017-02-08 12:19:06 -06:00
Aditya Atluri d191eac0e2 Fixed HIP for C runtime
1. fixed constant memory test
2. added runtime gcc/g++ test
3. added fix for runtime with gcc/g++

Change-Id: Ie14dad6633411b188bdeea044e140b6d5beabe16


[ROCm/clr commit: be6d01ebc3]
2017-02-07 13:15:36 -06:00
Maneesh Gupta c072d54c50 Disable broken test: hipC
Change-Id: I061aa125bbdc9f14bc870266ab0735593c861903


[ROCm/clr commit: 0bef861f1a]
2017-02-06 13:00:50 +05:30
Aditya Atluri 3937b398e9 changed __global__ attribute
1. Moved around tests and added them to HIT

Change-Id: I5d75280c42a5af852670ebabc7305ee56721ec7b


[ROCm/clr commit: f009f3533d]
2017-02-03 10:53:36 -06:00
Aditya Atluri f36992067b fixed symbol memcpy issue
Change-Id: I89d7401be51d194bcbf771020ba66e3d3b6a18f8


[ROCm/clr commit: 0d4e6ae60a]
2017-02-01 17:54:59 -06:00
Aditya Atluri 4085e01263 removed host math functions from math_functions.h
Change-Id: I90d8784e2d6b58c6fade9f0fa12c0db3ee417d3e


[ROCm/clr commit: a5d017e406]
2017-01-27 17:38:43 -06:00
Aditya Atluri 324409b2bb changed device code tests to work not to work as one
Change-Id: I0eec1eab19dda3b703bc3a0d778a6bbb2802a412


[ROCm/clr commit: eee520def9]
2017-01-27 09:20:14 -06:00
Aditya Atluri f6d09573aa added more test coverage for vector data types
Change-Id: I9f57a8b597bd2ee4b265eadfd0859531497a6ada


[ROCm/clr commit: fd2e6ac2f0]
2017-01-20 13:52:02 -06:00
Aditya Atluri b6f4fedaaf 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


[ROCm/clr commit: 02190736e3]
2017-01-20 09:49:11 -06:00
Aditya Atluri 69903887a4 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


[ROCm/clr commit: 41a46effef]
2017-01-18 11:53:47 -06:00
Aditya Atluri 42c627fbe8 Moved device code to mimic cuda header behavior
1. All fp32, fp64 math device/host functions should be in math_functions.h/.cpp
2. All fp32, fp64 fast math intrinsics for device/host functions should be in device_functions.h/.cpp
3. All the device code implementations should be in device_util.h/.cpp
4. Hence, made changes appropriately by moving code and creating new header files
5. Added math_functions.cpp/.h
6. Changed #ifndef signature to make sure no conflicts between headers with same names in hip/hip_runtime.h and hip/hcc_detail/hip_runtime.h
7. Changed tests to fit the code changes, making them to include appropriate headers
8. Added math_functions.cpp to CMakeLists.txt
9. Some of the tests are still broken, mostly host math functions will fix them in next commit
10. TODO: FIX compilation issues for host math functions

Change-Id: I7a17637d7e294a7d224ffba932c1a08668febd26


[ROCm/clr commit: d23b6b8694]
2017-01-17 14:57:51 -06:00
Aditya Atluri 968f2e9489 enabled integer intrinsics tests
Change-Id: I5d28d556f228240eda2fc0098121ed3b29b041e7


[ROCm/clr commit: 3f9a9d9318]
2017-01-17 09:59:08 -06:00
Aditya Atluri b864271e2c fixed broken tests and device code for integer intrinsics
1. Fixed build issues with new Integer intrinsics
2. Changed tests to work exactly as CUDA code
3. Still some integer intrinsics need to be supported

Change-Id: Ie6f4171259cf4da517436895d4f6f01e01f59b11


[ROCm/clr commit: f0ea51c786]
2017-01-17 09:00:09 -06:00
Aditya Atluri d9845446ef Added type conversion intrinsics
1. Added all type conversion intrinsics
2. NO TESTS have been added. (Will add in next commit)
3. Sanatized code in hip_runtime.h
4. Added passed() to hipTestHalf to make it pass on HIT

Change-Id: I0987963c802fc7ff4d7e07d7b88d86da35da53c9


[ROCm/clr commit: d496576b55]
2017-01-16 12:10:05 -06:00
Aditya Atluri d71fc0e60a added half2 log, log10, exp, exp10 math functions
1. Enabled tests for log, log10, exp, exp10 half2
2. h2rint is still disabled.

Change-Id: I01f6002f6992259919893c524c526db5ee09473a


[ROCm/clr commit: 5c5f5c1ad1]
2017-01-13 13:26:10 -06:00
Aditya Atluri 4756cf7c16 added half2 math operations
1. They use SDWA + LLVM IR
2. Added these functions to test
3. Need to do exp, exp10, log, log10, rint

Change-Id: I06176acc6cb8bb054495310531777406a41b54e4


[ROCm/clr commit: eff68c989a]
2017-01-13 12:27:11 -06:00
Aditya Atluri b2973b97e2 added math functions for half
1. Added math functions for half precision
2. HRCP is not available due to device code linking errors, will be enabled once it is fixed
3. Added math functions to half test file

Change-Id: Ie317ce70ef518a4fc3f27142143d01e0327f5df3


[ROCm/clr commit: fe38e9652b]
2017-01-13 12:05:29 -06:00
Aditya Atluri d3fe56550e added half2 cmp and conv, data movement device functions
1. Added half2 comparision functions
2. Added conversion and data movement half apis

Change-Id: Ia33c0e957d9deb1f2b7a8fde8e22168f4d41b88b


[ROCm/clr commit: 646f566bbf]
2017-01-13 10:56:07 -06:00
Aditya Atluri a52591d117 added comparision device functions for fp16
1. Added comparision device functions
2. Added test to check correct isa getting generated

Change-Id: I16732f5a1438bdce145f7bfcecd28198e3cc4b79


[ROCm/clr commit: 89998d436f]
2017-01-12 14:52:14 -06:00
Aditya Atluri d1f7c4e048 added packed math fp16 native device functions
1. Added SDWA implementation inside IR file
2. Added device functions to header + used them in test

Change-Id: Ib4e059a58eee201cc82438689e3e9bc5f9d26653


[ROCm/clr commit: eeef055469]
2017-01-12 14:10:51 -06:00
Aditya Atluri 5317ffa9ec Started adding native half math library support
1. Removed HIP_EXPERIMENTAL env variable so that device code will be accessed from LLVM IR
2. Removed soft support from headers and moved to hip_fp16.cpp
3. Added LLVM IR + inline asm to hip_ir.ll
4. Added test for fp16
5. Added barriers for hcc 3.5 and hcc 4.0 for half support
a. Which means, hcc 4.0 can parse __fp16 but hcc 3.5 cant
b. HCC 4.0 code is implemented now, hcc 3.5 will be added later

Change-Id: Ic37859b2688ebb02e168bab643d1882bf4727952


[ROCm/clr commit: c286bf6f8a]
2017-01-12 11:30:20 -06:00
Aditya Atluri e4189cab53 added test for vector data types
Change-Id: I0b6624886e474601cb1ef003c5f10adf399a21c9


[ROCm/clr commit: 1d8700096c]
2017-01-11 18:02:30 -06:00
Aditya Atluri fe2d13c861 Added proper device data types
Change-Id: I42029635ff68c3c13a764a3eda6447e6c77878c6


[ROCm/clr commit: 4e57822d95]
2017-01-11 15:06:25 -06:00
Ben Sander 9af6545379 Fix delete[]
[ROCm/clr commit: ff77106399]
2017-01-09 21:03:11 -06:00
Rahul Garg 8acb6a623d Mapped hipDevice_t to int
Change-Id: I6cfa56c42b7cd04aa0e0bce510c0d72d34ea211a


[ROCm/clr commit: bddaa0e81c]
2016-12-17 16:53:03 +05:30
Rahul Garg 74803e3ddf Simple test case for hipDeviceGetPCIBusId
Change-Id: I0fb6d1bef4739a5b6280928f7b349d95c1656431


[ROCm/clr commit: afc81c955e]
2016-12-06 17:31:54 +05:30
Maneesh Gupta 97d8017e80 Fix some broken directed tests
Change-Id: I40f2661a74951f8d91824df8fd9ef0cc2312d183


[ROCm/clr commit: 9cab8d4260]
2016-12-01 12:47:37 +05:30
Aditya Atluri 62e0623d18 added support for rcp for float and double
Change-Id: Ibeba3a9f64494fc0a176bcb4a854fb2f56567b55


[ROCm/clr commit: de89b25d52]
2016-11-23 20:01:18 -06:00
Aditya Atluri 11018fa0f6 added fma for double and float
1. Added fma intrinsic support for double and float
2. Added test for fma

Change-Id: I909fdbec34a3d12c03ba6eff3a39376a7128ee43


[ROCm/clr commit: cc1f8a1011]
2016-11-23 18:22:05 -06:00
Aditya Atluri a2ecbc2d6e Added fast math flag
1. Use -DHIP_FAST_MATH to make precise math functions compiled to fast math
2. Added double fast math functions for sqrt
3. Changed hipcc to parse -use_fast_math (not working)
4. Added passed tag to hipFloatMath test

Change-Id: I72884b2436b4efe61e9a9297346c1358fee38a2d


[ROCm/clr commit: c2f6ecf264]
2016-11-23 11:19:15 -06:00
Aditya Atluri 06969472b2 added fast math intrinsics to HIP
1. Added fast math intrinsics for single precision data types
2. Added test to check the intrinsics
3. Added HIP_PRECISE_MATH macro to enable precise math on fast math

Change-Id: Iadacbb6182c31252c5e3252854372d1b80dfd27b


[ROCm/clr commit: d9a3527769]
2016-11-22 15:26:00 -06:00
Aditya Atluri b8b1459a64 added fast math APIs
1. Added fast math apis for sin, cos, tan, sincos
2. Added test for trig math functions
3. Added logarithm fast math
4. Changed how hipGetDevice, hipDeviceGetCacheConfig emit errors

Change-Id: Ie6ab594ddd5853cbe85e39a2f6d3479a807fa323


[ROCm/clr commit: 1a85762f53]
2016-11-22 10:20:09 -06:00
Aditya Atluri a7bb6239ab fixed texture header on nvcc
Change-Id: Ibe19f94be5edf972b6b51dea263e1088b6c60c1d


[ROCm/clr commit: 2ded0ce302]
2016-11-21 13:53:28 -06:00
Aditya Atluri 59c423b36a fixed compilation bugs
1. Texture functions are now compiling fine
2. Fixed hipFuncCache to hipFuncCache_t

Change-Id: I8f815887e4de43ee115bbaff249905b236541c39


[ROCm/clr commit: 2611de2477]
2016-11-21 08:56:30 -06:00
Aditya Atluri df27127960 Fixed hipDeviceGetCacheConfig on nvcc path
1. Changed test macro to emit line numbers
2. Added getcacheconfig api test for nvcc path
3. Fixed hipFuncCache_t data type

TODO: With this commit, right now there are 2 func cache datatypes
a. hipFuncCache_t for runtime API
b. hipFuncCache for driver API

Map these to a single data type

Change-Id: Ia47c9f5d7c2633638051bf17b1103048a1ede973


[ROCm/clr commit: b3c16ea7b5]
2016-11-20 12:18:08 -06:00
Aditya Atluri 979e54d54e added new test for getting attribute
1. Added copyright to all new tests
2. Added test for hipDeviceGetAttribute

Change-Id: I7a070c5b8316ef6575b3f4c49bda2769aea2a7c4


[ROCm/clr commit: e0aba8647f]
2016-11-20 11:53:16 -06:00
Aditya Atluri 21470101c7 added tests to check nvcc runtime api output
Change-Id: Ifdd39b5d0a6a58d20a8e9745e59dd82d50a90e2f


[ROCm/clr commit: 6692ee09d7]
2016-11-19 21:36:28 -06:00
Maneesh Gupta 333e994a57 Fix broken tests due to d3559bf
Change-Id: I847c80f8462e1c955bdef957e6de2841a3a6ab29


[ROCm/clr commit: 1a67766dc5]
2016-11-18 12:20:47 +05:30
pensun dd1061b874 Add direct test case for threadfence_system workaround
Change-Id: I5b21b590e957c901044741ac94e816cd8b1426f9


[ROCm/clr commit: 992f94b3a1]
2016-11-11 15:09:43 -06:00
Ben Sander 5e354dcd77 Add option to deny peer access.
Also fix test.

Change-Id: I1b247f6c4271442b008e560669bca4daf8eb94c7


[ROCm/clr commit: d666fbaafe]
2016-11-10 23:12:48 -06:00
Ben Sander 0c66772f37 Doc change only - add comments to test.
Change-Id: Ie42087cf3c78e49337b18bb71f3f0e1e7950ee1b


[ROCm/clr commit: 85e65b55ff]
2016-11-10 23:12:48 -06:00
Ben Sander 3d4a76d560 Fix tests to read warpSize from device props.
Change-Id: I9583577793afad49f9eb1ee9069bd4c6963a6023


[ROCm/clr commit: a13ec441bf]
2016-11-06 04:26:28 -06:00
Ben Sander 03fcf556e9 Improve Peer support and testing.
Change-Id: Icadc65988aaf145a265587ab0357c5bf4d26f3eb


[ROCm/clr commit: f3d38c2615]
2016-11-06 03:22:36 -06:00