Commit Graph

37 Commits

Author SHA1 Message Date
Konstantin Zhuravlyov f8a81f869f [AMDGPU] Implement get_local_size for amdgcn--amdhsa triple
Differential Revision: http://reviews.llvm.org/D18284

llvm-svn: 265713
2016-04-07 19:54:19 +00:00
Matt Arsenault a48e15c6cb Split sources for amdgcn and r600
Most files remain in a common amdgpu directory.

Also switches barriers to to use convergent,
and use llvm.amdgcn.s.barrier.

This now requires 3.9/trunk to build amdgcn.

llvm-svn: 260777
2016-02-13 01:01:59 +00:00
Jan Vesely 6d870d2e84 configure: Remove llvm 3.6 defines
we require llvm 3.7

reviewer: tstellard
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 260304
2016-02-09 22:17:48 +00:00
Jan Vesely 46b7dd99e0 configure: Remove cl_khr_fp64 for device that don't support doubles
Also remove definitions if provided by clang (3.7+)
This halves the size of builtin.opt.{cedar,barts}.bc

reviewer: tstellard
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 260303
2016-02-09 22:17:46 +00:00
Jan Vesely 4f38142feb configure: Introduce per device defines
Make cl_khr_fp64 define per-device.
This patch does not change the generated Makefile (for llvm 3.6, 3.7)

v2: Make the device defines per LLVM version, 'all' for all versions

reviewer: tstellard
Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
llvm-svn: 260302
2016-02-09 22:17:45 +00:00
Tom Stellard 0abd28140b AMDGPU: Add aliases for all VI targets
llvm-svn: 255663
2015-12-15 18:37:04 +00:00
Tom Stellard 655680a22b AMDGPU: Add alias for tonga
Patch by: Vedran Mileti

llvm-svn: 255662
2015-12-15 18:37:02 +00:00
Jeroen Ketema d915739b38 Require LLVM >=3.7 and bump version to 0.2.0
v2: Also remove LLVM 3.6 traces from prepare-builtins.cpp

Patch by: EdB

llvm-svn: 244310
2015-08-07 08:31:37 +00:00
Tom Stellard ba742f58af Allow compilation depending to the LLVM version
It allows to keep temporary compatibilty with older version.
For exemple, this can be use when change are not to large.

Patch by: EdB

llvm-svn: 236113
2015-04-29 15:37:06 +00:00
Tom Stellard d9ca1f1596 configure: Add --enable-runtime-subnormal option
This makes it possible for runtime implementations to disable
subnormal handling at runtime.

When this flag is enabled, decisions about how to handle subnormals
in the library will be controlled by an external variable called
__CLC_SUBNORMAL_DISABLE.

Function implementations should use these new helpers for querying subnormal
support:
__clc_fp16_subnormals_supported();
__clc_fp32_subnormals_supported();
__clc_fp64_subnormals_supported();

In order for the library to link correctly with this feature,
users will be required to either:

1. Insert this variable into the module (if using the LLVM/Clang C++/C APIs).

2. Pass either subnormal_disable.bc or subnormal_use_default.bc to the
linker.  These files are distributed with liblclc and installed to
$(installdir).  e.g.:

llvm-link -o kernel-out.bc kernel.bc builtins-nosubnormal.bc subnormal_disable.bc

or

llvm-link -o kernel-out.bc kernel.bc builtins-nosubnormal.bc subnormal_use_default.bc

If you do not supply the --enable-runtime-subnormal then the library
behaves the same as it did before this commit.

In addition to these changes, the patch adds helper functions that
should be used when implementing library functions that need
special handling for denormals:

__clc_fp16_subnormals_supported();
__clc_fp32_subnormals_supported();
__clc_fp64_subnormals_supported();

llvm-svn: 235329
2015-04-20 18:49:50 +00:00
Tom Stellard 0f39721261 Use amdgcn triple for SI+ GPUs
llvm-svn: 225296
2015-01-06 20:42:12 +00:00
Tom Stellard 1d77071e0d Require LLVM 3.6 and bump version to 0.1.0
Some functions are implemented using hand-written LLVM IR, and
LLVM assembly format is allowed to change between versions, so we
should require a specific version of LLVM.

llvm-svn: 225041
2014-12-31 15:27:53 +00:00
Jan Vesely 62496142d5 configure: Add rpath to prepare-builtins util
v2: use space instead of '=' to make Mac happy

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Jeroen Ketema <j.ketema@imperial.ac.uk>
llvm-svn: 216922
2014-09-02 17:54:59 +00:00
Tom Stellard 425c181fb4 R600: Add aliases for hainan and mullins
llvm-svn: 216126
2014-08-20 21:23:01 +00:00
Tom Stellard d873c8e05a Add generic nvptx targets
These do not import the code specific to nvidiacl

Patch by: Jeroen Ketema

llvm-svn: 201431
2014-02-14 20:33:53 +00:00
Tom Stellard 55d3746dbe Revert "Enforce python2 for systems that use python3 as their default."
This reverts commit r200413.

This was breaking the build on systems where the python 2.x executable
was called python.

llvm-svn: 201239
2014-02-12 14:54:17 +00:00
Tom Stellard 8a3770ab97 Fixed rules names so they are unique when aliases are present.
This is necessary for building with Ninja because it does not allow
duplicate rule names.

Patch by: Dan Liew

llvm-svn: 200415
2014-01-29 20:03:27 +00:00
Tom Stellard 91d51db800 Fixed ninja build issues relating to use of $(DESTDIR)
We use ${DESTDIR} syntax now instead of $(DESTDIR) because that syntax
works both is the shell (at least it does for bash) and for make (at
least it does for GNU Make)

Patch By: Dan Liew

llvm-svn: 200414
2014-01-29 20:03:26 +00:00
Tom Stellard ac0fb621ce Enforce python2 for systems that use python3 as their default.
Patch by: Dan Liew

llvm-svn: 200413
2014-01-29 20:03:24 +00:00
Aaron Watry 8ef48d07ef Pass -fno-builtin flag to clang to silence warnings
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 198168
2013-12-29 16:39:55 +00:00
Aaron Watry b38037f7b7 Fix build with LLVM 3.5
Reviewed-by: Aaron Watry <awatry@gmail.com>
llvm-svn: 198167
2013-12-29 16:39:53 +00:00
Tom Stellard eedd3064de R600: Add aliases for Sea Islands GPUs
llvm-svn: 195023
2013-11-18 18:21:30 +00:00
Tom Stellard 9f48bb3b9a Make C++ compiler configurable
The C++ compiler used to build prepare-builtins
may differ from the llvm/clang for which we are
building libclc.

Use 'clang++' as the default compiler.

Patch by: Jeroen Ketema

llvm-svn: 193220
2013-10-23 02:49:27 +00:00
Tom Stellard f21e3ea972 Port pocl's gen_convert.py script to libclc
This script generates implementations for the entire set of convert_*
functions,

llvm-svn: 192385
2013-10-10 19:09:01 +00:00
Tom Stellard 6c7b86c106 Implement nextafter() builtin
There are two implementations of nextafter():
1. Using clang's __builtin_nextafter.  Clang replaces this builtin with
a call to nextafter which is part of libm.  Therefore, this
implementation will only work for targets with an implementation of
libm (e.g. most CPU targets).

2. The other implementation is written in OpenCL C.  This function is
known internally as __clc_nextafter and can be used by targets that
don't have access to libm.

llvm-svn: 192383
2013-10-10 19:08:51 +00:00
Tom Stellard 93d674f7b3 Place pkg-config file in $prefix/share/pkgconfig.
libclc is ABI-agnostic, and $prefix/lib/pkgconfig causes issues
on multilib setups. Using $prefix/share/pkgconfig allows us to reuse
a single libclc build across all system ABIs.

Patch by: Michał Górny

llvm-svn: 190107
2013-09-05 23:27:58 +00:00
Tom Stellard da920eab42 configure: Fix build when clang is installed to a non-standard prefix
llvm-svn: 188129
2013-08-10 03:40:26 +00:00
Tom Stellard 30f554b23d configure: Enable building separate libraries for target variants
llvm-svn: 184991
2013-06-26 18:20:38 +00:00
Tom Stellard 5668ea2d48 configure: fix out-of-source build
Patch by: Niels Ole Salscheider

llvm-svn: 184990
2013-06-26 18:20:35 +00:00
Tom Stellard d54f6ba7b8 Allow targets to override generic implementations
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

llvm-svn: 184982
2013-06-26 18:20:05 +00:00
Tom Stellard 1de7761ed9 Make libclc more Linux FHS conform.
- 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

llvm-svn: 184981
2013-06-26 18:20:03 +00:00
Tom Stellard 879327fcdc r600: Initial support
This includes a get_global_id() implementation and function stubs for
the other workitem and synchronization functions.

llvm-svn: 184975
2013-06-26 18:18:59 +00:00
Peter Collingbourne db47eac37d Fix build against recent versions of Clang. Based on patch by Alastair Donaldson!
llvm-svn: 169362
2012-12-05 07:39:02 +00:00
Peter Collingbourne bae6833419 configure.py: Add an install rule.
llvm-svn: 157821
2012-06-01 17:29:59 +00:00
Peter Collingbourne 9acfba3322 Enable cl_khr_fp64 when building the library, and fix several bugs
uncovered when doing so.

llvm-svn: 157617
2012-05-29 13:35:54 +00:00
Peter Collingbourne a3fc645609 Switch to the NVPTX backend.
llvm-svn: 157589
2012-05-28 20:42:15 +00:00
Peter Collingbourne d5395fbf03 Initial commit.
llvm-svn: 147756
2012-01-08 22:09:58 +00:00