커밋 그래프

152 커밋

작성자 SHA1 메시지 날짜
Evgeny Mankov 44007e29ab [HIPIFY][Blas] Sync with CUDA 9.1
[ROCm/hip commit: e5ba9668fc]
2018-05-08 20:42:30 +03:00
Evgeny Mankov ef30e4601e [HIPIFY] Sync with hipBLAS
[ROCm/hip commit: 7681775662]
2018-05-08 19:20:47 +03:00
Evgeny Mankov 057b13fc4a [HIPIFY] Initial cuDNN support
- based on https://github.com/ROCmSoftwarePlatform/hipDNN.
- lit testing was supplemented with CUDA_DNN_ROOT_DIR for cuDNN testing.
- single cuDNN test was added.


[ROCm/hip commit: fe421c89b2]
2018-05-03 11:33:40 +03:00
Evgeny Mankov 6b53512635 [HIPIFY] Sync with HIP (Execution Control, Surfaces, Memory)
Execution Control [deprecated since 7.0]:
cudaConfigureCall -> hipConfigureCall
cudaLaunch -> hipLaunchByPtr
cudaSetupArgument -> hipSetupArgument

Surfaces (partially):
cudaSurfaceObject_t -> hipSurfaceObject_t
cudaSurfaceBoundaryMode -> hipSurfaceBoundaryMode
cudaCreateSurfaceObject -> hipCreateSurfaceObject
cudaDestroySurfaceObject -> hipDestroySurfaceObject

Memory Management:
cudaMemset2DAsync -> hipMemset2DAsync


[ROCm/hip commit: 6f501f6094]
2018-04-24 22:04:04 +03:00
Evgeny Mankov 540fc92b1c [HIPIFY] undo apply .clangformat to HIPIFY source files
[ROCm/hip commit: 8499e79e2b]
2018-04-23 20:22:15 +03:00
Maneesh Gupta 4f42ee762d Apply .clangformat to all repo source files
Change-Id: I7e79c6058f0303f9a98911e3b7dd2e8596079344


[ROCm/hip commit: 1ba06f63c4]
2018-03-12 11:29:03 +05:30
Evgeny Mankov 1b25dfa25c [HIPIFY][fix] Build against llvm 3.8.0 fixed - missing include
/srv/HIP/hipify-clang/src/main.cpp:134:19: error: no member named 'dbgs' in namespace 'llvm'
      DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
            ~~~~~~^


[ROCm/hip commit: 7f3e89bbb0]
2018-02-08 00:08:47 +03:00
Evgeny Mankov 934c34e81f [HIPIFY][tests] add poisson-api-example test
+ add missing types
+ doc update


[ROCm/hip commit: 26095fe7c1]
2018-01-30 19:50:18 +03:00
Evgeny Mankov 58d1f6f3b3 [HIPIFY][tests] fix typo
[ROCm/hip commit: 9d814bef31]
2018-01-30 18:55:09 +03:00
Evgeny Mankov e455192444 [HIPIFY] InclusionDirective refactoring
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]
2018-01-29 18:33:47 +03:00
Evgeny Mankov 496111adc3 Merge pull request #340 from emankov/master
[HIPIFY][fix] Fix PragmaDirective

[ROCm/hip commit: fb1012a4b8]
2018-01-24 18:09:54 +03:00
Evgeny Mankov 369f3aa9fa [HIPIFY][fix] CUDA and cuBLAS main headers correct handling
[ROCm/hip commit: 600d5d7c06]
2018-01-23 23:43:36 +03:00
Evgeny Mankov fda3893ee3 [HIPIFY][fix] Fix PragmaDirective
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]
2018-01-23 23:06:55 +03:00
Evgeny Mankov 13aa621ce4 [HIPIFY] Sync with hipBLAS ToT and CUDA cuBLAS 9.1
[ROCm/hip commit: ddfb110080]
2018-01-22 17:12:02 +03:00
Evgeny Mankov 562844d378 [HIPIFY] cuRAND lib support (Device)
[ROCm/hip commit: 6e000adde4]
2018-01-19 21:29:05 +03:00
Evgeny Mankov 4d1fcf52e3 [HIPIFY] cuRAND lib support (partial - only Host)
[ROCm/hip commit: 8ff99eeadc]
2018-01-19 17:38:51 +03:00
Evgeny Mankov 5b9b271506 Merge pull request #332 from emankov/cudaMap_2
[HIPIFY] Add cudaMalloc3D support

[ROCm/hip commit: 0f7d687271]
2018-01-18 13:05:57 +03:00
Evgeny Mankov 7b7560f95d [HIPIFY] Add cudaMalloc3D support
[ROCm/hip commit: ff5f964c07]
2018-01-18 12:28:56 +03:00
Evgeny Mankov dad19d57f4 [HIPIFY] Add CUDA Driver API Texture Ref support (partial)
[ROCm/hip commit: 5788ac5d37]
2018-01-18 12:03:03 +03:00
Evgeny Mankov c52681edf2 [HIPIFY] Add more supported by HIP CUDA Driver API Arrays data types and functions
[ROCm/hip commit: 478fed74fe]
2018-01-16 21:07:50 +03:00
Evgeny Mankov 1506ec0d75 Merge pull request #319 from emankov/issue_211
[HIPIFY][fix][#211] Algorithm for explicit insert of hip include directive

[ROCm/hip commit: 88c5ffee3a]
2018-01-16 19:47:15 +03:00
Evgeny Mankov d5821de893 [HIPIFY] Add more supported by HIP CUDA RT API Textures and Arrays data types
[ROCm/hip commit: eb61038736]
2018-01-16 17:21:19 +03:00
Evgeny Mankov e09ba44b16 Update HipifyAction.cpp
dead code eliminate

[ROCm/hip commit: e54d9f3df0]
2018-01-16 15:08:08 +03:00
Evgeny Mankov 0e8da085bf [HIPIFY][fix][#211] Algorithm for explicit insert of hip include directive
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]
2018-01-15 21:05:05 +03:00
Evgeny Mankov 04bc4d4481 Merge pull request #307 from emankov/issue_306
[HIPIFY][FIX][#306] Eliminate second cuda main include directive

[ROCm/hip commit: f19cb9995e]
2018-01-10 22:50:44 +03:00
Evgeny Mankov 217d7031a4 [HIPIFY][fix][#306] Code improve
[ROCm/hip commit: 08662fba73]
2018-01-10 21:26:05 +03:00
Evgeny Mankov fd09c0eea8 [HIPIFY][#308][fix] Consume error returned by Replacements::add(...)
[ROCm/hip commit: 3f6de8bb10]
2018-01-09 20:03:53 +03:00
Evgeny Mankov 263cacc1c2 [HIPIFY][FIX][#306] Eliminate second cuda main include directive
// 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]
2017-12-26 20:54:54 +03:00
Evgeny Mankov 3233d633da [HIPIFY] Remove cudaBuiltin matcher
[ROCm/hip commit: 5d92a6c252]
2017-12-06 20:22:14 +03:00
Evgeny Mankov eefead5a1c [HIPIFY] Disable cudaBuiltin matcher.
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]
2017-12-05 20:28:51 +03:00
Evgeny Mankov af07df0b85 [HIPIFY] remove duplicates from CUDA_IDENTIFIER_MAP
[ROCm/hip commit: f24dfc6f36]
2017-12-05 19:46:53 +03:00
Evgeny Mankov 499e247586 Merge pull request #262 from ChrisKitching/frontendaction
[HIPIFY] Mostly fix preprocessor-or-template induced issues

[ROCm/hip commit: aa05b3d84e]
2017-11-27 17:30:11 +03:00
Chris Kitching 096136750c Use proper clang diagnostics for printing warnings
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]
2017-11-13 20:58:55 +00:00
Chris Kitching 78cf713140 Use a custom FrontendAction to simplify identifier translation
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]
2017-11-13 20:58:54 +00:00
Chris Kitching c2d54f0154 Add hipify mappings for all CUDA headers that have HIP equivalents
I'm particularly running into issues with `device_types.h` in real
CUDA code...


[ROCm/hip commit: 23b5d26582]
2017-11-13 17:20:07 +00:00
Evgeny Mankov 27c5e94c81 [HIPIFY] fix typo - missing )
[ROCm/hip commit: 44c74b6511]
2017-10-27 23:31:43 +03:00
Chris Kitching 54f786583b Remove commented else-block
A warning statement for _string literals_ seems a bit unhelpful.
There's no value in this being here.


[ROCm/hip commit: 20871a3a07]
2017-10-27 20:12:33 +01:00
Chris Kitching 05699e3779 Decouple the statistics system from the code translation
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]
2017-10-27 20:12:33 +01:00
Chris Kitching 1454bf9651 Copy-paste less in the statistics printing code
[ROCm/hip commit: 5699c18adc]
2017-10-27 20:12:33 +01:00
Chris Kitching e4569bc84e Inline updateCountersExt
[ROCm/hip commit: 50448aec3b]
2017-10-27 20:12:32 +01:00
Chris Kitching a3c1d30745 Update counter maps sanely
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]
2017-10-27 20:12:32 +01:00
Chris Kitching 51df5b20c9 Prefer references to pointers in updateCountersExt()
[ROCm/hip commit: 00bb447e55]
2017-10-27 20:12:32 +01:00
Chris Kitching 25517e41fd Move string utility functions into their own translation unit
[ROCm/hip commit: ee8e11a720]
2017-10-27 20:12:32 +01:00
Chris Kitching 85fd2e6f51 Extract LLVM compatibility code into its own translation unit
[ROCm/hip commit: 1bd837b4b1]
2017-10-27 20:12:32 +01:00
Chris Kitching 389fa2e68e Remove unused field
[ROCm/hip commit: 0c09bdf523]
2017-10-27 20:12:32 +01:00
Chris Kitching 48e7403762 Remove CUDA_EXCLUDES
An artefact from a now-defunct hack to avoid corrupting programs


[ROCm/hip commit: c6707ef33c]
2017-10-27 20:12:32 +01:00
Chris Kitching 199d75adc0 Make unsupported actually be a bool...
[ROCm/hip commit: 2f376c9b25]
2017-10-27 20:12:31 +01:00
Evgeny Mankov a9c3228cb1 Merge pull request #234 from ChrisKitching/warningSpam
[HIPIFY] Do not process __fetch_builtin_* in cudaCall()

[ROCm/hip commit: 9151a355c6]
2017-10-27 21:30:42 +03:00
Chris Kitching c15f6bdf5c Greatly enhance handling of macros in kernel launches
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]
2017-10-26 17:28:46 +01:00
Chris Kitching d0acfd5bde Simplify how kernel launch expressions get translated
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]
2017-10-26 17:28:30 +01:00