Commit Graph

302 Commits

Author SHA1 Message Date
Chris Kitching fd911e1839 Deduplicate preprocessor code
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.
2017-10-26 17:28:30 +01:00
Chris Kitching d1e26b2e7e Rewrite _all_ CUDA macro identifiers in the preprocessor
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.
2017-10-26 17:27:56 +01:00
Chris Kitching 4a794ed8c0 Don't special-case source locations for calls in macros
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...
2017-10-26 17:26:37 +01:00
Chris Kitching 35a892bc77 Prefer early-return to deep nesting
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.
2017-10-26 17:26:37 +01:00
Chris Kitching 2a5acac80e Do not process __fetch_builtin_* in cudaCall()
Fixes #205
2017-10-26 17:23:55 +01:00
Chris Kitching 778d6827f9 Refactor cudaCall to prefer early return to deep nesting
Sorry for the invasive refactor, but this was making reasoning
about this function more difficult.
2017-10-24 20:38:49 +01:00
Chris Kitching 695a1eb059 Split the giant lookup table into 3 smaller ones
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.
2017-10-24 20:38:49 +01:00
Chris Kitching a1d8340314 One matcher for type expressions to rule them all
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.
2017-10-24 20:38:49 +01:00
Chris Kitching 182d346356 Make control flow less insane
`while(false)` is certainly a bold choice.
2017-10-24 20:38:49 +01:00
Chris Kitching 4ab29113a1 Move giant lookup table into another translation unit
Also, rewrote it as a constant variable instead of a function
that imperatively fills a map. It's shorter, faster the compile,
and (depending on how badly the compiler screws it up) maybe
faster to run.

And, of course, it starts breaking up that giant .cpp file.
2017-10-24 20:38:25 +01:00
Chris Kitching 0207a29f78 Tweak some version numbers in clang version compatibility checks
Apparently a couple of those APIs changed in clang 5, not 4.

Drat.
2017-10-24 01:45:23 +01:00
Chris Kitching f5b273fc4f Avoid a double-free of HipifyPPCallbacks instance
This bug was present all along, but something changed in the order
of de-initialisation performed by llvm that makes it actually
crash now.

The constructor of HipifyPPCallbacks gives:
```
std::unique_ptr<HipifyPPCallbacks>(this)
```
to the LLVM Preprocessor instance. The Preprocessor instance
subsequently frees the HipifyPPCallbacks, which is then freed
again when we leave the stack frame at line 4340.

So: let's leak the HipifyPPCallbacks onto the heap, and leave
the LLVM Preprocessor object responsible for tidying it up.
2017-10-19 17:09:32 +01:00
Chris Kitching 6718519025 Be agnostic to the new getReplacements() API
See comment
2017-10-19 17:08:55 +01:00
Chris Kitching 24fc459f69 Use inline initialisers to set default field values
A trivial cleanup that helps in a moment..
2017-10-19 17:08:55 +01:00
Chris Kitching 7ae6a10c99 Omit now-removed Filename string arg from handleBeginSource 2017-10-19 17:08:55 +01:00
Chris Kitching 7e253365f1 Cope with Replacements now having llvm::Error returns 2017-10-19 17:08:54 +01:00
Chris Kitching 575bedb28c Cope with clang 4.0's rename of getNumArgs()
Sorry, this one I couldn't do in a perfectly elegant way ;)
2017-10-19 17:08:54 +01:00
Chris Kitching 0953a7887d Remove unnecessary call to Retain() on a smart pointer
The Preprocessor smart pointer is held by the CompilerInstance,
and therefore its reference count cannot reach zero until the
CompilerInstance itself is destroyed.

If the CompilerInstance is destroyed, you have more to worry about
than just the preprocessor being deallocated!

Newer versions of the LLVM/Clang API migrated to using
std::shared_ptr, so there is no `Retain()` function (by that
name, anyway). Eliminating this redundant use is a neat and
backward-compatible way to become compatible with newer versions
of the LLVM/Clang API.
2017-10-19 17:08:54 +01:00
Chris Kitching 2975d00edc Refer to clang::StringLiteral explicitly
Newer versions of llvm/clang mean there is both an
llvm::StringLiteral and a clang::StringLiteral. Since we're
dumping both namespaces wholesale into the global namespace with
`using` declarations, this creates a name collision, which must be
resolved.

This change is backwards-compatible, and fixes a problem you
encounter when using newer versions of the llvm/clang API.
2017-10-19 17:08:54 +01:00
Chris Kitching 3a2fe40f78 Fix two faulty LLVM version checks
What we actually want to do here is use the StringRef version in
versions newer than 3.8, and the void one in 3.8 and older.

Checking "major-version >= 3 && minor-version >= 9" does not do
what we want. Consider what this will do for version 4.0, for
which minor-version is zero...
2017-10-19 17:08:54 +01:00
Chris Kitching 4ba2b2e5b8 Inline a redundant std::vector
Looks like a partial refactor..?
2017-10-19 14:11:25 +01:00
Chris Kitching 742526706b Allow user-provided -std=* argument to override the default
This means any user-provided -std=* argument overrides this one,
allowing use of newer standards.
2017-10-19 14:11:15 +01:00
Evgeny Mankov db7661ccad Merge pull request #210 from ChrisKitching/output-hipify
[HIPIFY] Alter output file behaviour slightly
2017-10-18 21:46:31 +03:00
Chris Kitching 296c5a33ce Unconditionally append .hip as the default output filename 2017-10-18 19:19:18 +01:00
Chris Kitching e4e17a56bb If an output path is given _use it_
Don't append .hip to a user-provided output file...
2017-10-16 17:32:57 +01:00
Chris Kitching 2263cb9f72 Operate on a temp file in-place
This copies to the output after operation, instead of working
_on_ the output. This allows includes to work correctly, while
supporting output paths anywhere on the filesystem.

Fixes #208
Fixes #206
2017-10-16 17:32:56 +01:00
Evgeny Mankov 7ee8e2d51b [HIPIFY][#199][Partial fix] Fix for cudaLaunchKernel transformation
[Synopsis]
If any of kernel arguments is a MACRO its location calculation is wrong (location of its definition is actually calculated).
Thus garbage code is being produced on the place of such a MACRO starting from the end of its actual definition.

[Solution]
Add isMacroBodyExpansion and isMacroArgExpansion checks on kernel arguments.
2017-10-12 20:05:42 +03:00
Evgeny Mankov 53ce40ebab [HIPIFY] CUDA RT Textures and Arrays support update 2017-10-05 16:38:36 +03:00
Evgeny Mankov 28a300aed2 [HIPIFY] CUDA RT Textures support update 2017-09-18 17:52:50 +03:00
Evgeny Mankov d614f53ad7 Update Cuda2Hip.cpp 2017-07-28 01:45:24 +03:00
Evgeny Mankov ffcf5d32a0 [HIPIFY] CUDA RT Textures support 2017-07-27 18:39:40 +03:00
Evgeny Mankov 368aebc4bc [HIPIFY] CUDA Stream Callback support 2017-07-26 19:29:50 +03:00
Evgeny Mankov 85e6593a1b [HIPIFY] Finally finished syncing with CUDA 8.0.61 Driver and Runtime… (#105)
[hipify-clang] Finally finished syncing with CUDA 8.0.61 Driver and Runtime API (including missing data types, D3D, OpenGL, VDPAU and EGL interop).

    + All the Modules are supported now:
	1) 4.1 – 4.31 from CUDA 8.0.61 Driver API
	2) CUDA_Driver_API_functions_supported_by_HIP.md updated accordingly
	3) 4.1 – 4.31 from CUDA 8.0.61 Runtime API
	4) CUDA_Runtime_API_functions_supported_by_HIP.md updated accordingly
    + Typos fixing
    + Annotating
2017-07-26 09:38:03 +05:30
Evgeny Mankov b4a39664f0 [HIPIFY] Sync more CUDA Driver API functions.
+ 4.14. Event Management
+ 4.15. Execution Control

ToDo: 4.16 - 4.31 modules of CUDA Driver API.
2017-06-22 21:53:32 +03:00
Evgeny Mankov 0208fa4e70 [HIPIFY] Sync HIPIFY with HIP by CUDA Driver API functions.
+ 4.12. Unified Addressing
+ 4.13. Stream Management

ToDo: 4.14 - 4.31 modules of CUDA Driver API.
2017-06-14 19:55:55 +03:00
Evgeny Mankov f805bd0c1d [HIPIFY] Initial sync HIPIFY with HIP by CUDA Driver API functions.
+ CUDA_Driver_API_functions_supported_by_HIP.md update.
+ Initial update of HIPIFY with CUDA driver API functions:
  1.Error Handling, 2.Initialization, 3.Version Management, 5-6.Device Management, 7.Primary Context, 8-9.Context, 10.Module Management, 11.Memory Management.
+ Sync HIP functions against CUDA Driver and Runtime API functions.
+ Typo fixes.

ToDo: 12-30 modules of CUDA Driver API.
2017-06-13 20:25:11 +03:00
Evgeny Mankov 4a5484c616 [HIPIFY] All CUDA 8.0.44 API functions update
(for both Driver and Runtime APIs)

1) P2P
cuDeviceGetP2PAttribute   cudaDeviceGetP2PAttribute

2) Memory Mngmnt
cuMemPrefetchAsync        cudaMemPrefetchAsync
cuMemAdvise               cudaMemAdvise
cuMemRangeGetAttribute    cudaMemRangeGetAttribute
cuMemRangeGetAttributes   cudaMemRangeGetAttributes

3) Streams (Driver API only, no analogues in Runtime API)
cuStreamWaitValue32
cuStreamWaitValue32
cuStreamWriteValue32

4) Texture Reference Mngmnt (Driver API only, no analogues in Runtime API)
cuTexRefSetBorderColor
cuTexRefGetBorderColor
2017-06-01 21:08:33 +03:00
Evgeny Mankov 6e99e388ea [HIP] [HIPIFY] CUDA Driver API 8.0.44 JIT options support. 2017-05-31 18:55:29 +03:00
Evgeny Mankov 063539308e [HIPIFY] Add the rest CUDA Runtime API 8.0.44 Data structures.
+ sync with corresponding CUDA Driver API Data structures.

P.S.
There is no any new changes in CUDA Runtime API 8.0.61 Data structures since 8.0.44.
2017-05-30 19:45:59 +03:00
Evgeny Mankov 1cc5f42e34 [HIPIFY] Add the rest CUDA Driver API 8.0.44 Data structures.
+ Memory advise values
+ Memory Range Attributes
+ P2P Attributes

P.S.
There is no any new changes in CUDA Driver API 8.0.61 Data structures since 8.0.44.
2017-05-30 18:29:14 +03:00
Evgeny Mankov 3e99bc23e7 [HIPIFY] Add more CUDA Driver API 8.0.44 Data structures. 2017-05-30 17:58:13 +03:00
Evgeny Mankov c863215611 [FIX] [HIPIFY] Add matchers for function return types.
https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/issues/73

Examples (https://github.com/thrust/thrust/blob/master/thrust/system/cuda/detail/trivial_copy.inl):

template<typename System1,
         typename System2>
cudaStream_t cuda_memcpy_stream(const thrust::cpp::execution_policy<System1> &,
                                const thrust::cuda::execution_policy<System2> &exec)

template<typename System1,
         typename System2>
cudaMemcpyKind cuda_memcpy_kind(const thrust::cuda::execution_policy<System1> &,
                                const thrust::cpp::execution_policy<System2> &)
2017-05-24 18:25:40 +03:00
Evgeny Mankov 3d973dc4da [FIX] [HIPIFY] Matcher for new operator is missing.
https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/issues/80

Example from CUDA 8.0.44 sample (CUDASamples\0_Simple\matrixMulDrv\matrixMulDrv.cpp):
    CUjit_option *jitOptions = new CUjit_option[jitNumOptions];
where CUjit_option is enum, should be:
    hipJitOption *jitOptions = new hipJitOption[jitNumOptions];
2017-05-23 19:45:38 +03:00
Evgeny Mankov b5a1d47e68 [HIPIFY] [FIX] [HIPIFY] Matcher for pointer to enum var declaration is missing.
https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/issues/79

Example from CUDA 8.0.44 sample (CUDASamples\0_Simple\matrixMulDrv\matrixMulDrv.cpp):
    CUjit_option *jitOptions = new CUjit_option[jitNumOptions];
where CUjit_option is enum, should be:
    hipJitOption *jitOptions = new hipJitOption[jitNumOptions];

[TODO]
1. new CUjit_option -> new hipJitOption.
   Matcher for new operator is missing:
https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/issues/79
2. Merge matchers cudaEnumDecl and cudaEnumVarPtr.
2017-05-19 17:39:09 +03:00
Evgeny Mankov 270f643c9c [HIP] [HIPIFY] [FIX] cuModuleLoadDataEx -> hipModuleLoadDataEx
https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/issues/81

1.	Do not use JIT options on HCC path, call hipModuleLoadData instead.
2.	NVCC path is unchanged, to call cuModuleLoadDataEx with all options.
3.	Get rid of manual hipification, based on #ifdef #else for NVCC/HIP.
4.	Update documentation accordingly.
2017-05-19 17:22:14 +03:00
Evgeny Mankov 12d8c53c90 [HIPIFY] cudaMallocPitch -> hipMallocPitch 2017-05-16 18:21:25 +03:00
Evgeny Mankov a97cb6810c [HIPIFY] CUDA Driver API: Primary Context Management support. 2017-05-15 15:35:52 +03:00
Evgeny Mankov 25d470c380 [HIPIFY] Fix string routines.
Some Clang tooling functions return std::string, some return StringRef. Assigning of returning std::string to StringRef variables leads to garbage in it.
DEBUG build is always affected.
2017-05-11 21:50:36 +03:00
Evgeny Mankov 05be936fd6 [HIPIFY] LLVM 3.9 support
3.8 and 3.9 are both supported. 3.8 is stable, 3.9 needs more testing.
2017-05-05 21:28:02 +03:00
Evgeny Mankov 9b65358c68 [HIPIFY] CUDA RT memcpy functions update.
cudaMemcpyFromSymbol -> hipMemcpyFromSymbol
cudaMemcpyFromSymbolAsync -> hipMemcpyFromSymbolAsync
cudaMemcpy2DAsync -> hipMemcpy2DAsync
2017-05-03 23:05:44 +03:00