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

621 Commits

Author SHA1 Message Date
Jan Vesely 58fdb3b09a rootn: Use denormal path only 
It's OK to either flush to 0 or return denormal result if the device
does not support denormals. See sec 7.2 and 7.5.3 of OCL specs
Use 0.0f explicitly intead of relying on GPU to flush it.
Fixes CTS on carrizo and turks

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-by: Aaron Watry <awatry@gmail.com>
Tested-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 332324
 2018-05-15 04:22:43 +00:00
Jan Vesely 21e77037c0 remquo: Flush denormals if not supported 
It's OK to either flush to 0 or return denormal result if the device
does not support denormals. See sec 7.2 and 7.5.3 of OCL specs.
Fixes CTS on carrizo and turks.

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry <awatry@gmail.com>
llvm-svn: 331435
 2018-05-03 05:44:28 +00:00
Jan Vesely 8db45e4cf1 remquo: Port from amd builtins 
double version passes on carrizo. float version fails on denormals.

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry <awatry@gmail.com>
llvm-svn: 331434
 2018-05-03 05:44:26 +00:00
Jan Vesely 6146eda75d math: Add helper function to flush denormals if not supported. 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry <awatry@gmail.com>
llvm-svn: 331433
 2018-05-03 05:44:22 +00:00
Jan Vesely 616a38a693 clc_sqrt: Reuse unary_decl.inc 
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: 331366
 2018-05-02 16:06:52 +00:00
Jan Vesely 1647e50359 relational/select: Condition types for half are short/ushort, not char/uchar 
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 330851
 2018-04-25 17:36:36 +00:00
Jan Vesely e7d567ee0d log10: Use sw implementation from amd builtins 
Add missing table.
Fixes log10d CTS on carrizo.
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Acked-by: Aaron Watry <awatry@gmail.com>
Tested-by: Aaron Watry <awatry@gmail.com>

llvm-svn: 330649
 2018-04-23 21:10:42 +00:00
Jan Vesely 96591b6202 powr: Use denormal path only 
It's OK to either flush to 0 or return denormal result if the device
does not support denormals. See sec 7.2 and 7.5.3 of OCL specs
Fixes CTS on carrizo and turks.
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry <awatry@gmail.com>

llvm-svn: 330207
 2018-04-17 19:35:32 +00:00
Jan Vesely 4388d2883c pown: Use denormal path only 
It's OK to either flush to 0 or return denormal result if the device
does not support denormals. See sec 7.2 and 7.5.3 of OCL specs
Fixes CTS on carrizo and turks.
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry <awatry@gmail.com>

llvm-svn: 330206
 2018-04-17 19:35:30 +00:00
Jan Vesely 0d92f3047f pow: Use denormal path only 
It's OK to either flush to 0 or return denormal result if the device
does not support denormals. See sec 7.2 and 7.5.3 of OCL specs
Fixes CTS on carrizo and turks.
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewer: Aaron Watry <awatry@gmail.com>

llvm-svn: 330205
 2018-04-17 19:35:28 +00:00
Jan Vesely 8fa100dfe3 amdgcn/fmin: Fix typos that reduced precision 
Not sure how these sneaked in.
Fixes fminD and few other tests(fractD, cosD) on carrizo
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>

llvm-svn: 330198
 2018-04-17 18:11:29 +00:00
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