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!
Только зарегистрированные пользователи имеют доступ к сервису!
Для получения аккаунта, обратитесь к администратору.
Clang allows third party shared libraries to provide user-defined
extensions. For example, a custom libTemplateInstantiation.so can
visualize all template instantiation chains in clang. To enable it
one needs to pass a set of options such as -fplugin.
Cling should be able to inherently work with clang plugins. However,
cling still does not make full use of the clang driver where the plugin
setup is handled.
This patch enables plugins in cling and extends them in some aspects.
In particular, cling allows loading of plugins from shared libraries
but also if they are linked to the same library where cling is. This is
very useful in cases where cling runs itself in a shared library (eg
libCling). Users of libCling (such as ROOT) prefer to keep all llvm and
clang related symbols local to avoid symbol clashes if there is another
version of clang and llvm linked against a package. This can be done by
dlopen-ing libCling with RTLD_LOCAL visibility mode. Then the only way
for clang plugins to work in this scenario is to be linked to libCling.
Patch by Aleksandr Efremov and me.
- 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
- 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
Add llvm module pass to generate unique cuda module ctor/dtor names.
This llvm module pass address the follow problem. Every llvm module has a cuda ctor and dtor (if a cuda fatbinary exist), with at least a function call to register the fatbinary. The ctor/dtor can also include function calls to register global functions and variables at runtime, depending on user's code. The lazy compilation detects functions by the name. If the name (symbol) already exists it uses the existing translation. Otherwise it translates the function on first use (but it never translates twice). Without the module pass, Cling will always use the translation of the first module.
The testcase use the reflection of the gCling interpreter object. It takes two random modules and compare the symbols of the cuda module ctor and dtor.
Also add function, which change the symbol of the cuda module ctor and dtor to preprocessor compliant symbols.
Using unary operator address of (eg. MyClass m; &m) takes into account
overloaded operators which may not give us the precide address of the
allocated storage.
This patch teaches cling to use std::addressof instead.
