IF YOU WOULD LIKE TO GET AN ACCOUNT, please write an
email to Administrator. User accounts are meant only to access repo
and report issues and/or generate pull requests.
This is a purpose-specific Git hosting for
BaseALT
projects. Thank you for your understanding!
Только зарегистрированные пользователи имеют доступ к сервису!
Для получения аккаунта, обратитесь к администратору.
Problem:
If a c++ statement with a concrete template parameter will type in, cling pass the (unwrapped) input line and the specialization of the statement, generated by compilerinstance to the clang nvptx. This causes an explicit-specialization-after-instantiation-error.
Solution:
Template specialization declaration will be save two times at a transaction. Once with the type kCCIHandleCXXImplicitFunctionInstantiation and once with the type kCCIHandleTopLevelDecl. To avoid sending a template specialization to the clang nvptx and causing an explicit-specialization-after-instantiation-error it have to check, which kCCIHandleTopLevelDecl declaration is also a kCCIHandleCXXImplicitFunctionInstantiation declaration.
As input of clang nvptx there are two possible sources. The raw input of `IncrementalParser::parseInternal()` or the AST-Printer of the `llvm::module`. Now, to decide, which input source should be use, I use the value `Transaction::ConsumerCallInfo::kCCIHandleTopLevelDecl`. This allows more C++ constructs to use.
Fixed a small bug. The AST-printer prints `setValueNoAlloc()` without the scope `cling::runtime::internal`. This caused an error, because clang could not find the declaration in the global space. The solution are dummy declaration in the global space.
At the moment, templates doesn't works.
- little changes at comments and code style
- try to use const in IncrementalCUDADeviceCompiler, where is possible
- move CUDA device code compiler instance to IncrementalParser
- change the members of CuArgs to const and adjust the setCuArgs method
- use std::vector<string> instead llvm::Smallvector<const char *> to build argv for executeAndWait
- improve the error messages of generatePCH(), generatePTX() and generateFatbinary()
- replace m_Counter with a copy in IncrementalCUDADeviceCompiler to avoid involuntary changes
Before this commit, the CUDA device code Compiler class overwrited source code .cu files, if the translation get wrong. Now, it renames the faulty files, so you can check, which source code cause compiler errors.
- add cudaDeviceSynchronize() at every kernel launch
- remove small address bug at cudaMemcpy, if host array is used
- in parallel test cases, replace fixes thread number with variable
- overworked shared memory kernel
- CUDA __constant__ memory
- CUDA global __device__ memory
- CUDA __host__ prefix
- CUDA kernel launch with arguments
- CUDA templated kernels
- CUDA shared memory with dynamic runtime
- CUDA Streams
- test if CUDA device is available
Before, it was not possible to find the clang++, which is contained in the cling, if we don't start the clang from the bin folder ('./cling -xcuda'). Now, for example it is possible to start the cling with 'bin/cling -xcuda' .
Fix a Bug, which avoid to start './cling -xcuda -fsyntax-only'.
In some cases, the path of the cling temp folder contains some non printable chars at the end.
Change the handling of the path string, to solve this problem.
Now, it is possible to declare variables on the prompt, which are visible for other statements.
The problems was, that cling wrapped all statements in a function, to get valid input. So, every variable is just visible inside in its own wrapper function. To solve the problem, cling change the local variable declaration to a global declaration.
The implementation for the CUDA compiler checks, if the unwrapping is happened. If it happened, the c++ code of the unwrapped variable declaration (AST-printer) instead the raw input will be written to the .cu-file.
At the moment, to extend the AST-tree of the device code, we use PCH-files to extend the exist device code with new lines of code. In detail, if we want to create a new PTX-file, we use the CUDA code (.cu file) and a PCH-file with the existing AST as input and generate an new PCH-file, which contains the whole AST. Then, the PCH-file will compiled to a PTX-file.
A bug in clang prevent, that we can’t generate more than 5 new PCH-files. The bug is not easy to fix, so I write a small workaround. Instead using a PCH-file, which contains the AST, we generate a new complete AST from all .cu-files every time.
The workaround is temporary and should removed, if clang is patched.
Now, it is possible to set some arguments of the clang nvptx and fatbinary via arguments at cling start. The arguments are filtered. So not every argument is possible at the moment. The Arguments can’t changed during runtime, because the PCH-files forbid it. For Example, the calng nvptx use the optimization level, which is set at start of cling.
At the moment, the debug options of clang nvptx are simple. If any debug option is detected, just a -g will add to the clang nvptx.
Additional PTX options for clang nvptx doesn’t works at the moment. There is a problem at parsing at the start of cling.
I replaced copies of the include paths with a pointer to the headerSearchOptions. Now, explicit handling of the include paths is not more necessary. Add include paths, which was declared via argument at start also works.
The class IncrementalCUDADeviceCompiler use external tools to generate PTX and cuda fatbin files. It runs the tools clang and fatbinary via llvm::sys::ExecuteAndWait. The class also handle to include new code in existing code. The steps of the compiler pipeline are:
- clang: CUDA C++ + previous PCH -> PCH
- clang: PCH -> PTX
- fatbinary: PTX -> fatbin
There is no selection of code. Every input of the cling will pass to the IncrementalCUDADeviceCompiler.
Now, it is possible to define functions with c++ attributes without the .rawInput mode. For example functions like `[[ noreturn ]] foo() { ... }` or `[[deprecated]] [[nodiscard]] int bar(){ … }`.
COMPILER="ccache clang" gets lost in CMake; using ccache does not work as there is no ccache-wrapper for clang-3.9.
So just use clang-3.9 without ccache.
This reverts commit 011aa8200277cd31957e222afd9b37415458b31f.
This is a revert of revert. I reverted the first commit because adding
"." to prebuiltmodulepath was causing failure in runtime modules, but
now we're skipping "." in TCling::LazyFunctionCreatorAutoloadForModule so
doesn't matter even if we have ".".
We had test failures in runtime nightlies such as this one:
https://epsft-jenkins.cern.ch/view/ROOT/job/root-nightly-runtime-cxxmodules/95/BUILDTYPE=Debug,COMPILER=gcc62,LABEL=slc6/testReport/junit/projectroot.roottest.root.math/smatrix/roottest_root_math_smatrix_testKalman/
Failures were due to what @pcanal commented in #2135, that some so files in
roottest doesn't have external linkage. (It means that if you call
dlopen(libfoo.so), linux kernel can't find dependency libraries and it
emits "undefined symbol" error when they try to initialize global
variables in libfoo.so but couldn't find symbol definition)
With pch, rootmap files were providing information about the depending library.
However we stopped generating rootmap files in #2127 and that's why we
got these failures. To fix this issue, I implemented a callback to
TCling which gets called when DynamicLibraryManager fails. The callback
pass error message to TCling and it handles message if it contains "undefined error".
