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