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

8 Commits

Author SHA1 Message Date
Yaxun Liu 1282889347 [HIP] Support new kernel launching API 
Differential Revision: https://reviews.llvm.org/D67947

llvm-svn: 372773
 2019-09-24 19:16:40 +00:00
Artem Belevich c62214da3d [CUDA] add support for the new kernel launch API in CUDA-9.2+. 
Instead of calling CUDA runtime to arrange function arguments,
the new API constructs arguments in a local array and the kernels
are launched with __cudaLaunchKernel().

The old API has been deprecated and is expected to go away
in the next CUDA release.

Differential Revision: https://reviews.llvm.org/D57488

llvm-svn: 352799
 2019-01-31 21:34:03 +00:00
Yaxun Liu 887c569bcb [HIP] Add hip input kind and codegen for kernel launching 
HIP is a language similar to CUDA (https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_kernel_language.md ).
The language syntax is very similar, which allows a hip program to be compiled as a CUDA program by Clang. The main difference
is the host API. HIP has a set of vendor neutral host API which can be implemented on different platforms. Currently there is open source
implementation of HIP runtime on amdgpu target (https://github.com/ROCm-Developer-Tools/HIP).

This patch adds support of input kind and language standard hip.

When hip file is compiled, both LangOpts.CUDA and LangOpts.HIP is turned on. This allows compilation of hip program as CUDA
in most cases and only special handling of hip program is needed LangOpts.HIP is checked.

This patch also adds support of kernel launching of HIP program using HIP host API.

When -x hip is not specified, there is no behaviour change for CUDA.

Patch by Greg Rodgers.
Revised and lit test added by Yaxun Liu.

Differential Revision: https://reviews.llvm.org/D44984

llvm-svn: 330790
 2018-04-25 01:10:37 +00:00
Reid Kleckner a769fd50ba Avoid depending on test inputes that aren't in Inputs 
Some people have weird CI systems that run each test subdirectory
independently without access to other parallel trees.

Unfortunately, this means we have to suffer some duplication until Art
can sort out how to share these types.

llvm-svn: 270164
 2016-05-20 00:38:25 +00:00
Justin Lebar 3039a593db [CUDA] Make printf work. 
Summary:
The code in CGCUDACall is largely based on a patch written by Eli
Bendersky:
http://lists.llvm.org/pipermail/llvm-commits/Week-of-Mon-20140324/210218.html

That patch implemented an LLVM pass lowering printf to vprintf; this
one does something similar, but in Clang codegen.

Reviewers: echristo

Subscribers: cfe-commits, jhen, tra, majnemer

Differential Revision: http://reviews.llvm.org/D16372

llvm-svn: 258642
 2016-01-23 21:28:14 +00:00
Artem Belevich 5d40ae3a46 Allow linking multiple bitcode files. 
Linking options for particular file depend on the option that specifies the file.
Currently there are two:

* -mlink-bitcode-file links in complete content of the specified file.
* -mlink-cuda-bitcode links in only the symbols needed by current TU.
   Linked symbols are internalized. This bitcode linking mode is used to
   link device-specific bitcode provided by CUDA.

Files are linked in order they are specified on command line.

-mlink-cuda-bitcode replaces -fcuda-uses-libdevice flag.

Differential Revision: http://reviews.llvm.org/D13913

llvm-svn: 251427
 2015-10-27 17:56:59 +00:00
Artem Belevich 7cb25c9b69 [CUDA] Postprocess bitcode linked in during device-side CUDA compilation. 
Link in and internalize the symbols we need from supplied bitcode library.

Differential Revision: http://reviews.llvm.org/D11664

llvm-svn: 247317
 2015-09-10 18:24:23 +00:00
Eli Bendersky 3468d9d929 Move all CUDA testing inputs to Inputs/ subdirectory inside the tests. 
llvm-svn: 207453
 2014-04-28 22:21:28 +00:00