670 Commits

Author SHA1 Message Date
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
Aaron Watry
e62f5fa64d Add native_recip(x) as ((1)/(x))
Signed-off-by: Aaron Watry <awatry@gmail.com>
Acked-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 313107
2017-09-13 01:40:25 +00:00
Aaron Watry
415a60f303 integer: Add popcount implementation using ctpop intrinsic
Also copy/modify the unary_intrin.inc from math/ to make the
intrinsic declaration somewhat reusable.

Passes CL CTS integer_ops/test_integer_ops popcount tests for CL 1.2

Tested-by on GCN 1.0 (Pitcairn)

Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 312854
2017-09-09 02:23:54 +00:00
Jan Vesely
285d2fb85c Implement vload_half{,n} and vload(half)
v2: add vload(half) as well
    make helpers amdgpu specific (NVPTX uses different private AS numbering)
    use clang builtin on clang >= 6

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Tom Stellard <tstellar@redhat.com>
llvm-svn: 312839
2017-09-08 23:59:00 +00:00
Jan Vesely
661ac03a1b vstore: Cleanup and add vstore(half)
Add missing undefs
Make helpers amdgpu specific (NVPTX uses different numbering for private AS)
Use clang builtins on clang >= 6

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Tom Stellard <tstellar@redhat.com>
llvm-svn: 312838
2017-09-08 23:58:57 +00:00
Jan Vesely
b9dbaae3fb configure.py: Simplify compatibility sources
Just add the SOURCE_X.Y list to the list of sources if X.Y is the current llvm version.

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Tom Stellard <tstellar@redhat.com>
llvm-svn: 312837
2017-09-08 23:58:53 +00:00
Jan Vesely
3d1db3de74 amdgcn,waitcnt: Add datalayout info
This file is only compiled for GCN which all share the same layout

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 312493
2017-09-04 15:52:07 +00:00
Jan Vesely
e337b30c7d r600: Cleanup barrier implementation.
We don't have memory fences for r600 so just call group barrier directly
Make sure that barrier is called even with 0 flags

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 312492
2017-09-04 15:52:05 +00:00
Jan Vesely
1796d590c1 Fixup clc.h comment
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 312491
2017-09-04 15:52:03 +00:00
Aaron Watry
0bf96b1712 relational: Implement shuffle2 builtin
This was added in CL 1.1

Tested with a Radeon HD 7850 (Pitcairn) using the CL CTS via:
test_conformance/relationals/test_relationals shuffle_built_in_dual_input

v2: Add half support to shuffle2
    Move shuffle2 to misc/

Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 312404
2017-09-02 02:23:28 +00:00
Aaron Watry
880f15dae6 relational: Implement shuffle builtin
This was added in CL 1.1

Tested with a Radeon HD 7850 (Pitcairn) using the CL CTS via:
test_conformance/relationals/test_relationals shuffle_built_in

v2: Add half-precision support to shuffle when available.
    Move to misc/ and add section 6.12.12 to clc.h

Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 312403
2017-09-02 02:23:26 +00:00
Aaron Watry
da8dfefd1c Add halfN types and enable fp16 when generating builtin declarations
Uses the same mechanism to enable fp16 as we use for fp64 when
processing clc.h

Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 312402
2017-09-02 02:23:16 +00:00
Jan Vesely
999b1d9426 amdgcn: rewrite barrier() using fence and clang __builtin_amdgcn_s_barrier
Specs require using fences when barrier() is invoked:
"The barrier function will either flush any variables stored in local memory
or queue a memory fence to ensure correct ordering of memory operations to local memory."
and
"The barrier function will queue a memory fence to ensure correct ordering
of memory operations to global memory."

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: 311022
2017-08-16 17:09:00 +00:00
Jan Vesely
1977092dc3 amdgcn: Implement {read_,write_,}mem_fence builtin
v2: add more detailed comment about waitcnt instruction

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: 311021
2017-08-16 17:08:56 +00:00
Jan Vesely
7fc4c79fa5 configure.py: Drop explicit import of int builtin
I can't reproduce the error that made me add this.

Reported-by: Kim Gräsman <kim.grasman@gmail.com>
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Kim Gräsman <kim.grasman@gmail.com>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 310968
2017-08-15 22:24:05 +00:00
Jan Vesely
a4a20cd2f3 configure.py: Make python3 friendly
mostly prints and exceptions.
Few behavioral changes are documented in the text
Generated Makefile is identical between python2 and python3

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 309820
2017-08-02 15:00:59 +00:00
Jan Vesely
09f0a560e1 add __kernel_exec macros
also consolidate macros into one file, and rename to clcmacros.h

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 309358
2017-07-28 03:39:03 +00:00
Jan Vesely
2f2a3bc0dc generic: add missing get_work_dim include
Fixes few piglits since clang r304193

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 304556
2017-06-02 15:58:35 +00:00
Jan Vesely
9f7172965c math: Implement sinh function
mostly copied form amd_builtins

llvm-svn: 296233
2017-02-25 02:46:53 +00:00
Jan Vesely
c3868c8f8d .gitignore: Ignore amdgcn-mesa object directory
llvm-svn: 296164
2017-02-24 20:32:18 +00:00
Aaron Watry
dfec3c8e95 math: Add native_tan as wrapper to tan
Trivially define native_tan as a redirect to tan.

If there are any targets with a native implementation, we can deal with it later.

Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Matt Arsenault <arsenm2@gmail.com>
llvm-svn: 295920
2017-02-23 01:46:57 +00:00
Jeroen Ketema
80d2e8ffc1 Move BufferPtr into the block where it it being used
The previous location outside the block would crash prepare-builtins
when no the builtins file accidentially not passed on the command line.

llvm-svn: 294916
2017-02-12 21:33:49 +00:00
Jeroen Ketema
ed98e8d099 Add the correct prefixes to the cl_khr_fp64 pragma
llvm-svn: 294915
2017-02-12 21:31:41 +00:00
Matt Arsenault
9df2b9781c math: Add native_rsqrt builtin function
Trivial define to rsqrt.

Patch by Vedran Miletić <vedran@miletic.net>

llvm-svn: 294608
2017-02-09 18:39:26 +00:00
Aaron Watry
c606efabb7 math: Add logb builtin
Ported from the amd-builtins branch.

Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Matt Arsenault <Matthew.Arsenault@amd.com>
CC: Tom Stellard <thomas.stellard@amd.com>
llvm-svn: 292335
2017-01-18 03:14:10 +00:00
Aaron Watry
900bd7eb7f math: Add expm1 builtin function
Ported from the amd-builtins branch.

Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Matt Arsenault <Matthew.Arsenault@amd.com>
CC: Tom Stellard <thomas.stellard@amd.com>
llvm-svn: 292334
2017-01-18 03:13:37 +00:00
Tom Stellard
d83eb34ee7 Fix build since r286752.
llvm-svn: 286839
2016-11-14 16:06:33 +00:00
Tom Stellard
088faab429 Fix build since llvm r286566 and require at least llvm 4.0
llvm-svn: 286634
2016-11-11 21:34:47 +00:00
Jan Vesely
0a5aac3fc4 Provide vstore_half helper to workaround clc restrictions
clang won't accept half precision loads and stores without cl_khr_fp16 since r281904

llvm-svn: 282106
2016-09-21 20:15:55 +00:00
Tom Stellard
6b195ece57 configure: Add amdgcn-mesa-mesa3d target
llvm-svn: 281793
2016-09-16 22:43:33 +00:00
Tom Stellard
f19cf403c4 amdgcn-amdhsa: Add get_num_groups implementation
llvm-svn: 281792
2016-09-16 22:43:31 +00:00
Tom Stellard
e7ad23bad3 amdgcn-amdhsa: Add get_global_size() implementation
llvm-svn: 281791
2016-09-16 22:43:29 +00:00
Aaron Watry
af569547fa math: Implement tgamma
Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Tom Stellard <thomas.stellard@amd.com>
llvm-svn: 281566
2016-09-15 00:17:34 +00:00
Aaron Watry
e9009cdd21 math: Implement lgamma
Just use lgamma_r and ignore the value returned in the second argument

Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Tom Stellard <thomas.stellard@amd.com>
llvm-svn: 281565
2016-09-15 00:17:31 +00:00
Aaron Watry
0ab07e1bde math: Implement lgamma_r
Ported from the amd-builtins branch, which is itself based on the
Sun Microsystems implementation.

Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Tom Stellard <thomas.stellard@amd.com>
llvm-svn: 281564
2016-09-15 00:17:28 +00:00
Aaron Watry
f969413a82 Add ADDR_SPACE parameter to _CLC_V_V_VP_VECTORIZE
This macro is currently unused, but I plan to use it shortly.

The previous form did casts of pointers without an address space, which
doesn't work so well for CL 1.x.

Signed-off-by: Aaron Watry <awatry@gmail.com>
Reviewed-by: Tom Stellard <thomas.stellard@amd.com>
llvm-svn: 281563
2016-09-15 00:17:22 +00:00
Matt Arsenault
fbfd828d2a Replace nextafter implementation
This one passes conformance.

llvm-svn: 280961
2016-09-08 16:37:56 +00:00
Jan Vesely
eade17271a Avoid ambiguity in calling atom_add functions.
clang (since r280553) allows pointer casts in function overloads,
so we need to disambiguate the second argument.

clang might be smarter about overloads in the future
see https://reviews.llvm.org/D24113, but let's be safe in libclc anyway.

llvm-svn: 280871
2016-09-07 22:11:02 +00:00
Niels Ole Salscheider
63f71057c0 configure.py: Add polaris10 and polaris11
llvm-svn: 280121
2016-08-30 18:00:41 +00:00
Matt Arsenault
958fce3192 amdgcn: Fix return type of get_num_groups
llvm-svn: 279723
2016-08-25 07:31:40 +00:00
Matt Arsenault
7ef7e6aacd Strip opencl.ocl.version metadata
This should be uniqued when linking, but right now it creates
a lot of metadata spam listing the same version. This should also
probably be reporting the compiled version of the user program,
which may differ from the library. Currently the library IR files report
1.0 while 1.1/1.2 are the default for user programs.

llvm-svn: 279692
2016-08-25 00:25:10 +00:00