Age | Commit message (Collapse) | Author | Files | Lines |
|
|
|
From: Tom Stellard <thomas.stellard at amd.com>
|
|
Implement mad24 and mul24 builtins
|
|
Conflicts:
generic/lib/shared/vload.cl
r600/lib/SOURCES
|
|
The assembly optimizations were making unsafe assumptions about which
address spaces had which identifiers.
Signed-off-by: Aaron Watry <awatry@gmail.com>
|
|
This was broken previously on Radeon SI.
|
|
libclc was defining and undefing GENTYPE and several other macros with
common names in its header files. This was preventing applications from
defining macros with identical names as command line arguments to the
compiler, because the definitions in the header files were masking the
macros defined as compiler arguements.
Reviewed-by: Aaron Watry <awatry@gmail.com>
|
|
Reviewed-by: Aaron Watry <awatry@gmail.com>
|
|
Conflicts:
configure.py
generic/include/clc/clc.h
generic/lib/SOURCES
generic/lib/shared/clamp.inc
generic/lib/shared/vload.cl
generic/lib/shared/vstore.cl
r600/lib/SOURCES
utils/prepare-builtins.cpp
|
|
The assembly should be generic, but at least currently R600 only supports
32-bit stores of [u]int1/4, and I believe that only global is well-supported.
R600 lowers the 8/16 component stores to multiple 4-component stores.
The unoptimized C versions of the other stuff is left in place.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185009 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
The assembly should be generic, but at least currently R600 only supports
32-bit loads of int1/4, and I believe that only global is well-supported.
R600 lowers the 8/16 component vectors to multiple 4-bit loads.
The unoptimized C versions of the other stuff is left in place.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185008 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Assumes that the target supports byte-addressable stores.
Completely unoptimized.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185007 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Should work for all targets and data types. Completely unoptimized.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185006 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185005 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Squashed commit of the following:
commit a0df0a0e86c55c1bdc0b9c0f5a739e5adef4b056
Author: Aaron Watry <awatry@gmail.com>
Date: Mon Apr 15 18:42:04 2013 -0500
libclc: Rename clz.ll to clz_if.ll to ensure it gets built.
configure.py treats files that have the same name with the .cl and .ll
extensions as overriding eachother.
E.g. If you have clz.cl and clz.ll both specified to be built in the same
SOURCES file, only the first file listed will actually be built.
Since the contents of clz.ll were an interface that is implemented in
clz_impl.ll, rename clz.ll to clz_if.ll to make sure that the interface is
built.
commit 931b62bed05c58f737de625bd415af09571a6a5a
Author: Aaron Watry <awatry@gmail.com>
Date: Sat Apr 13 12:32:54 2013 -0500
libclc: llvm assembly implementation of clz
Untested... currently crashes in the same manner as add_sat.
commit 6ef0b7b0b6d2e5584086b4b9a9243743b2e0538f
Author: Aaron Watry <awatry@gmail.com>
Date: Sat Mar 23 12:35:27 2013 -0500
libclc: Add stub clz builtin
For scalar int/uint, attempt to use the clz llvm builtin.. for all others
return 0 until an actual implementation is finished.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185004 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
For any GENTYPE that isn't scalar, we need to implement a mixed
vector/scalar version of clamp/max.
This depends on the min() patches I sent to the list a few minutes ago.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185003 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Checks if the current GENTYPE is scalar, and if not, then defines a separate
implementation of the function which casts the second arg to vector before
proceeding.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185002 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
This doesn't handle the integer cases for min(vector, scalar).
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185001 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
configure.py allows overloading *.cl with *.ll, but will only ever build
the first file listed in SOURCES of ${file}.cl and ${file}.ll
add_sat, sub_sat, (and the soon to be submitted clz) all define interfaces in
${function_name}.ll which are implemented in ${function_name}_impl.ll.
Renaming the interface files is enough to get them to build again, fixing
CL usage of these functions.
Tested on clover/r600g.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@185000 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184999 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184998 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Much more understandable/readable as a result, and probably more efficient.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184997 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
This implementation does a lot of bit shifting and masking. Suffice to say,
this is somewhat suboptimal... but it does look to produce correct results
(after the piglit tests were corrected for sign extension issues).
Someone who knows LLVM better than I could re-write this more efficiently.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184996 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Max(x,y) is available for all integer/floating types.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184995 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Created under a new shared/ directory for functions which are available for
both integer and floating point types.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184994 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184993 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Adds this function for both int and floating data types.
Patch by: Aaron Watry
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184992 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184991 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Patch by: Niels Ole Salscheider
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184990 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Patch by: Niels Ole Salscheider
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184989 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184988 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184987 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
static functions are not allowed in OpenCL C
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184986 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
These functions were not being compiled
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184985 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184984 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184983 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
Targets can override generic implementations by adding a file called
OVERRIDES in $(TARGET_DIR)/lib and listing the generic implementations
that it wants to override. For example, to override get_group_id() and
get_global_size() you would add these lines to the OVERRIDES file:
workitem/get_group_id.cl
workitem/get_global_size.cl
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184982 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
- First introducing a versioning scheme
- Add --libexecdir, --includedir and --pkgconfigdir and prefill them as well as --prefix
- Build all targets by default
- Create clc.pc and install it in $pkgconfigdir
- Use clang++ instead of c++
- Rename builtins.bc to built_libs/$triple.bc and install them in $libexecdir
- Includes are installed recursively to $includedir
- Finally add $(DESTDIR) for 'make install'
Patch by: Johannes Obermayr
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184981 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
This allows libclc to be built for R600 with upstream clang and LLVM.
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184980 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
to lib, and add header files in generic. Incorporates a patch by
Tom Stellard!
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184979 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184978 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184977 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184976 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
This includes a get_global_id() implementation and function stubs for
the other workitem and synchronization functions.
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@184975 91177308-0d34-0410-b5e6-96231b3b80d8
|
|
The bug in R600 which prevented them from working has been mitigated.
Signed-off-by: Aaron Watry<awatry@gmail.com>
|
|
From: Tom Stellard <thomas.stellard@amd.com>
|
|
Conflicts:
generic/lib/SOURCES
generic/lib/shared/vload.cl
generic/lib/shared/vstore.cl
|
|
From: Tom Stellard <thomas.stellard at amd.com>
|
|
The assembly should be generic, but at least currently R600 only supports
32-bit stores of [u]int1/4, and I believe that only global is well-supported.
R600 lowers the 8/16 component stores to multiple 4-component stores.
The unoptimized C versions of the other stuff is left in place.
|
|
The assembly should be generic, but at least currently R600 only supports
32-bit loads of int1/4, and I believe that only global is well-supported.
R600 lowers the 8/16 component vectors to multiple 4-bit loads.
The unoptimized C versions of the other stuff is left in place.
|