Due to support of cuRAND headers.
+ compound test on all headers is added;
+ missing entities are added with updating the doc;
+ a couple cuRAND tests are added (https://github.com/ROCmSoftwarePlatform/rocRAND/tree/master/benchmark):
- the following CUDA entities are still unsupported by hipRAND:
curandMakeMTGP32Constants
curandMakeMTGP32KernelState
curandGetDirectionVectors32
curandDirectionVectorSet_t
CURAND_DIRECTION_VECTORS_32_JOEKUO6
curandStateSobol64_t
curandStateScrambledSobol64_t
curandGenerateLongLong
- and the following - by HIP:
cudaRuntimeGetVersion
- those entities are handled by CHECK-NOT directive for now.
[ROCm/hip commit: 02e23c4d87]
File location have to be verified, otherwise location of the first found '#pragma once' in any included header even system will be erroneously handled, which might lead to attempt to including hip_runtime.h in it.
[ROCm/hip commit: 77f807b597]
If in source CUDA file main header (cuda_runtime.h or cuda.h) is not presented, corresponding HIP main header (hip_runtime.h) should be explicitly included in output hipified file.
[Algorithm]
1. If #pragma once is presented, HIP main header should be placed just after it;
2. Otherwise if any other (not CUDA main) header is presented, HIP main header should be placed just before it;
3. Otherwise HIP main header should be placed in the beginning of output file.
P.S.
There might be one more situation when #ifndef #define ... #endif guard for the entire file is presented (make sense for *.h, *.hpp, *.cuh files). In this case HIP main include should be placed just after such #ifdef, or after #pragma once, if it is also presented. This situation will be handled in a separate change.
[ROCm/hip commit: 09655a0853]
// hipified to #include<hip/hip_runtime.h>
#include<cuda.h> // 1st cuda main include (Driver API)
// to eliminate
#include<cuda_runtime.h> // 2nd cuda main include (Runtime API)
HIP has one header hip_runtime.h for both CUDA APIs, thus second cuda main include directive is eliminated entirely.
[ROCm/hip commit: 7e7cfa10cc]
As HIP has started to support vanilla CUDA syntax for threadIdx, blockIdx, blockDim and gridDim.
Other CUDA builtins are not tracked for now.
[ROCm/hip commit: 71d2fb20c8]
Much pretty. Very wow
This gives users all the usual power when it comes to manipulating
clang diagnostics. People can pass -Werror can have hipify fail if
it doesn't completely translate a file, for example. Much nicer
than reinventing the wheel.
[ROCm/hip commit: 6b767a59ba]
Most of what hipify does is really just replacing CUDA idenitifers
with HIP ones. CUDA function calls, preprocessor macro calls,
enum references, types, etc.
This is problematic: calls/types/enum-refs require name resolution
for the AST matcher to work. This fails in the presence of code
deleted by the preprocessor, and in two-pass template compilation.
Instead, we can simply hook the lexer and have it rewrite the
identifiers for us.
This approach means identifier transformations will work correctly
regardless of where they appear (and we get to delete lots of code)
- Fixes#260
- Helps a bit with #207 - it will still fail to translate kernel
calls in preprocessor-ignored code, but everything except kerel
launches should translate correctly now, even in
preprocessor-deleted code.
[ROCm/hip commit: 24cdc5e1d3]
The original implementation had the statistics system woken very
tightly into things like PPCallbacks, with counters duplicated
in two places, and all the output code duplicated. This made it
very difficult to alter the structure of the program without
breaking the statistics system.
Since the planned approach for solving the remaining preprocessor
bugs needs the introduction of a custom FrontendAction, and such
a restructure was incompatible with the way the statistics system
was set up, this rewrite was required.
'tis rather simpler now, mind you :D
This commit also fixes an issue where some stats were counted
twice, and allows `-print-stats` to operate independently of
`-stat-output`, allowing you to print stats to a file without
printing them to a terminal (or vice-versa).
[ROCm/hip commit: b303ffe53e]
operator[] default-constructs the map value if no value exists
for that key. Default-construction of int yields a zero. So all
the manual faffing around is just unnecessary.
[ROCm/hip commit: d8beee8918]
All but the most contrived use of macros is now properly handled -
have a look at the new testcases this commit adds. You can have
macros in kernel calls, macros spanning chunks of your arguments,
the call, call parameters, or callee can all be macros or
partially macros.
[ROCm/hip commit: 094b2b9b05]
It seems like there was a lot of machinery here that is no longer
needed now we have hipLaunchKernelGGL (which doesn't require us
to insert an extra argument into kernel functions). We no longer
need to waste cycles scanning the AST for callees.
We can literally just do "Take the callee expression, and dump
it into the first argument of hipLaunchKernelGGL()".
[ROCm/hip commit: eff86d975b]
There's three functions here that all do the same thing...
There was also logic that looks for numeric literals and works
backwards to find the macro name from which they are expanded.
I previously introduced code that rewrites macro references at
expand-time in the `MacroExpands` callback, so that code is no
longer doing anything useful.
[ROCm/hip commit: fd911e1839]
Calls to macros that were themselves CUDA API calls were often
being missed - this applies the identifier transform to macro
names at the callsites, too.
[ROCm/hip commit: d1e26b2e7e]
The source location for a call that's inside a macro body will,
by default, point into the macro definition itself. The original
logic was causing macro invocations to be overwritten, as I
explain here:
https://github.com/ROCm-Developer-Tools/HIP/issues/207#issuecomment-337521851
The existing PPCallbacks code is correctly rewriting macro
definitions, so the practical effect of this change is that AST
rewrites on code that's expanded from macros are no-ops.
It might be a performance optimisation to put a short-circiut at
the top of the AST callbacks to abort when faced with code that
was expanded from macros.
It might yet prove wise to do absolutely everything at lex-time...
[ROCm/hip commit: 4a794ed8c0]
A chain of 7 closing braces is never a great sign :D
In the process it became apparant that the unsupported flag
was being silently ignored, causing users to be left with cuda
API calls in their programs with no warning given. This has been
rectified for consistency.
[ROCm/hip commit: 35a892bc77]
Instead of having a single, enormous LUT for all CUDA names, let's
have separate ones for different types of entity. We often know
that we're looking at a typename, or a function name, or a macro
name - so we can be more efficient (and resilient to name
collisions) by having smaller lookup tables for each of those
classes of entity).
Here we start that off by having three LUTs:
- Header names
- Type names
- Everything else
Future work could usefully split "everything else" into:
- enum values
- macro names
- function names
- everything else
It's worth noting that the "needs new matcher" todos I delete here
were actually resolved with the previous commit. It no longer
naively searches for things that start with "cu*" - it will find
exactly those things that are present in our lookup tables.
[ROCm/hip commit: 695a1eb059]
Previously, there were different AST matchers for each
language construct that contains a type reference, and custom
logic to perform the transformation within each of those
structures.
Since the transformation in all such cases was only replacing
CUDA types with hip ones, we can instead use an AST matcher
that finds and updates the type references directly.
This simplifies the program considerably, and it won't fail
when it finds a language feature (or complicated type expression)
that nobody wrote custom logic for yet.
[ROCm/hip commit: a1d8340314]
Evgeny Mankov
Evgeny Mankov
Chris Kitching