maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
373,802 Commits 23 Branches 0 Tags 2.2 GiB
Commit Graph

610 Commits

Author SHA1 Message Date
Jan Vesely 15c388cd79 exp10: Port from amd builtins 
Passes CTS on carrizo and turks.
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed and Tested (on RX 580) by: Aaron Watry <awatry@gmail.com>

llvm-svn: 330197
 2018-04-17 18:08:08 +00:00
Jan Vesely 4be0339023 hypot: Port from amd builtins 
v2: Fix whitespace errors

Use only subnormal path.
Passes CTS on carrizo and turks.
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry <awatry@gmail.com>

llvm-svn: 329647
 2018-04-10 00:11:58 +00:00
Jan Vesely 4c1112612c select: simplify implementation and fix fp16 
Fix half precision implementation
Vector ?: operator should behave exactly as select
Passes CTS on carrizo

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
llvm-svn: 329462
 2018-04-06 22:00:00 +00:00
Jan Vesely 93af966747 fmod: Port from amd_builtins 
Uses only denormal path for fp32.
Passes CTS on carrizo and turks.

v2: whitespace fix

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry <awatry@gmail.com>
llvm-svn: 329433
 2018-04-06 17:43:08 +00:00
Jan Vesely 92357a2336 r600: Update datalayout after LLVM r328656 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 329291
 2018-04-05 14:47:57 +00:00
Jan Vesely fd11db19c2 amdgcn: Update datalayout after LLVM r328656 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 329290
 2018-04-05 14:47:44 +00:00
Jan Vesely 5b10494fa8 remainder: Port from amd builtins 
Mostly ported from amd_builtins, uses only denormal path for fp32.
Passes CTS on carrizo and turks

Reviewer: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 327818
 2018-03-19 01:01:10 +00:00
Jan Vesely b672f7a251 nan: Implement 
Passes CTS on carrizo and turks

Reviewer: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 327324
 2018-03-12 19:46:52 +00:00
Jan Vesely 08c96acb27 travis: Add build using llvm-6 
Acked-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 327323
 2018-03-12 19:46:48 +00:00
Jan Vesely f96b1b88f8 amdgcn/fmax: fcanonicalize operands 
v_max instruction needs canonicalized operands.
Passes CTS on carrizo

Reviewer: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 327076
 2018-03-08 23:01:01 +00:00
Jan Vesely e724e346ab amdgcn/fmin: fcanonicalize operands 
v_min instruction needs canonicalized operands.
Passes CTS on carrizo

Reviewer: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 327075
 2018-03-08 23:00:58 +00:00
Jan Vesely 04a46bf0a2 amdgcn,popcount: Workaround broken llvm.ctpop intrinsic on some GCN ASICs 
This is only really needed for VI+ ASICs. However, llvm would cast the value to
i32 for older asics anyway. The proper fix is in LLVM-7 (r326535).
Fixes CTS popcount on carrizo.

Reviewer: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 327044
 2018-03-08 18:58:07 +00:00
Jan Vesely 0883c4d365 integer/gentype: Add __CLC_VECSIZE macro 
Reviewer: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 327043
 2018-03-08 18:58:05 +00:00
Jan Vesely 17e8679493 popcount: Provide function implementation rather than intrinsic redirect 
amdgcn will need to override this

Reviewer: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 327042
 2018-03-08 18:58:00 +00:00
Jan Vesely c15a48dd9c lgamma_r: Move code from .inc to .cl file 
Reviewed-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 326821
 2018-03-06 17:48:47 +00:00
Jan Vesely 2dcb382efc frexp: Reuse types provided by gentype.inc 
v2: Use select instead of bitselect to consolidate scalar and vector
versions

Passes CTS on Carrizo

Reviewed-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 326820
 2018-03-06 17:48:45 +00:00
Jan Vesely ae156b66f8 select: Add vector implementation 
Passes CTS on Carrizo

Reviewed-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 326819
 2018-03-06 17:48:43 +00:00
Jan Vesely 44f21978a2 minmag: Condition variable needs to be the same bitwidth as operands 
No changes wrt CTS

Reviewed-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 326818
 2018-03-06 17:48:40 +00:00
Jan Vesely 4e72300929 maxmag: Condition variable needs to be the same bitwidth as operands 
No changes wrt CTS

Reviewed-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 326817
 2018-03-06 17:48:38 +00:00
Jan Vesely dbaf6d0f7c Move cl_khr_fp64 exntension enablement to gentype include lists 
This will make adding cl_khr_fp16 support easier

Reviewed-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 326816
 2018-03-06 17:48:35 +00:00
Jan Vesely 86db4302e0 utils: Adapt to llvm r325155 
r325155 ("Pass a reference to a module to the bitcode writer.")
changed bit writer interface from pointer to reference

Reviewer: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325867
 2018-02-23 07:37:03 +00:00
Jan Vesely 1ad6a94676 amdgcn: Fix build after GDS/const AS swap in r325030 
Acked-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325866
 2018-02-23 07:37:01 +00:00
Jan Vesely eda1872d04 amdgcn: Fix datalayout after addition of 32bit const AS in r324747 
Acked-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325865
 2018-02-23 07:36:54 +00:00
Jan Vesely 83cd840010 r600: Fix datalayout after clang r324101 
r324101 switched around AS numbering

Acked-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325864
 2018-02-23 07:36:51 +00:00
Jan Vesely 911666f3fa amdgcn: Fix datalayout after clang r324101 
r324101 switched around AS numbering

Acked-by: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325863
 2018-02-23 07:36:39 +00:00
Jan Vesely b424954682 amdgpu/half_recip: Switch implementation to native_recip 
Reviewer: Tom Stellard <tstellar@redhat.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325061
 2018-02-13 22:09:46 +00:00
Jan Vesely ed28c4458a amdgpu/half_log2: Switch implementation to native_log2 
Reviewer: Tom Stellard <tstellar@redhat.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325060
 2018-02-13 22:09:44 +00:00
Jan Vesely 86cbf56a4b amdgpu/half_log10: Switch implementation to native_log10 
Reviewer: Tom Stellard <tstellar@redhat.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325059
 2018-02-13 22:09:42 +00:00
Jan Vesely 65fd65efbf amdgpu/half_log: Switch implementation to native_log 
Reviewer: Tom Stellard <tstellar@redhat.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325058
 2018-02-13 22:09:41 +00:00
Jan Vesely 2d3b6dfdca amdgpu/half_exp2: Switch implementation to native_exp2 
Reviewer: Tom Stellard <tstellar@redhat.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325057
 2018-02-13 22:09:38 +00:00
Jan Vesely 021264c75a amdgpu/half_exp10: Switch implementation to native_exp10 
Reviewer: Tom Stellard <tstellar@redhat.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325056
 2018-02-13 22:09:37 +00:00
Jan Vesely 4879dd7471 amdgpu/half_exp: Switch implementation to native_exp 
Reviewer: Tom Stellard <tstellar@redhat.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325055
 2018-02-13 22:09:35 +00:00
Jan Vesely bca92445ba amdgpu/half_sqrt: Switch implementation to native_sqrt 
Reviewer: Tom Stellard <tstellar@redhat.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325054
 2018-02-13 22:09:33 +00:00
Jan Vesely aad28681c2 amdgpu/half_rsqrt: Switch implementation to native_rsqrt 
Reviewer: Tom Stellard <tstellar@redhat.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 325053
 2018-02-13 22:09:31 +00:00
Jan Vesely 1c570566c3 Add vstore_half_rte implementation 
Passes CTS on carrizo

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 324376
 2018-02-06 18:44:50 +00:00
Jan Vesely f2d876ae83 Add vstore_half_rtp implementation 
Passes CTS on carrizo

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 324375
 2018-02-06 18:44:47 +00:00
Jan Vesely 2655312c69 Add vstore_half_rtn implementation 
Passes CTS on carrizo

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 324374
 2018-02-06 18:44:45 +00:00
Jan Vesely d526a2b6e8 Add vstore_half_rtz implementation 
Passes CTS on carrizo

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 324373
 2018-02-06 18:44:43 +00:00
Jan Vesely 4475aca172 vstore_half: Consolidate declarations 
Add support for rounding suffix

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 324372
 2018-02-06 18:44:41 +00:00
Jan Vesely 187ec00556 vstore_half: Add support for custom rounding functions 
Add another layer of indirection
This will be used for specific rounding modes

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 324371
 2018-02-06 18:44:39 +00:00
Jan Vesely 87036d2701 vstore_half: Make sure the helper function is always inline 
Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 324370
 2018-02-06 18:44:35 +00:00
Jan Vesely 3b8b4eb64d half_powr: Implement using powr 
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 323942
 2018-02-01 03:00:35 +00:00
Jan Vesely a75677c2b7 math.h: Use logical operations instead of bit operations for readability 
Trivial.

Reported-by: Roman Lebedev <lebedev.ri@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 323920
 2018-01-31 21:53:42 +00:00
Jan Vesely 0ecb5e511e math.h: Set HAVE_HW_FMA32 based on compiler provided macro 
Fixes sin/cos piglits on non-FMA capable asics.
Bugzilla: https://bugs.llvm.org/show_bug.cgi?id=35983

Reviewer: Tom Stellard
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 323677
 2018-01-29 19:05:08 +00:00
Jan Vesely 7013857f95 tanpi: Port from amd_builtins 
Passes piglit on turks and carrizo.
Passes CTS on carrizo.

Acked-By: Aaron Watry <awatry@gmail.com>
Tested-By: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322980
 2018-01-19 18:57:22 +00:00
Jan Vesely 03937bdec3 tan: Port from amd_builtins 
v2: fixup constant precision
Passes piglit on turks and carrizo.
Passes CTS on carrizo
Fixes half_tan to pass CTS on carrizo

Acked-By: Aaron Watry <awatry@gmail.com>
Tested-By: Aaron Watry <awatry@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322979
 2018-01-19 18:57:19 +00:00
Jan Vesely 44e0522c09 half_divide: Implement using x/y 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322899
 2018-01-18 21:12:06 +00:00
Jan Vesely 2813b4f8d9 half_tan: Implement using tan 
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322898
 2018-01-18 21:12:04 +00:00
Jan Vesely bf38fae8de half_sin: Implement using sin 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322897
 2018-01-18 21:12:01 +00:00
Jan Vesely 398108b91e half_recip: Implement using 1/x 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322896
 2018-01-18 21:11:58 +00:00
Jan Vesely a1aba44ffa half_log2: Implement using log2 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322895
 2018-01-18 21:11:56 +00:00
Jan Vesely b3b72af4b9 half_log10: Implement using log10 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322894
 2018-01-18 21:11:53 +00:00
Jan Vesely 6852023802 half_log: Implement using log 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322893
 2018-01-18 21:11:50 +00:00
Jan Vesely aa4c3899b5 half_exp10: Implement using exp10 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322892
 2018-01-18 21:11:48 +00:00
Jan Vesely 3c0e19b61a half_exp2: Implement using exp2 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322891
 2018-01-18 21:11:45 +00:00
Jan Vesely caa9000b1c half_exp: Implement using exp 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322890
 2018-01-18 21:11:43 +00:00
Jan Vesely b5d556061d half_cos: Implement using cos 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322889
 2018-01-18 21:11:40 +00:00
Jan Vesely e53ae3b596 half_sqrt: Cleanup implementation 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322888
 2018-01-18 21:11:38 +00:00
Jan Vesely a95db14461 half_rsqrt: Cleanup implementation 
Passes CTS on carrizo
v2: Use full precision implementation

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322887
 2018-01-18 21:11:35 +00:00
Jan Vesely fe8e00bc3c rootn: Port from amd_builtins 
Passes piglit on turks and carrizo
fp64 passes ctx on carrizo

v2: fix formatting
    check fp32 denormal support at runtime

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322763
 2018-01-17 21:22:14 +00:00
Jan Vesely c45ec604f5 powr: Port from amd_builtins 
Passes piglit on turks and carrizo
fp64 passes cts on carrizo

v2: fix formatting
    check fp32 denormal support at runtime

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322762
 2018-01-17 21:22:06 +00:00
Jan Vesely 5efc8fe321 pown: Port from amd_builtins 
Passes piglit on turks and carrizo
fp64 passes CTS on carrizo

v2: fix formatting
    check fp32 denormal support at runtime

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322761
 2018-01-17 21:22:03 +00:00
Jan Vesely cc5c65b2c2 pow: Port from amd_builtins 
Passes piglit on turks and carrizo
fp64 passes CTS on carrizo

v2: fix formatting
    check fp32 denormal support at runtime

Reviewer: Jeroen Ketema <j.ketema@xs4all.nl>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 322760
 2018-01-17 21:21:35 +00:00
Vedran Miletic 79b7f4c125 configure.py: Add gfx900 (Vega, Raven) 
Sort amdgcn-- and amdgcn--amdhsa in a consistent way.

llvm-svn: 319017
 2017-11-27 11:14:06 +00:00
Jan Vesely fe7c045753 math: Implement minmag 
Reviewer: Aaron Watry
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 318265
 2017-11-15 04:10:39 +00:00
Jan Vesely 7ba243cc3d math: Implement maxmag 
Reviewer: Aaron Watry
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 318264
 2017-11-15 04:10:37 +00:00
Jan Vesely 383fbd050c native_powr: Switch implementation to native_exp2 and native_log2 
v2: don't use assume
    check only for x<0, the other conditions are handled transparently
v3: don't check inputs at all, nan propagation works as expected

Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 318204
 2017-11-14 21:55:41 +00:00
Jan Vesely f38b40daf7 native_divide: provide function implementation instead of macro 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 318067
 2017-11-13 18:28:56 +00:00
Jan Vesely 1b9825f982 native_recip: provide function implementation instead of macro 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 318066
 2017-11-13 18:28:53 +00:00
Jan Vesely a6758c94ef native_rsqrt: Switch implementation to 1 / native_sqrt 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 318065
 2017-11-13 18:28:51 +00:00
Jan Vesely 541a3f0758 native_tan: Switch implementation to use native_sin/native_cos 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 318064
 2017-11-13 18:28:48 +00:00
Jan Vesely 79b7566210 math: Use precomputed constant for log2(10.0) 
exp10 CTS fails with or without this change

Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 318063
 2017-11-13 18:28:45 +00:00
Jan Vesely 6b4a625438 native_exp10: Switch implementation to llvm intrinsic 
v2: Use native_log2 instead of wrong constant

Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 317941
 2017-11-10 22:16:41 +00:00
Jan Vesely 4301e6d0c9 native_sqrt: Switch implementation to llvm intrinsic 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 317940
 2017-11-10 22:16:39 +00:00
Jan Vesely 1f34c851e0 native_sin: Switch implementation to llvm intrinsic 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 317939
 2017-11-10 22:16:36 +00:00
Jan Vesely 0750b7df51 native_cos: Switch implementation to llvm intrinsic 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 317938
 2017-11-10 22:16:33 +00:00
Jan Vesely edbde58de0 native_exp2: Switch implementation to llvm intrinsic 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 317937
 2017-11-10 22:16:31 +00:00
Jan Vesely 504f85c551 native_exp: Switch implementation to llvm intrinsic 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 317936
 2017-11-10 22:16:28 +00:00
Jan Vesely 8dc6e98d47 amdgpu: Add workaround for unimplemented llvm.exp intrinsic 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 317935
 2017-11-10 22:16:25 +00:00
Jan Vesely adc1eaedf8 native_log10: Switch to generic native intrinsic inc file 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 317934
 2017-11-10 22:16:22 +00:00
Jan Vesely 086e796053 native_log: Switch to generic native intrinsic inc file 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 317933
 2017-11-10 22:16:20 +00:00
Jan Vesely f58dee9f3a native_log2: Switch to generic native intrinsic inc file 
v2: Add __CLC_XCONCAT instead of function name redirection
    Use __CLC_XCONCAT for intrinsic functions as well

Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 317932
 2017-11-10 22:16:15 +00:00
Jan Vesely 39ef293533 tgamma: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317260
 2017-11-02 19:49:00 +00:00
Jan Vesely 91d7b92d8a tanh: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317259
 2017-11-02 19:48:58 +00:00
Jan Vesely e4d5d10076 tan: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317258
 2017-11-02 19:48:57 +00:00
Jan Vesely b0fab2696a sqrt: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317257
 2017-11-02 19:48:55 +00:00
Jan Vesely e3802356e2 sinpi: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317256
 2017-11-02 19:48:53 +00:00
Jan Vesely 4708b10878 sinh: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317255
 2017-11-02 19:48:51 +00:00
Jan Vesely a3febe3fa9 sin: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317254
 2017-11-02 19:48:50 +00:00
Jan Vesely 25671b40d7 native_log: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317253
 2017-11-02 19:48:48 +00:00
Jan Vesely fd13434d83 native_log2: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317252
 2017-11-02 19:48:46 +00:00
Jan Vesely d6ad07687d native_log10: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317251
 2017-11-02 19:48:44 +00:00
Jan Vesely 139185dfc7 log: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317250
 2017-11-02 19:48:43 +00:00
Jan Vesely 27dffff6e8 logb: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317249
 2017-11-02 19:48:41 +00:00
Jan Vesely 7fc23fbdcb log2: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317248
 2017-11-02 19:48:39 +00:00
Jan Vesely a9132ce347 log1p: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317247
 2017-11-02 19:48:37 +00:00
Jan Vesely 4e062cb74e lgamma: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317246
 2017-11-02 19:48:35 +00:00
Jan Vesely 4cb612e140 exp2: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317245
 2017-11-02 19:48:33 +00:00
Jan Vesely e99ba9a23d cospi: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317244
 2017-11-02 19:48:31 +00:00
Jan Vesely c708278f13 cosh: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317243
 2017-11-02 19:48:30 +00:00
Jan Vesely f76371d948 cos: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317242
 2017-11-02 19:48:27 +00:00
Jan Vesely 50a3cccdbe cbrt: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317241
 2017-11-02 19:48:25 +00:00
Jan Vesely a4df39bcad atanpi: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317240
 2017-11-02 19:48:23 +00:00
Jan Vesely 1bd2ac257a atanh: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317239
 2017-11-02 19:48:22 +00:00
Jan Vesely 0942b5e1bf atan: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317238
 2017-11-02 19:48:20 +00:00
Jan Vesely d3d5e322e3 asinpi: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317237
 2017-11-02 19:48:18 +00:00
Jan Vesely 48bda32986 asinh: Use unary_dec instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317236
 2017-11-02 19:48:16 +00:00
Jan Vesely ba4b98c691 asin: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317235
 2017-11-02 19:48:15 +00:00
Jan Vesely 61171847b7 acospi: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317234
 2017-11-02 19:48:13 +00:00
Jan Vesely 720783d9f5 acosh: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317233
 2017-11-02 19:48:11 +00:00
Jan Vesely caca914218 acos: Use unary_decl instead of custom inc file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-By: Aaron Watry <awatry@gmail.com>
llvm-svn: 317232
 2017-11-02 19:48:06 +00:00
Jan Vesely 47e093da9b math: Implement native_log10 
Use llvm instrinsic by default
Provide amdgpu workaround

v2: drop old amd copyrights

Reviewer: Aaron Watry
Reviewed-by: Vedran Miletić <vedran@miletic.net>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 316588
 2017-10-25 16:49:22 +00:00
Jan Vesely 9fedbb9d8e amdgpu/math: Don't use llvm instrinsic for native_log 
AMDGPU targets don't have insturction for it,
so it'll be expanded to C * log2 anyway.

v2: use native_log2 instead of the more precise sw implementation
v3: move to amdgpu
v4: drop old AMD copyright

Reviewer: Aaron Watry
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 316587
 2017-10-25 16:49:17 +00:00
Jan Vesely 7ab2d0bdcd shared: Implement aligned vector stores (vstorea_half) 
Float version passes newly posted piglit tests on turks, float and double pass on carrizo.
v2: scalar vstorea_half
v3: fix typo

Reviewer: Aaron Watry
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 316291
 2017-10-22 14:21:59 +00:00
Jan Vesely 12061c7125 shared: Implement aligned vector loads (vloada_half) 
Passes newly posted piglits on turks and carrizo
v2: add scalar vloada_half
v3: fix typo

Reviewer: Aaron Watry
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 316290
 2017-10-22 14:21:56 +00:00
Jan Vesely c420b61b26 amdgcn: Add missing datalayout info to .ll files 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 316239
 2017-10-20 21:10:18 +00:00
Jan Vesely 66b32ad9ad r600: Add missing datalayout to .ll files 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 316238
 2017-10-20 21:00:31 +00:00
Jan Vesely 577c52b9c7 travis: enable checks of nvptx libraries 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315343
 2017-10-10 18:10:25 +00:00
Jan Vesely 2601429bac travis: Enable external function call checks on llvm-{4,5} 
Reviewer: Aaron Watry
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315342
 2017-10-10 18:10:24 +00:00
Jan Vesely 3d349ea98e Make image builtins r600/llvm-3.9 only 
The implementation uses r600 sepcific intrinsics
LLVM-4 switched to _ro_t and _rw_t image types
Portions of the code can be moved back as more targets/llvm versions add image support

Reviewer: Aaron Watry
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315341
 2017-10-10 18:10:21 +00:00
Jeroen Ketema 1364d268a4 Implement mem_fence on ptx 
PTX does not differentiate between read and write fences. Hence, these a
lowered to a mem_fence call. The mem_fence function compiles to the
“member.cta” instruction, which commits all outstanding reads and writes
of a thread such that these become visible to all other threads in the same
CTA (i.e., work-group). The instruction does not differentiate between
global and local memory. Hence, the flags parameter is ignored, except
for deciding whether a “member.cta” instruction should be issued at all.

Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315235
 2017-10-09 19:43:04 +00:00
Jeroen Ketema 4f5a3d5d6f Make ptx barrier work irrespective of the cl_mem_fence_flags 
This generates a "bar.sync 0” instruction, which not only causes the
threads to wait, but does acts as a memory fence, as required by
OpenCL. The fence does not differentiate between local and global
memory. Unfortunately, there is no similar instruction which does
not include a memory fence. Hence, we cannot optimize the case
where neither CLK_LOCAL_MEM_FENCE nor CLK_GLOBAL_MEM_FENCE is
passed.

llvm-svn: 315228
 2017-10-09 18:36:48 +00:00
Jan Vesely 3c51ae5bd9 travis: Make sure we report failure even if only earlier checked files fail 
for loop would only report status of the last command
v2: return '1'
    call test instead of '['

Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315193
 2017-10-08 20:07:58 +00:00
Jan Vesely 136381dc38 check_external_calls.sh: Print number of calls in tested file. 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315192
 2017-10-08 20:07:56 +00:00
Jan Vesely 80bb52ae75 ptx: Use __clc_nextafter to implement nextafter 
using clang builtin results in external library call

Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315191
 2017-10-08 19:34:00 +00:00
Jan Vesely 1de1444d62 Do not include clc_nextafter header globally 
Drop unused clc/math/clc_nextafter.h header

Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315190
 2017-10-08 19:33:58 +00:00
Jan Vesely 6a5c8ddb3a math/nextafter: Use custom declaration inc file 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315189
 2017-10-08 19:33:55 +00:00
Jan Vesely 72be1cc0be math/binary_decl.inc: Do not declare mixed float/double functions 
fmin/fmax only need vector/scalar mix

Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315188
 2017-10-08 19:33:53 +00:00
Jan Vesely beb6591753 ldexp: Fix double precision function return type 
Fixes ~1200 external calls from nvtpx library.

Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315170
 2017-10-08 06:56:14 +00:00
Jan Vesely 391305638c configure: Fix handling of directories with compats only source lists 
Reviewer: Jeroen Ketema
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 315018
 2017-10-05 20:16:28 +00:00
Jeroen Ketema 957151bd86 Add vload_half helpers for ptx 
The removes the vload_half unresolved calls from the nvptx libraries.

Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314998
 2017-10-05 18:17:40 +00:00
Jeroen Ketema feefb0870f Add vstore_half helpers for ptx 
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314925
 2017-10-04 19:07:48 +00:00
Jan Vesely a02d0e2c50 integer/sub_sat: Use clang builtin instead of llvm asm 
reviewer: Tom Stellard

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314703
 2017-10-02 18:39:03 +00:00
Jan Vesely 1964df8fad integer/add_sat: Use clang builtin instead of llvm asm 
reviewer: Tom Stellard

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314702
 2017-10-02 18:39:00 +00:00
Jan Vesely 943057a288 integer/clz: Use clang builtin instead of llvm asm 
The generated llvm IR mostly identical. char/uchar case is a bit worse.

reviewer: Tom Stellard

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314701
 2017-10-02 18:38:57 +00:00
Jeroen Ketema fe9fa89854 Let get_work_dim take exactly 0 arguments 
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314634
 2017-10-01 20:11:46 +00:00
Jeroen Ketema 17fdf263c5 Do no circularly define NULL 
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314633
 2017-10-01 20:10:14 +00:00
Jan Vesely 2b7fa1c6f6 Fix amdgcn-amdhsa on llvm-3.9 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 314548
 2017-09-29 19:06:52 +00:00
Jan Vesely aee030f284 travis: Check built libraries on llvm-3.9 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 314547
 2017-09-29 19:06:50 +00:00
Jan Vesely 8c8c287adf Add script to check for unresolved function calls 
v2: add shell shebang
    improve error checks and reporting
v3: fix typo

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 314546
 2017-09-29 19:06:48 +00:00
Jan Vesely 41b1500db0 geometric: geometric functions are only supported for vector lengths <=4 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 314545
 2017-09-29 19:06:47 +00:00
Jan Vesely 8d08f01eff travis: add build using llvm-3.9 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 314544
 2017-09-29 19:06:45 +00:00
Jan Vesely ce29e8cde1 Restore support for llvm-3.9 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 314543
 2017-09-29 19:06:41 +00:00
Jan Vesely 3bb50f6f7b Add missing HAVE_LLVM define to fix build with latest llvm 
Broken since r314111

V2: pointed out by Jan Vesely
   - Use format() instead of % formating

Patch-by: Pavel Ondračka <pavel.ondracka@gmail.com>
Signed-off-by: Pavel Ondračka <pavel.ondracka@gmail.com>
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314261
 2017-09-26 23:15:54 +00:00
Jan Vesely 1fa727d615 Rework atomic ops to use clang builtins rather than llvm asm 
reviewer: Aaron Watry

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314112
 2017-09-25 16:07:34 +00:00
Jan Vesely 760052047b prepare_builtins: Fix compile breakage with older LLVM 
Fixes r314050

reviewer: Tom Stellard

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 314111
 2017-09-25 16:04:37 +00:00
Reid Kleckner 3fc649cb76 [Support] Rename tool_output_file to ToolOutputFile, NFC 
This class isn't similar to anything from the STL, so it shouldn't use
the STL naming conventions.

llvm-svn: 314050
 2017-09-23 01:03:17 +00:00
Jan Vesely c9bbbe2403 Implement cl_khr_int64_extended_atomics builtins 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
Tested-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 313811
 2017-09-20 20:42:19 +00:00
Jan Vesely 1c81f4b0e3 Implement cl_khr_int64_base_atomics builtins 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
Tested-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 313810
 2017-09-20 20:42:14 +00:00
Jan Vesely d0320d5289 Add travis CI configuration file 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 313773
 2017-09-20 17:28:58 +00:00