summaryrefslogtreecommitdiffstats
path: root/clang/lib/Driver/ToolChains/Cuda.cpp
Commit message (Collapse)AuthorAgeFilesLines
* [CUDA] Warn about unsupported CUDA SDK version only if it's used.Artem Belevich2020-03-231-13/+19
| | | | | | | | | This fixes an issue with clang issuing a warning about unknown CUDA SDK if it's detected during non-CUDA compilation. Differential Revision: https://reviews.llvm.org/D76030 (cherry picked from commit eb2ba2ea953b5ea73cdbb598f77470bde1c6a011)
* [CUDA] Assume the latest known CUDA version if we've found an unknown one.Artem Belevich2020-01-291-28/+15
| | | | | | | | | | | | | | | | | This makes clang somewhat forward-compatible with new CUDA releases without having to patch it for every minor release without adding any new function. If an unknown version is found, clang issues a warning (can be disabled with -Wno-cuda-unknown-version) and assumes that it has detected the latest known version. CUDA releases are usually supersets of older ones feature-wise, so it should be sufficient to keep released clang versions working with minor CUDA updates without having to upgrade clang, too. Differential Revision: https://reviews.llvm.org/D73231 (cherry picked from commit 12fefeef203ab4ef52d19bcdbd4180608a4deae1)
* [Clang] Bypass distro detection on non-Linux hostsAlexandre Ganea2019-11-281-1/+2
| | | | | | Skip distro detection when we're not running on Linux, or when the target triple is not Linux. This saves a few OS calls for each invocation of clang.exe. Differential Revision: https://reviews.llvm.org/D70467
* [Clang][OpenMP Offload] Add new tool for wrapping offload device binariesSergey Dmitriev2019-10-091-2/+0
| | | | | | | | | | This patch removes the remaining part of the OpenMP offload linker scripts which was used for inserting device binaries into the output linked binary. Device binaries are now inserted into the host binary with a help of the wrapper bit-code file which contains device binaries as data. Wrapper bit-code file is dynamically created by the clang driver with a help of new tool clang-offload-wrapper which takes device binaries as input and produces bit-code file with required contents. Wrapper bit-code is then compiled to an object and resulting object is appended to the host linking by the clang driver. This is the second part of the patch for eliminating OpenMP linker script (please see https://reviews.llvm.org/D64943). Differential Revision: https://reviews.llvm.org/D68166 llvm-svn: 374219
* [HIP] Use option -nogpulib to disable linking device libYaxun Liu2019-10-031-2/+2
| | | | | | Differential Revision: https://reviews.llvm.org/D68300 llvm-svn: 373649
* [Clang] Migrate llvm::make_unique to std::make_uniqueJonas Devlieghere2019-08-141-3/+3
| | | | | | | | | | Now that we've moved to C++14, we no longer need the llvm::make_unique implementation from STLExtras.h. This patch is a mechanical replacement of (hopefully) all the llvm::make_unique instances across the monorepo. Differential revision: https://reviews.llvm.org/D66259 llvm-svn: 368942
* [CUDA][Clang][Bugfix] Add missing CUDA 9.2 caseGheorghe-Teodor Bercea2019-05-031-0/+3
| | | | | | | | | | | | | | | | | | | | | | | | Summary: The bug was reported on the OpenMP-dev list: .../obj-release/lib/clang/9.0.0/include/__clang_cuda_intrinsics.h:173:35: error: '__nvvm_shfl_sync_idx_i32' needs target feature ptx60|ptx61|ptx63|ptx64 __MAKE_SYNC_SHUFFLES(__shfl_sync, __nvvm_shfl_sync_idx_i32, This problem occurs when trying to compile a .cu file that requires a newer ptx version (>ptx60 in this case) than ptx42. Reviewers: tra, ABataev, caomhin Reviewed By: tra Subscribers: jdoerfert, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D61474 llvm-svn: 359910
* [CUDA] Do not pass deprecated option fo fatbinaryArtem Belevich2019-05-021-1/+2
| | | | | | | | | CUDA 10.1 tools deprecated some command line options. fatbinary no longer needs --cuda. Differential Revision: https://reviews.llvm.org/D61470 llvm-svn: 359838
* [CUDA] Implemented _[bi]mma* builtins.Artem Belevich2019-04-251-13/+19
| | | | | | | | | | | | | | | | These builtins provide access to the new integer and sub-integer variants of MMA (matrix multiply-accumulate) instructions provided by CUDA-10.x on sm_75 (AKA Turing) GPUs. Also added a feature for PTX 6.4. While Clang/LLVM does not generate any PTX instructions that need it, we still need to pass it through to ptxas in order to be able to compile code that uses the new 'mma' instruction as inline assembly (e.g used by NVIDIA's CUTLASS library https://github.com/NVIDIA/cutlass/blob/master/cutlass/arch/mma.h#L101) Differential Revision: https://reviews.llvm.org/D60279 llvm-svn: 359248
* Basic CUDA-10 support.Artem Belevich2019-02-051-0/+2
| | | | | | Differential Revision: https://reviews.llvm.org/D57771 llvm-svn: 353232
* [CUDA] Propagate detected version of CUDA to cc1Artem Belevich2019-01-311-1/+5
| | | | | | | | | | | | | ..and use it to control that parts of CUDA compilation that depend on the specific version of CUDA SDK. This patch has a placeholder for a 'new launch API' support which is in a separate patch. The list will be further extended in the upcoming patch to support CUDA-10.1. Differential Revision: https://reviews.llvm.org/D57487 llvm-svn: 352798
* Update the file headers across all of the LLVM projects in the monorepoChandler Carruth2019-01-191-4/+3
| | | | | | | | | | | | | | | | | to reflect the new license. We understand that people may be surprised that we're moving the header entirely to discuss the new license. We checked this carefully with the Foundation's lawyer and we believe this is the correct approach. Essentially, all code in the project is now made available by the LLVM project under our new license, so you will see that the license headers include that license only. Some of our contributors have contributed code under our old license, and accordingly, we have retained a copy of our old license notice in the top-level files in each project and repository. llvm-svn: 351636
* [CUDA][OPENMP][NVPTX]Improve logic of the debug info support.Alexey Bataev2018-12-121-26/+53
| | | | | | | | | | | | | | | | | | | | | | Summary: Added support for the -gline-directives-only option + fixed logic of the debug info for CUDA devices. If optimization level is O0, then options --[no-]cuda-noopt-device-debug do not affect the debug info level. If the optimization level is >O0, debug info options are used + --no-cuda-noopt-device-debug is used or no --cuda-noopt-device-debug is used, the optimization level for the device code is kept and the emission of the debug directives is used. If the opt level is > O0, debug info is requested + --cuda-noopt-device-debug option is used, the optimization is disabled for the device code + required debug info is emitted. Reviewers: tra, echristo Subscribers: aprantl, guansong, JDevlieghere, cfe-commits Differential Revision: https://reviews.llvm.org/D51554 llvm-svn: 348930
* [CUDA] Fix nvidia-cuda-toolkit detection on UbuntuJoel E. Denny2018-12-061-1/+1
| | | | | | | | | | This just extends D40453 (r319317) to Ubuntu. Reviewed By: Hahnfeld, tra Differential Revision: https://reviews.llvm.org/D55269 llvm-svn: 348504
* Lift VFS from clang to llvm (NFC)Jonas Devlieghere2018-10-101-1/+1
| | | | | | | | | | | | | | | | | | | This patch moves the virtual file system form clang to llvm so it can be used by more projects. Concretely the patch: - Moves VirtualFileSystem.{h|cpp} from clang/Basic to llvm/Support. - Moves the corresponding unit test from clang to llvm. - Moves the vfs namespace from clang::vfs to llvm::vfs. - Formats the lines affected by this change, mostly this is the result of the added llvm namespace. RFC on the mailing list: http://lists.llvm.org/pipermail/llvm-dev/2018-October/126657.html Differential revision: https://reviews.llvm.org/D52783 llvm-svn: 344140
* [HIP] Support early finalization of device code for -fno-gpu-rdcYaxun Liu2018-10-021-4/+4
| | | | | | | | | | | | | | | | | | | | | | | | This patch renames -f{no-}cuda-rdc to -f{no-}gpu-rdc and keeps the original options as aliases. When -fgpu-rdc is off, clang will assume the device code in each translation unit does not call external functions except those in the device library, therefore it is possible to compile the device code in each translation unit to self-contained kernels and embed them in the host object, so that the host object behaves like usual host object which can be linked by lld. The benefits of this feature is: 1. allow users to create static libraries which can be linked by host linker; 2. amortized device code linking time. This patch modifies HIP action builder to insert actions for linking device code and generating HIP fatbin, and pass HIP fatbin to host backend action. It extracts code for constructing command for generating HIP fatbin as a function so that it can be reused by early finalization. It also modifies codegen of HIP host constructor functions to embed the device fatbin when it is available. Differential Revision: https://reviews.llvm.org/D52377 llvm-svn: 343611
* [OpenMP] Improve search for libomptarget-nvptxJonas Hahnfeld2018-09-271-6/+14
| | | | | | | | | | | | | | | When looking for the bclib Clang considered the default library path first while it preferred directories in LIBRARY_PATH when constructing the invocation of nvlink. The latter actually makes more sense because during development it allows using a non-default runtime library. So change the search for the bclib to start looking in directories given by LIBRARY_PATH. Additionally add a new option --libomptarget-nvptx-path= which will be searched first. This will be handy for testing purposes. Differential Revision: https://reviews.llvm.org/D51686 llvm-svn: 343230
* [CUDA] Added basic support for compiling with CUDA-10.0Artem Belevich2018-09-241-1/+6
| | | | llvm-svn: 342924
* Rename -mlink-cuda-bitcode to -mlink-builtin-bitcodeMatt Arsenault2018-08-201-2/+2
| | | | | | | The same semantics work for OpenCL, and probably any offload language. Keep the old name around as an alias. llvm-svn: 340193
* [DEBUGINFO] Disable unsupported debug info options for NVPTX target.Alexey Bataev2018-07-271-0/+12
| | | | | | | | | | | | | | | | | | Summary: Some targets support only default set of the debug options and do not support additional debug options, like NVPTX target. Patch introduced virtual function supportsDebugInfoOptions() that can be overloaded by the toolchain, checks if the target supports some debug options and emits warning when an unsupported debug option is found. Reviewers: echristo Subscribers: aprantl, JDevlieghere, cfe-commits Differential Revision: https://reviews.llvm.org/D49148 llvm-svn: 338155
* [CUDA] Fixed the list of GPUs supported by CUDA-9.Artem Belevich2018-05-231-2/+2
| | | | | | Differential Revision: https://reviews.llvm.org/D47268 llvm-svn: 333098
* [CUDA] Added -f[no-]cuda-short-ptr optionArtem Belevich2018-05-091-2/+4
| | | | | | | | | The option enables use of 32-bit pointers for accessing const/local/shared memory. The feature is disabled by default. Differential Revision: https://reviews.llvm.org/D46148 llvm-svn: 331938
* [CUDA] Enable CUDA compilation with CUDA-9.2Artem Belevich2018-04-241-0/+2
| | | | | | Differential Revision: https://reviews.llvm.org/D45827 llvm-svn: 330753
* [NVPTX, CUDA] Added support for m8n32k16 and m32n8k16 variants of wmma ↵Artem Belevich2018-04-181-10/+12
| | | | | | | | | | instructions. The new instructions were added added for sm_70+ GPUs in CUDA-9.1. Differential Revision: https://reviews.llvm.org/D45068 llvm-svn: 330296
* [NVPTX] Emit debug info in DWARF-2 by default for Cuda devices.Alexey Bataev2018-04-181-3/+36
| | | | | | | | | | | | | | Summary: NVPTX target supports debug info in DWARF-2 format. Patch adds emission of debug info in DWARF-2 by default. Reviewers: tra, jlebar Subscribers: aprantl, JDevlieghere, cfe-commits Differential Revision: https://reviews.llvm.org/D42581 llvm-svn: 330272
* [CUDA] Added --[no-]cuda-include-ptx=sm_XX|all option.Artem Belevich2018-04-101-0/+19
| | | | | | | | | | | | | | Currently we always include PTX into the fatbin along with the GPU code.It about doubles the size of the GPU binary we need to carry in the executable. These options allow control inclusion of PTX into GPU binary. This patch does not change the defaults, though we may consider making no-PTX the default in the future. Differential Revision: https://reviews.llvm.org/D45495 llvm-svn: 329737
* Fix typos in clangAlexander Kornienko2018-04-061-1/+1
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Found via codespell -q 3 -I ../clang-whitelist.txt Where whitelist consists of: archtype cas classs checkk compres definit frome iff inteval ith lod methode nd optin ot pres statics te thru Patch by luzpaz! (This is a subset of D44188 that applies cleanly with a few files that have dubious fixes reverted.) Differential revision: https://reviews.llvm.org/D44188 llvm-svn: 329399
* [OpenMP] Add flag for linking runtime bitcode libraryGheorghe-Teodor Bercea2018-03-131-0/+38
| | | | | | | | | | | | | | Summary: This patch adds an additional flag to the OpenMP device offloading toolchain to link in the runtime library bitcode. Reviewers: Hahnfeld, ABataev, carlo.bertolli, caomhin, grokos, hfinkel Reviewed By: ABataev, grokos Subscribers: jholewinski, guansong, cfe-commits Differential Revision: https://reviews.llvm.org/D43197 llvm-svn: 327460
* Revert revision 327438.Gheorghe-Teodor Bercea2018-03-131-38/+0
| | | | llvm-svn: 327447
* [OpenMP] Add flag for linking runtime bitcode libraryGheorghe-Teodor Bercea2018-03-131-0/+39
| | | | | | | | | | | | | | Summary: This patch adds an additional flag to the OpenMP device offloading toolchain to link in the runtime library bitcode. Reviewers: Hahnfeld, ABataev, carlo.bertolli, caomhin, grokos, hfinkel Reviewed By: ABataev, grokos Subscribers: jholewinski, guansong, cfe-commits Differential Revision: https://reviews.llvm.org/D43197 llvm-svn: 327438
* [CUDA] Add option to generate relocatable device codeJonas Hahnfeld2018-02-121-5/+15
| | | | | | | | | | | As a first step, pass '-c/--compile-only' to ptxas so that it doesn't complain about references to external function. This will successfully generate object files, but they won't work at runtime because the registration routines need to adapted. Differential Revision: https://reviews.llvm.org/D42921 llvm-svn: 324878
* [CUDA] Detect installation in PATHJonas Hahnfeld2018-01-311-15/+50
| | | | | | | | | | | | | | | | If the CUDA toolkit is not installed to its default locations in /usr/local/cuda, the user is forced to specify --cuda-path. This is tedious and the driver can be smarter if well-known tools (like ptxas) can already be found in the PATH environment variable. Add option --cuda-path-ignore-env if the user wants to ignore set environment variables. Also use it in the tests to make sure the driver always finds the same CUDA installation, regardless of the user's environment. Differential Revision: https://reviews.llvm.org/D42642 llvm-svn: 323848
* [CUDA] Added partial support for CUDA-9.1Artem Belevich2018-01-301-5/+11
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Clang can use CUDA-9.1 now, though new APIs (are not implemented yet. The major change is that headers in CUDA-9.1 went through substantial changes that started in CUDA-9.0 which required substantial changes in the cuda compatibility headers provided by clang. There are two major issues: * CUDA SDK no longer provides declarations for libdevice functions. * A lot of device-side functions have become nvcc's builtins and CUDA headers no longer contain their implementations. This patch changes the way CUDA headers are handled if we compile with CUDA 9.x. Both 9.0 and 9.1 are affected. * Clang provides its own declarations of libdevice functions. * For CUDA-9.x clang now provides implementation of device-side 'standard library' functions using libdevice. This patch should not affect compilation with CUDA-8. There may be some observable differences for CUDA-9.0, though they are not expected to affect functionality. Tested: CUDA test-suite tests for all supported combinations of: CUDA: 7.0,7.5,8.0,9.0,9.1 GPU: sm_20, sm_35, sm_60, sm_70 Differential Revision: https://reviews.llvm.org/D42513 llvm-svn: 323713
* Fix function call to fix buildIsmail Donmez2017-11-291-1/+1
| | | | | | | | | ../tools/clang/lib/Driver/ToolChains/Cuda.cpp:80:18: error: reference to non-static member function must be called; did you mean to call it with no arguments? if (Distro(D.getVFS).IsDebian()) ~~^~~~~~ () llvm-svn: 319322
* Follow up of r319317, add the missing header fileSylvestre Ledru2017-11-291-0/+1
| | | | llvm-svn: 319319
* Add the nvidia-cuda-toolkit Debian package path to search pathSylvestre Ledru2017-11-291-0/+5
| | | | | | | | | | | | | | | | | | | Summary: Reported here: http://bugs.debian.org/882505 Patch by Andreas Beckmann Reviewers: Hahnfeld, tra Reviewed By: tra Subscribers: jlebar, cfe-commits Differential Revision: https://reviews.llvm.org/D40453 llvm-svn: 319317
* [OpenMP] Consistently use cubin extension for nvlinkJonas Hahnfeld2017-11-211-9/+17
| | | | | | | | | | | | This was previously done in some places, but for example not for bundling so that single object compilation with -c failed. In addition cubin was used for all file types during unbundling which is incorrect for assembly files that are passed to ptxas. Tighten up the tests so that we can't regress in that area. Differential Revision: https://reviews.llvm.org/D40250 llvm-svn: 318763
* [CUDA] Print an error if you try to compile with < sm_30 on CUDA 9.Justin Lebar2017-10-251-7/+9
| | | | | | | | | | | | | | Summary: CUDA 9's minimum sm is sm_30. Ideally we should also make sm_30 the default when compiling with CUDA 9, but that seems harder than it should be. Subscribers: sanjoy Differential Revision: https://reviews.llvm.org/D39109 llvm-svn: 316611
* [CMake][OpenMP] Customize default offloading archJonas Hahnfeld2017-10-171-10/+5
| | | | | | | | | For the shuffle instructions in reductions we need at least sm_30 but the user may want to customize the default architecture. Differential Revision: https://reviews.llvm.org/D38883 llvm-svn: 315996
* [CUDA] Require libdevice only if neededJonas Hahnfeld2017-10-161-13/+4
| | | | | | | | | If the user passes -nocudalib, we can live without it being present. Simplify the code by just checking whether LibDeviceMap is empty. Differential Revision: https://reviews.llvm.org/D38901 llvm-svn: 315902
* [OpenMP] Don't throw cudalib not found error if only front-end is required.Gheorghe-Teodor Bercea2017-09-261-0/+4
| | | | | | | | | | | | | | Summary: If we only use the compiler front-end, do not throw an error about the cuda device library not being found. This allows the front-end to be run on systems where no Cuda installation is found. Reviewers: Hahnfeld, ABataev, carlo.bertolli, caomhin, tra Reviewed By: tra Subscribers: hfinkel, tra, cfe-commits Differential Revision: https://reviews.llvm.org/D37914 llvm-svn: 314217
* [OpenMP] Enable the existing nocudalib flag for OpenMP offloading toolchain.Gheorghe-Teodor Bercea2017-09-251-3/+3
| | | | | | | | | | | | | | Summary: Enable the -nocudalib flag for the OpenMP device offloading toolchain as well. Currently it can only be used for the CUDA toolchain. Reviewers: Hahnfeld, ABataev, carlo.bertolli, caomhin, hfinkel, tra Reviewed By: tra Subscribers: hfinkel, cfe-commits Differential Revision: https://reviews.llvm.org/D37913 llvm-svn: 314164
* [OpenMP] Bugfix: output file name drops the absolute path where full path is ↵Gheorghe-Teodor Bercea2017-09-251-1/+1
| | | | | | | | | | | | | | | | needed. Summary: When composing the output file name, the path to the file is being dropped. The full path is required. Reviewers: Hahnfeld, ABataev, caomhin, carlo.bertolli, hfinkel, tra Reviewed By: tra Subscribers: hfinkel, tra, cfe-commits Differential Revision: https://reviews.llvm.org/D37912 llvm-svn: 314156
* Revert commit with wrong message.Gheorghe-Teodor Bercea2017-09-251-1/+1
| | | | llvm-svn: 314154
* [OpenMP] Don't throw cudalib not found error if only front-end is required.Gheorghe-Teodor Bercea2017-09-251-1/+1
| | | | | | | | | | | | | | Summary: If we only use the compiler front-end, do not throw an error about the cuda device library not being found. This allows the front-end to be run on systems where no Cuda installation is found. Reviewers: Hahnfeld, ABataev, carlo.bertolli, caomhin, tra Reviewed By: tra Subscribers: hfinkel, tra, cfe-commits Differential Revision: https://reviews.llvm.org/D37914 llvm-svn: 314150
* [NVPTX] Implemented shfl.sync instruction and supporting intrinsics/builtins.Artem Belevich2017-09-201-5/+11
| | | | | | Differential Revision: https://reviews.llvm.org/D38090 llvm-svn: 313820
* [CUDA] Added rudimentary support for CUDA-9 and sm_70.Artem Belevich2017-09-071-37/+51
| | | | | | | | | | | | | For now CUDA-9 is not included in the list of CUDA versions clang searches for, so the path to CUDA-9 must be explicitly passed via --cuda-path=. On LLVM side NVPTX added sm_70 GPU type which bumps required PTX version to 6.0, but otherwise is equivalent to sm_62 at the moment. Differential Revision: https://reviews.llvm.org/D37576 llvm-svn: 312734
* [OpenMP] Enable previously successful offloading tests.Gheorghe-Teodor Bercea2017-08-111-0/+13
| | | | | | | | | | | | Create a separate test file to contain all tests for OpenMP offloading to GPUs. Make libdevice checking more robust by accounting for the case in which no libdevice is found. This changes are in connrection with diff: D29660 llvm-svn: 310718
* [OpenMP] Delete tests in openmp-offload.c which cuase failuresGheorghe-Teodor Bercea2017-08-101-3/+7
| | | | | | | | until a better way to perform these tests is figured out. Change connected to diff: D29654 llvm-svn: 310625
* Revert r310489 and follow-up commits r310505, r310519, r310537 and r310549Alex Lorenz2017-08-101-14/+4
| | | | | | | | | | Commit r310489 caused 'openmp-offload.c' test failures on Darwin and other platforms: http://lab.llvm.org:8080/green/job/clang-stage1-cmake-RA-incremental_check/39230/testReport/junit/Clang/Driver/openmp_offload_c/ The follow-up commits tried to fix the test, but the test is still failing. llvm-svn: 310580
OpenPOWER on IntegriCloud