From 78826f551260033da2d7e6739d0b04f494144d7e Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Tue, 17 Oct 2017 15:30:38 +0100
Subject: [PATCH 01/28] Do not process __fetch_builtin_* in cudaCall()

Fixes #205


[ROCm/hip commit: 2a5acac80e2a4fe45d3aa02c05d55edc7af10c66]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 10 +++++++++-
 1 file changed, 9 insertions(+), 1 deletion(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index 4bffcdaed9..375ab8e8b2 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -814,7 +814,15 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac
           isExpansionInMainFile(),
           callee(
               functionDecl(
-                  matchesName("cu.*")
+                  matchesName("cu.*"),
+                  unless(
+                      // Clang generates structs with functions on them to represent things like
+                      // threadIdx.x. We have other logic to handle those builtins directly, so
+                      // we need to suppress the call-handling.
+                      // We can't handle those directly in the call-handler without special-casing
+                      // it unpleasantly, since the names of the functions are unique only per-struct.
+                      matchesName("__fetch_builtin.*")
+                  )
               )
           )
       ).bind("cudaCall"),

From 32cbe68a936666afa571c3e6347d4c20d4fc9224 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 18 Oct 2017 10:32:06 +0100
Subject: [PATCH 02/28] 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.


[ROCm/hip commit: 35a892bc77b37f6c1bcf2b87c4cc74b9539ae8ab]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 268 ++++++++++++---------
 1 file changed, 148 insertions(+), 120 deletions(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index 4bffcdaed9..046b3823ec 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -291,147 +291,175 @@ public:
                                   const FileEntry *file, StringRef search_path,
                                   StringRef relative_path,
                                   const clang::Module *imported) override {
-    if (_sm->isWrittenInMainFile(hash_loc)) {
-      if (is_angled) {
-        const auto found = CUDA_INCLUDE_MAP.find(file_name);
-        if (found != CUDA_INCLUDE_MAP.end()) {
-          updateCounters(found->second, file_name.str());
-          if (!found->second.unsupported) {
-            StringRef repName = found->second.hipName;
-            DEBUG(dbgs() << "Include file found: " << file_name << "\n"
-                         << "SourceLocation: "
-                         << filename_range.getBegin().printToString(*_sm) << "\n"
-                         << "Will be replaced with " << repName << "\n");
-            SourceLocation sl = filename_range.getBegin();
-            SourceLocation sle = filename_range.getEnd();
-            const char *B = _sm->getCharacterData(sl);
-            const char *E = _sm->getCharacterData(sle);
-            SmallString<128> tmpData;
-            Replacement Rep(*_sm, sl, E - B, Twine("<" + repName + ">").toStringRef(tmpData));
-            FullSourceLoc fullSL(sl, *_sm);
-            insertReplacement(Rep, fullSL);
-          }
-        } else {
-//          llvm::outs() << "[HIPIFY] warning: the following reference is not handled: '" << file_name << "' [inclusion directive].\n";
-        }
-      }
+    if (!_sm->isWrittenInMainFile(hash_loc) || !is_angled) {
+      return; // We're looking to rewrite angle-includes in the main file to point to hip.
     }
+
+    const auto found = CUDA_INCLUDE_MAP.find(file_name);
+    if (found == CUDA_INCLUDE_MAP.end()) {
+      // Not a CUDA include - don't touch it.
+      return;
+    }
+
+    updateCounters(found->second, file_name.str());
+    if (found->second.unsupported) {
+      // An unsupported CUDA header? Oh dear. Print a warning.
+      printHipifyMessage(*_sm, hash_loc, "Unsupported CUDA header used: " + file_name.str());
+      return;
+    }
+
+    StringRef repName = found->second.hipName;
+    DEBUG(dbgs() << "Include file found: " << file_name << "\n"
+                 << "SourceLocation: "
+                 << filename_range.getBegin().printToString(*_sm) << "\n"
+                 << "Will be replaced with " << repName << "\n");
+    SourceLocation sl = filename_range.getBegin();
+    SourceLocation sle = filename_range.getEnd();
+    const char *B = _sm->getCharacterData(sl);
+    const char *E = _sm->getCharacterData(sle);
+    SmallString<128> tmpData;
+    Replacement Rep(*_sm, sl, E - B, Twine("<" + repName + ">").toStringRef(tmpData));
+    FullSourceLoc fullSL(sl, *_sm);
+    insertReplacement(Rep, fullSL);
   }
 
   virtual void MacroDefined(const Token &MacroNameTok,
                             const MacroDirective *MD) override {
-    if (_sm->isWrittenInMainFile(MD->getLocation()) &&
-        MD->getKind() == MacroDirective::MD_Define) {
-      for (auto T : MD->getMacroInfo()->tokens()) {
-        if (T.isAnyIdentifier()) {
-          StringRef name = T.getIdentifierInfo()->getName();
-          const auto found = CUDA_RENAMES_MAP().find(name);
-          if (found != CUDA_RENAMES_MAP().end()) {
-            updateCounters(found->second, name.str());
-            if (!found->second.unsupported) {
-              StringRef repName = found->second.hipName;
-              SourceLocation sl = T.getLocation();
-              DEBUG(dbgs() << "Identifier " << name << " found in definition of macro "
-                           << MacroNameTok.getIdentifierInfo()->getName() << "\n"
-                           << "will be replaced with: " << repName << "\n"
-                           << "SourceLocation: " << sl.printToString(*_sm) << "\n");
-              Replacement Rep(*_sm, sl, name.size(), repName);
-              FullSourceLoc fullSL(sl, *_sm);
-              insertReplacement(Rep, fullSL);
-            }
-          } else {
-            // llvm::outs() << "[HIPIFY] warning: the following reference is not handled: '" << name << "' [macro].\n";
-          }
-        }
+    if (!_sm->isWrittenInMainFile(MD->getLocation()) ||
+        MD->getKind() != MacroDirective::MD_Define) {
+      return;
+    }
+
+    for (auto T : MD->getMacroInfo()->tokens()) {
+      // We're looking for CUDA identifiers in the macro definition to rewrite...
+      if (!T.isAnyIdentifier()) {
+        continue;
       }
+
+      StringRef name = T.getIdentifierInfo()->getName();
+      const auto found = CUDA_RENAMES_MAP().find(name);
+      if (found == CUDA_RENAMES_MAP().end()) {
+        // So it's an identifier that isn't CUDA? Boring.
+        continue;
+      }
+
+      updateCounters(found->second, name.str());
+      SourceLocation sl = T.getLocation();
+      if (found->second.unsupported) {
+        // An unsupported identifier? Curses! Warn the user.
+        printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier used: " + name.str());
+        continue;
+      }
+
+      StringRef repName = found->second.hipName;
+      DEBUG(dbgs() << "Identifier " << name << " found in definition of macro "
+                   << MacroNameTok.getIdentifierInfo()->getName() << "\n"
+                   << "will be replaced with: " << repName << "\n"
+                   << "SourceLocation: " << sl.printToString(*_sm) << "\n");
+      Replacement Rep(*_sm, sl, name.size(), repName);
+      FullSourceLoc fullSL(sl, *_sm);
+      insertReplacement(Rep, fullSL);
     }
   }
 
   virtual void MacroExpands(const Token &MacroNameTok,
                             const MacroDefinition &MD, SourceRange Range,
                             const MacroArgs *Args) override {
-    if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) {
-        // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics
-        // remain unchanged.
+
+    if (!_sm->isWrittenInMainFile(MacroNameTok.getLocation())) {
+      return; // Macros in headers are not our concern.
+    }
+
+    // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics
+    // remain unchanged.
 #if LLVM_VERSION_MAJOR > 4
-        #define GET_NUM_ARGS() getNumParams()
+    #define GET_NUM_ARGS() getNumParams()
 #else
-        #define GET_NUM_ARGS() getNumArgs()
+    #define GET_NUM_ARGS() getNumArgs()
 #endif
-      for (unsigned int i = 0; Args && i < MD.getMacroInfo()->GET_NUM_ARGS(); i++) {
-        std::vector<Token> toks;
-        // Code below is a kind of stolen from 'MacroArgs::getPreExpArgument'
-        // to workaround the 'const' MacroArgs passed into this hook.
-        const Token *start = Args->getUnexpArgument(i);
-        size_t len = Args->getArgLength(start) + 1;
+
+    for (unsigned int i = 0; Args && i < MD.getMacroInfo()->GET_NUM_ARGS(); i++) {
+      std::vector<Token> toks;
+      // Code below is a kind of stolen from 'MacroArgs::getPreExpArgument'
+      // to workaround the 'const' MacroArgs passed into this hook.
+      const Token *start = Args->getUnexpArgument(i);
+      size_t len = Args->getArgLength(start) + 1;
 #if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
-        _pp->EnterTokenStream(start, len, false, false);
+      _pp->EnterTokenStream(start, len, false, false);
 #else
-        _pp->EnterTokenStream(ArrayRef<Token>(start, len), false);
+      _pp->EnterTokenStream(ArrayRef<Token>(start, len), false);
 #endif
-        do {
-          toks.push_back(Token());
-          Token &tk = toks.back();
-          _pp->Lex(tk);
-        } while (toks.back().isNot(tok::eof));
-        _pp->RemoveTopOfLexerStack();
-        // end of stolen code
-        for (auto tok : toks) {
-          if (tok.isAnyIdentifier()) {
-            StringRef name = tok.getIdentifierInfo()->getName();
+      do {
+        toks.push_back(Token());
+        Token &tk = toks.back();
+        _pp->Lex(tk);
+      } while (toks.back().isNot(tok::eof));
+
+      _pp->RemoveTopOfLexerStack();
+      // end of stolen code
+
+      for (auto tok : toks) {
+        if (tok.isAnyIdentifier()) {
+          StringRef name = tok.getIdentifierInfo()->getName();
+          SourceLocation sl = tok.getLocation();
+
+          const auto found = CUDA_RENAMES_MAP().find(name);
+          if (found == CUDA_RENAMES_MAP().end()) {
+            // It's not a CUDA identifier. We have nothing to do.
+            continue;
+          }
+
+          updateCounters(found->second, name.str());
+          if (found->second.unsupported) {
+            // We know about it, but it isn't supported. Warn the user.
+            printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier: " + name.str());
+            continue;
+          }
+
+          StringRef repName = found->second.hipName;
+          DEBUG(dbgs() << "Identifier " << name
+                       << " found as an actual argument in expansion of macro "
+                       << MacroNameTok.getIdentifierInfo()->getName() << "\n"
+                       << "will be replaced with: " << repName << "\n");
+          size_t length = name.size();
+          if (_sm->isMacroBodyExpansion(sl)) {
+            LangOptions DefaultLangOptions;
+            SourceLocation sl_macro = _sm->getExpansionLoc(sl);
+            SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions);
+            length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro);
+            sl = sl_macro;
+          }
+          Replacement Rep(*_sm, sl, length, repName);
+          FullSourceLoc fullSL(sl, *_sm);
+          insertReplacement(Rep, fullSL);
+        } else if (tok.isLiteral()) {
+          SourceLocation sl = tok.getLocation();
+          if (_sm->isMacroBodyExpansion(sl)) {
+            LangOptions DefaultLangOptions;
+            SourceLocation sl_macro = _sm->getExpansionLoc(sl);
+            SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions);
+            size_t length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro);
+            StringRef name = StringRef(_sm->getCharacterData(sl_macro), length);
+
             const auto found = CUDA_RENAMES_MAP().find(name);
-            if (found != CUDA_RENAMES_MAP().end()) {
-              updateCounters(found->second, name.str());
-              if (!found->second.unsupported) {
-                StringRef repName = found->second.hipName;
-                DEBUG(dbgs() << "Identifier " << name
-                             << " found as an actual argument in expansion of macro "
-                             << MacroNameTok.getIdentifierInfo()->getName() << "\n"
-                             << "will be replaced with: " << repName << "\n");
-                size_t length = name.size();
-                SourceLocation sl = tok.getLocation();
-                if (_sm->isMacroBodyExpansion(sl)) {
-                  LangOptions DefaultLangOptions;
-                  SourceLocation sl_macro = _sm->getExpansionLoc(sl);
-                  SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions);
-                  length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro);
-                  name = StringRef(_sm->getCharacterData(sl_macro), length);
-                  sl = sl_macro;
-                }
-                Replacement Rep(*_sm, sl, length, repName);
-                FullSourceLoc fullSL(sl, *_sm);
-                insertReplacement(Rep, fullSL);
-              }
-            } else {
-              // llvm::outs() << "[HIPIFY] warning: the following reference is not handled: '" << name << "' [macro expansion].\n";
+            if (found == CUDA_RENAMES_MAP().end()) {
+              continue; // Not CUDA, we don't care.
             }
-          } else if (tok.isLiteral()) {
-            SourceLocation sl = tok.getLocation();
-            if (_sm->isMacroBodyExpansion(sl)) {
-              LangOptions DefaultLangOptions;
-              SourceLocation sl_macro = _sm->getExpansionLoc(sl);
-              SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions);
-              size_t length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro);
-              StringRef name = StringRef(_sm->getCharacterData(sl_macro), length);
-              const auto found = CUDA_RENAMES_MAP().find(name);
-              if (found != CUDA_RENAMES_MAP().end()) {
-                updateCounters(found->second, name.str());
-                if (!found->second.unsupported) {
-                  StringRef repName = found->second.hipName;
-                  sl = sl_macro;
-                  Replacement Rep(*_sm, sl, length, repName);
-                  FullSourceLoc fullSL(sl, *_sm);
-                  insertReplacement(Rep, fullSL);
-                }
-              } else {
-                // llvm::outs() << "[HIPIFY] warning: the following reference is not handled: '" << name << "' [literal macro expansion].\n";
-              }
-            } else {
-              if (tok.is(tok::string_literal)) {
-                StringRef s(tok.getLiteralData(), tok.getLength());
-                processString(unquoteStr(s), *_sm, tok.getLocation());
-              }
+
+            updateCounters(found->second, name.str());
+            if (found->second.unsupported) {
+              printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier: " + name.str());
+              continue;
             }
+
+            StringRef repName = found->second.hipName;
+            sl = sl_macro;
+            Replacement Rep(*_sm, sl, length, repName);
+            FullSourceLoc fullSL(sl, *_sm);
+            insertReplacement(Rep, fullSL);
+          } else if (tok.is(tok::string_literal)) {
+            StringRef s(tok.getLiteralData(), tok.getLength());
+            processString(unquoteStr(s), *_sm, tok.getLocation());
           }
         }
       }

From cbf786a8fd2f3ed82da789837b72b0806bb591f6 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 18 Oct 2017 11:25:17 +0100
Subject: [PATCH 03/28] 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...


[ROCm/hip commit: 4a794ed8c091df88565e0d2cc1b937083b2c509f]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 26 ++++------------------
 1 file changed, 4 insertions(+), 22 deletions(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index 046b3823ec..a3ec11ff1a 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -531,28 +531,10 @@ private:
     }
 
     size_t length = name.size();
-    bool bReplace = true;
-    if (SM->isMacroArgExpansion(sl)) {
-      sl = SM->getImmediateSpellingLoc(sl);
-    } else if (SM->isMacroBodyExpansion(sl)) {
-      LangOptions DefaultLangOptions;
-      SourceLocation sl_macro = SM->getExpansionLoc(sl);
-      SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *SM, DefaultLangOptions);
-      length = SM->getCharacterData(sl_end) - SM->getCharacterData(sl_macro);
-      StringRef macroName = StringRef(SM->getCharacterData(sl_macro), length);
-      if (CUDA_EXCLUDES.end() != CUDA_EXCLUDES.find(macroName)) {
-        bReplace = false;
-      } else {
-        sl = sl_macro;
-      }
-    }
-
-    if (bReplace) {
-      updateCounters(found->second, name);
-      Replacement Rep(*SM, sl, length, hipCtr.hipName);
-      FullSourceLoc fullSL(sl, *SM);
-      insertReplacement(Rep, fullSL);
-    }
+    updateCounters(found->second, name);
+    Replacement Rep(*SM, sl, length, hipCtr.hipName);
+    FullSourceLoc fullSL(sl, *SM);
+    insertReplacement(Rep, fullSL);
 
     return true;
   }

From ca913bb196cce952567ec2dd001590249baa6044 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 18 Oct 2017 14:25:14 +0100
Subject: [PATCH 04/28] 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.


[ROCm/hip commit: d1e26b2e7e0fb604e0c7f0912afeca83d7c12580]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 27 ++++++++++++++++++++++
 1 file changed, 27 insertions(+)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index a3ec11ff1a..8edc87bc2f 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -362,6 +362,29 @@ public:
     }
   }
 
+  void RewriteMacroIdentifier(const Token &MacroNameTok) {
+    std::string macroName = MacroNameTok.getIdentifierInfo()->getName();
+
+    // TODO: LUT just for macro names, to improve performance.
+    const auto found = CUDA_RENAMES_MAP().find(macroName);
+    if (found == CUDA_RENAMES_MAP().end()) {
+      // Not a CUDA macro. Moving on...
+      return;
+    }
+
+    SourceLocation sl = MacroNameTok.getLocation();
+    if (found->second.unsupported) {
+      // We know about it, but it isn't supported. Warn the user.
+      printHipifyMessage(*_sm, sl, "Unsupported CUDA macro: " + macroName);
+      return;
+    }
+
+    StringRef repName = found->second.hipName;
+    Replacement Rep(*_sm, sl, macroName.size(), repName);
+    FullSourceLoc fullSL(sl, *_sm);
+    insertReplacement(Rep, fullSL);
+  }
+
   virtual void MacroExpands(const Token &MacroNameTok,
                             const MacroDefinition &MD, SourceRange Range,
                             const MacroArgs *Args) override {
@@ -370,6 +393,9 @@ public:
       return; // Macros in headers are not our concern.
     }
 
+    // Is the macro itself a CUDA identifier? If so, rewrite it
+    RewriteMacroIdentifier(MacroNameTok);
+
     // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics
     // remain unchanged.
 #if LLVM_VERSION_MAJOR > 4
@@ -378,6 +404,7 @@ public:
     #define GET_NUM_ARGS() getNumArgs()
 #endif
 
+    // If it's a macro with arguments, rewrite all the arguments as hip, too.
     for (unsigned int i = 0; Args && i < MD.getMacroInfo()->GET_NUM_ARGS(); i++) {
       std::vector<Token> toks;
       // Code below is a kind of stolen from 'MacroArgs::getPreExpArgument'

From 59713c2459d171b8a7838746880ac207aa1a11a0 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 18 Oct 2017 15:16:02 +0100
Subject: [PATCH 05/28] 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.


[ROCm/hip commit: fd911e18395c21e91c4b0fe46a61e7dbeb4be774]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 157 ++++++---------------
 1 file changed, 42 insertions(+), 115 deletions(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index 8edc87bc2f..882b90f4d1 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -323,6 +323,45 @@ public:
     insertReplacement(Rep, fullSL);
   }
 
+  /**
+   * Look at, and consider altering, a given token.
+   *
+   * If it's not a CUDA identifier, nothing happens.
+   * If it's an unsupported CUDA identifier, a warning is emitted.
+   * Otherwise, the source file is updated with the corresponding hipification.
+   */
+  void RewriteToken(Token t) {
+    // String literals containing CUDA references need fixing...
+    if (t.is(tok::string_literal)) {
+      StringRef s(t.getLiteralData(), t.getLength());
+      processString(unquoteStr(s), *_sm, t.getLocation());
+      return;
+    } else if (!t.isAnyIdentifier()) {
+      // If it's neither a string nor an identifier, we don't care.
+      return;
+    }
+
+    StringRef name = t.getIdentifierInfo()->getName();
+    const auto found = CUDA_RENAMES_MAP().find(name);
+    if (found == CUDA_RENAMES_MAP().end()) {
+      // So it's an identifier, but not CUDA? Boring.
+      return;
+    }
+    updateCounters(found->second, name.str());
+
+    SourceLocation sl = t.getLocation();
+    if (found->second.unsupported) {
+      // An unsupported identifier? Curses! Warn the user.
+      printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier used: " + name.str());
+      return;
+    }
+
+    StringRef repName = found->second.hipName;
+    Replacement Rep(*_sm, sl, name.size(), repName);
+    FullSourceLoc fullSL(sl, *_sm);
+    insertReplacement(Rep, fullSL);
+  }
+
   virtual void MacroDefined(const Token &MacroNameTok,
                             const MacroDirective *MD) override {
     if (!_sm->isWrittenInMainFile(MD->getLocation()) ||
@@ -331,60 +370,10 @@ public:
     }
 
     for (auto T : MD->getMacroInfo()->tokens()) {
-      // We're looking for CUDA identifiers in the macro definition to rewrite...
-      if (!T.isAnyIdentifier()) {
-        continue;
-      }
-
-      StringRef name = T.getIdentifierInfo()->getName();
-      const auto found = CUDA_RENAMES_MAP().find(name);
-      if (found == CUDA_RENAMES_MAP().end()) {
-        // So it's an identifier that isn't CUDA? Boring.
-        continue;
-      }
-
-      updateCounters(found->second, name.str());
-      SourceLocation sl = T.getLocation();
-      if (found->second.unsupported) {
-        // An unsupported identifier? Curses! Warn the user.
-        printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier used: " + name.str());
-        continue;
-      }
-
-      StringRef repName = found->second.hipName;
-      DEBUG(dbgs() << "Identifier " << name << " found in definition of macro "
-                   << MacroNameTok.getIdentifierInfo()->getName() << "\n"
-                   << "will be replaced with: " << repName << "\n"
-                   << "SourceLocation: " << sl.printToString(*_sm) << "\n");
-      Replacement Rep(*_sm, sl, name.size(), repName);
-      FullSourceLoc fullSL(sl, *_sm);
-      insertReplacement(Rep, fullSL);
+      RewriteToken(T);
     }
   }
 
-  void RewriteMacroIdentifier(const Token &MacroNameTok) {
-    std::string macroName = MacroNameTok.getIdentifierInfo()->getName();
-
-    // TODO: LUT just for macro names, to improve performance.
-    const auto found = CUDA_RENAMES_MAP().find(macroName);
-    if (found == CUDA_RENAMES_MAP().end()) {
-      // Not a CUDA macro. Moving on...
-      return;
-    }
-
-    SourceLocation sl = MacroNameTok.getLocation();
-    if (found->second.unsupported) {
-      // We know about it, but it isn't supported. Warn the user.
-      printHipifyMessage(*_sm, sl, "Unsupported CUDA macro: " + macroName);
-      return;
-    }
-
-    StringRef repName = found->second.hipName;
-    Replacement Rep(*_sm, sl, macroName.size(), repName);
-    FullSourceLoc fullSL(sl, *_sm);
-    insertReplacement(Rep, fullSL);
-  }
-
   virtual void MacroExpands(const Token &MacroNameTok,
                             const MacroDefinition &MD, SourceRange Range,
                             const MacroArgs *Args) override {
@@ -394,7 +383,7 @@ public:
     }
 
     // Is the macro itself a CUDA identifier? If so, rewrite it
-    RewriteMacroIdentifier(MacroNameTok);
+    RewriteToken(MacroNameTok);
 
     // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics
     // remain unchanged.
@@ -426,69 +415,7 @@ public:
       // end of stolen code
 
       for (auto tok : toks) {
-        if (tok.isAnyIdentifier()) {
-          StringRef name = tok.getIdentifierInfo()->getName();
-          SourceLocation sl = tok.getLocation();
-
-          const auto found = CUDA_RENAMES_MAP().find(name);
-          if (found == CUDA_RENAMES_MAP().end()) {
-            // It's not a CUDA identifier. We have nothing to do.
-            continue;
-          }
-
-          updateCounters(found->second, name.str());
-          if (found->second.unsupported) {
-            // We know about it, but it isn't supported. Warn the user.
-            printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier: " + name.str());
-            continue;
-          }
-
-          StringRef repName = found->second.hipName;
-          DEBUG(dbgs() << "Identifier " << name
-                       << " found as an actual argument in expansion of macro "
-                       << MacroNameTok.getIdentifierInfo()->getName() << "\n"
-                       << "will be replaced with: " << repName << "\n");
-          size_t length = name.size();
-          if (_sm->isMacroBodyExpansion(sl)) {
-            LangOptions DefaultLangOptions;
-            SourceLocation sl_macro = _sm->getExpansionLoc(sl);
-            SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions);
-            length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro);
-            sl = sl_macro;
-          }
-          Replacement Rep(*_sm, sl, length, repName);
-          FullSourceLoc fullSL(sl, *_sm);
-          insertReplacement(Rep, fullSL);
-        } else if (tok.isLiteral()) {
-          SourceLocation sl = tok.getLocation();
-          if (_sm->isMacroBodyExpansion(sl)) {
-            LangOptions DefaultLangOptions;
-            SourceLocation sl_macro = _sm->getExpansionLoc(sl);
-            SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions);
-            size_t length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro);
-            StringRef name = StringRef(_sm->getCharacterData(sl_macro), length);
-
-            const auto found = CUDA_RENAMES_MAP().find(name);
-            if (found == CUDA_RENAMES_MAP().end()) {
-              continue; // Not CUDA, we don't care.
-            }
-
-            updateCounters(found->second, name.str());
-            if (found->second.unsupported) {
-              printHipifyMessage(*_sm, sl, "Unsupported CUDA identifier: " + name.str());
-              continue;
-            }
-
-            StringRef repName = found->second.hipName;
-            sl = sl_macro;
-            Replacement Rep(*_sm, sl, length, repName);
-            FullSourceLoc fullSL(sl, *_sm);
-            insertReplacement(Rep, fullSL);
-          } else if (tok.is(tok::string_literal)) {
-            StringRef s(tok.getLiteralData(), tok.getLength());
-            processString(unquoteStr(s), *_sm, tok.getLocation());
-          }
-        }
+        RewriteToken(tok);
       }
     }
   }

From d0acfd5bdeac3d4dc9e6362735d02f4b1ab89cd5 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Fri, 20 Oct 2017 12:46:39 +0100
Subject: [PATCH 06/28] 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: eff86d975bedea76a60ef94855da6a98e39837f7]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 65 ++++------------------
 1 file changed, 11 insertions(+), 54 deletions(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index 882b90f4d1..d5fd29688f 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -435,28 +435,6 @@ private:
 
 class Cuda2HipCallback : public MatchFinder::MatchCallback, public Cuda2Hip {
 private:
-  void convertKernelDecl(const FunctionDecl *kernelDecl, const MatchFinder::MatchResult &Result) {
-    SourceManager *SM = Result.SourceManager;
-    LangOptions DefaultLangOptions;
-    SmallString<40> XStr;
-    raw_svector_ostream OS(XStr);
-    SourceLocation sl = kernelDecl->getNameInfo().getEndLoc();
-    SourceLocation kernelArgListStart = Lexer::findLocationAfterToken(sl, tok::l_paren, *SM, DefaultLangOptions, true);
-    DEBUG(dbgs() << kernelArgListStart.printToString(*SM));
-    if (kernelDecl->getNumParams() > 0) {
-      const ParmVarDecl *pvdFirst = kernelDecl->getParamDecl(0);
-      const ParmVarDecl *pvdLast =  kernelDecl->getParamDecl(kernelDecl->getNumParams() - 1);
-      SourceLocation kernelArgListStart(pvdFirst->getLocStart());
-      SourceLocation kernelArgListEnd(pvdLast->getLocEnd());
-      SourceLocation stop = Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions);
-      size_t repLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart);
-      OS << StringRef(SM->getCharacterData(kernelArgListStart), repLength);
-      Replacement Rep0(*(Result.SourceManager), kernelArgListStart, repLength, OS.str());
-      FullSourceLoc fullSL(sl, *(Result.SourceManager));
-      insertReplacement(Rep0, fullSL);
-    }
-  }
-
   bool cudaCall(const MatchFinder::MatchResult &Result) {
     const CallExpr *call = Result.Nodes.getNodeAs<CallExpr>("cudaCall");
     if (!call) {
@@ -498,30 +476,19 @@ private:
     if (const CUDAKernelCallExpr *launchKernel = Result.Nodes.getNodeAs<CUDAKernelCallExpr>(refName)) {
       SmallString<40> XStr;
       raw_svector_ostream OS(XStr);
-      StringRef calleeName;
-      const FunctionDecl *kernelDecl = launchKernel->getDirectCallee();
-      if (kernelDecl) {
-        calleeName = kernelDecl->getName();
-        convertKernelDecl(kernelDecl, Result);
-      } else {
-        const Expr *e = launchKernel->getCallee();
-        if (const UnresolvedLookupExpr *ule =
-          dyn_cast<UnresolvedLookupExpr>(e)) {
-          calleeName = ule->getName().getAsIdentifierInfo()->getName();
-          owner->addMatcher(functionTemplateDecl(hasName(calleeName))
-            .bind("unresolvedTemplateName"), this);
-        }
-      }
-      XStr.clear();
-      if (calleeName.find(',') != StringRef::npos) {
-        SmallString<128> tmpData;
-        calleeName = Twine("(" + calleeName + ")").toStringRef(tmpData);
-      }
-      OS << "hipLaunchKernelGGL(" << calleeName << ",";
+
+      LangOptions DefaultLangOptions;
+      SourceManager *SM = Result.SourceManager;
+
+      const Expr *e = launchKernel->getCallee();
+
+      // Grab the characters for the callee expression and dump them into hipLaunchKernelGGL's
+      // first argument.
+      StringRef exprSource = Lexer::getSourceText(CharSourceRange::getTokenRange(e->getSourceRange()), *SM, LangOptions(), 0);
+
+      OS << "hipLaunchKernelGGL(" << exprSource << ",";
       const CallExpr *config = launchKernel->getConfig();
       DEBUG(dbgs() << "Kernel config arguments:" << "\n");
-      SourceManager *SM = Result.SourceManager;
-      LangOptions DefaultLangOptions;
       for (unsigned argno = 0; argno < config->getNumArgs(); argno++) {
         const Expr *arg = config->getArg(argno);
         if (!isa<CXXDefaultArgExpr>(arg)) {
@@ -724,15 +691,6 @@ private:
     return false;
   }
 
-  bool unresolvedTemplateName(const MatchFinder::MatchResult &Result) {
-    if (const FunctionTemplateDecl *templateDecl = Result.Nodes.getNodeAs<FunctionTemplateDecl>("unresolvedTemplateName")) {
-      FunctionDecl *kernelDecl = templateDecl->getTemplatedDecl();
-      convertKernelDecl(kernelDecl, Result);
-      return true;
-    }
-    return false;
-  }
-
   bool stringLiteral(const MatchFinder::MatchResult &Result) {
     if (const clang::StringLiteral *sLiteral = Result.Nodes.getNodeAs<clang::StringLiteral>("stringLiteral")) {
       if (sLiteral->getCharByteWidth() == 1) {
@@ -759,7 +717,6 @@ public:
     if (cudaLaunchKernel(Result)) return;
     if (cudaSharedIncompleteArrayVar(Result)) return;
     if (stringLiteral(Result)) return;
-    if (unresolvedTemplateName(Result)) return;
   }
 
 private:

From c15f6bdf5caa9ab4ccb53866ec81218f1d067a2f Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Fri, 20 Oct 2017 13:10:31 +0100
Subject: [PATCH 07/28] 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: 094b2b9b0503c1e2935863a1d596d1045b71e7e4]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 126 +++++++++++++--------
 projects/hip/tests/hipify-clang/axpy.cu    |  35 +++++-
 2 files changed, 111 insertions(+), 50 deletions(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index d5fd29688f..fd23344f9f 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -471,6 +471,49 @@ private:
     return true;
   }
 
+  SourceRange getReadRange(clang::SourceManager &SM, const SourceRange &exprRange) {
+    SourceLocation begin = exprRange.getBegin();
+    SourceLocation end = exprRange.getEnd();
+
+    bool beginSafe = !SM.isMacroBodyExpansion(begin) || Lexer::isAtStartOfMacroExpansion(begin, SM, LangOptions{});
+    bool endSafe = !SM.isMacroBodyExpansion(end) || Lexer::isAtEndOfMacroExpansion(end, SM, LangOptions{});
+
+    if (beginSafe && endSafe) {
+      return {SM.getFileLoc(begin), SM.getFileLoc(end)};
+    } else {
+      return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)};
+    }
+  }
+
+  SourceRange getWriteRange(clang::SourceManager &SM, const SourceRange &exprRange) {
+    SourceLocation begin = exprRange.getBegin();
+    SourceLocation end = exprRange.getEnd();
+
+    // If the range is contained within a macro, update the macro definition.
+    // Otherwise, use the file location and hope for the best.
+    if (!SM.isMacroBodyExpansion(begin) || !SM.isMacroBodyExpansion(end)) {
+      return {SM.getFileLoc(begin), SM.getFileLoc(end)};
+    }
+
+    return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)};
+  }
+
+  StringRef readSourceText(clang::SourceManager& SM, const SourceRange& exprRange) {
+    return Lexer::getSourceText(CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM, LangOptions(), nullptr);
+  }
+
+  /**
+   * Get a string representation of the expression `arg`, unless it's a defaulting function
+   * call argument, in which case get a 0. Used for building argument lists to kernel calls.
+   */
+  std::string stringifyZeroDefaultedArg(SourceManager& SM, const Expr* arg) {
+    if (isa<CXXDefaultArgExpr>(arg)) {
+      return "0";
+    } else {
+      return readSourceText(SM, arg->getSourceRange());
+    }
+  }
+
   bool cudaLaunchKernel(const MatchFinder::MatchResult &Result) {
     StringRef refName = "cudaLaunchKernel";
     if (const CUDAKernelCallExpr *launchKernel = Result.Nodes.getNodeAs<CUDAKernelCallExpr>(refName)) {
@@ -480,59 +523,46 @@ private:
       LangOptions DefaultLangOptions;
       SourceManager *SM = Result.SourceManager;
 
-      const Expr *e = launchKernel->getCallee();
+      const Expr& calleeExpr = *(launchKernel->getCallee());
+      OS << "hipLaunchKernelGGL(" << readSourceText(*SM, calleeExpr.getSourceRange()) << ", ";
 
-      // Grab the characters for the callee expression and dump them into hipLaunchKernelGGL's
-      // first argument.
-      StringRef exprSource = Lexer::getSourceText(CharSourceRange::getTokenRange(e->getSourceRange()), *SM, LangOptions(), 0);
+      // Next up are the four kernel configuration parameters, the last two of which are optional and default to zero.
+      const CallExpr& config = *(launchKernel->getConfig());
 
-      OS << "hipLaunchKernelGGL(" << exprSource << ",";
-      const CallExpr *config = launchKernel->getConfig();
-      DEBUG(dbgs() << "Kernel config arguments:" << "\n");
-      for (unsigned argno = 0; argno < config->getNumArgs(); argno++) {
-        const Expr *arg = config->getArg(argno);
-        if (!isa<CXXDefaultArgExpr>(arg)) {
-          const ParmVarDecl *pvd = config->getDirectCallee()->getParamDecl(argno);
-          SourceLocation sl(arg->getLocStart());
-          SourceLocation el(arg->getLocEnd());
-          SourceLocation stop = Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions);
-          StringRef outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl));
-          DEBUG(dbgs() << "args[ " << argno << "]" << outs << " <" << pvd->getType().getAsString() << ">\n");
-          if (pvd->getType().getAsString().compare("dim3") == 0) {
-            OS << " dim3(" << outs << "),";
-          } else {
-            OS << " " << outs << ",";
-          }
-        } else {
-          OS << " 0,";
-        }
+      // Copy the two dimensional arguments verbatim.
+      OS << "dim3(" << readSourceText(*SM, config.getArg(0)->getSourceRange()) << "), ";
+      OS << "dim3(" << readSourceText(*SM, config.getArg(1)->getSourceRange()) << "), ";
+
+      // The stream/memory arguments default to zero if omitted.
+      OS << stringifyZeroDefaultedArg(*SM, config.getArg(2)) << ", ";
+      OS << stringifyZeroDefaultedArg(*SM, config.getArg(3));
+
+      // If there are ordinary arguments to the kernel, just copy them verbatim into our new call.
+      int numArgs = launchKernel->getNumArgs();
+      if (numArgs > 0) {
+        OS << ", ";
+
+        // Start of the first argument.
+        SourceLocation argStart = launchKernel->getArg(0)->getLocStart();
+
+        // End of the last argument.
+        SourceLocation argEnd = launchKernel->getArg(numArgs - 1)->getLocEnd();
+
+        OS << readSourceText(*SM, {argStart, argEnd});
       }
-      for (unsigned argno = 0; argno < launchKernel->getNumArgs(); argno++) {
-        const Expr *arg = launchKernel->getArg(argno);
-        SourceLocation sl(arg->getLocStart());
-        if (SM->isMacroBodyExpansion(sl)) {
-          sl = SM->getExpansionLoc(sl);
-        } else if (SM->isMacroArgExpansion(sl)) {
-          sl = SM->getImmediateSpellingLoc(sl);
-        }
-        SourceLocation el(arg->getLocEnd());
-        if (SM->isMacroBodyExpansion(el)) {
-          el = SM->getExpansionLoc(el);
-        } else if (SM->isMacroArgExpansion(el)) {
-          el = SM->getImmediateSpellingLoc(el);
-        }
-        SourceLocation stop = Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions);
-        std::string outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl));
-        DEBUG(dbgs() << outs << "\n");
-        OS << " " << outs << ",";
-      }
-      XStr.pop_back();
+
       OS << ")";
+
+      SourceRange replacementRange = getWriteRange(*SM, {launchKernel->getLocStart(), launchKernel->getLocEnd()});
+      SourceLocation launchStart = replacementRange.getBegin();
+      SourceLocation launchEnd = replacementRange.getEnd();
+
       size_t length = SM->getCharacterData(Lexer::getLocForEndOfToken(
-                        launchKernel->getLocEnd(), 0, *SM, DefaultLangOptions)) -
-                        SM->getCharacterData(launchKernel->getLocStart());
-      Replacement Rep(*SM, launchKernel->getLocStart(), length, OS.str());
-      FullSourceLoc fullSL(launchKernel->getLocStart(), *SM);
+                        launchEnd, 0, *SM, DefaultLangOptions)) -
+                        SM->getCharacterData(launchStart);
+
+      Replacement Rep(*SM, launchStart, length, OS.str());
+      FullSourceLoc fullSL(launchStart, *SM);
       insertReplacement(Rep, fullSL);
       hipCounter counter = {"hipLaunchKernelGGL", CONV_KERN, API_RUNTIME};
       updateCounters(counter, refName.str());
diff --git a/projects/hip/tests/hipify-clang/axpy.cu b/projects/hip/tests/hipify-clang/axpy.cu
index 8c6b0e0d8d..2fd62ac344 100644
--- a/projects/hip/tests/hipify-clang/axpy.cu
+++ b/projects/hip/tests/hipify-clang/axpy.cu
@@ -2,11 +2,23 @@
 
 #include <iostream>
 
-__global__ void axpy(float a, float* x, float* y) {
+
+#define TOKEN_PASTE(X, Y) X ## Y
+#define ARG_LIST_AS_MACRO a, device_x, device_y
+#define KERNEL_CALL_AS_MACRO axpy<float><<<1, kDataLen>>>
+#define KERNEL_NAME_MACRO axpy<float>
+
+// CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y)
+#define COMPLETE_LAUNCH axpy<<<1, kDataLen>>>(a, device_x, device_y)
+
+
+template<typename T>
+__global__ void axpy(T a, T *x, T *y) {
   // CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x];
   y[threadIdx.x] = a * x[threadIdx.x];
 }
 
+
 int main(int argc, char* argv[]) {
   const int kDataLen = 4;
 
@@ -27,10 +39,29 @@ int main(int argc, char* argv[]) {
   // CHECK: hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice);
   cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice);
 
-  // Launch the kernel.
+  // Launch the kernel in numerous different strange ways to exercise the prerocessor.
   // CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
   axpy<<<1, kDataLen>>>(a, device_x, device_y);
 
+  // CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
+  axpy<float><<<1, kDataLen>>>(a, device_x, device_y);
+
+  // CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y);
+  axpy<float><<<1, kDataLen>>>(a, TOKEN_PASTE(device, _x), device_y);
+
+  // CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
+  axpy<float><<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
+
+  // CHECK: hipLaunchKernelGGL(KERNEL_NAME_MACRO, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
+  KERNEL_NAME_MACRO<<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
+
+  // CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
+  KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO);
+
+  // CHECK: COMPLETE_LAUNCH;
+  COMPLETE_LAUNCH;
+
+
   // Copy output data to host.
   // CHECK: hipDeviceSynchronize();
   cudaDeviceSynchronize();

From 7947e406c7398dacfe0e76728ed43e890a202f73 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Tue, 24 Oct 2017 20:56:55 +0100
Subject: [PATCH 08/28] Make BUILD_HIPIFY_CLANG a cmake option

Instead of deciding whether to build hipify-clang based on
the presence of an LLVM path on the command line, have an
explicit option.

Do we want this default-on or default-off? I've defaulted it to
on for now, but maybe we want the opposite?


[ROCm/hip commit: a4ecd4eb3124d5cc7e414178b5a56709d8cea547]
---
 projects/hip/CMakeLists.txt              | 9 ++++++++-
 projects/hip/hipify-clang/CMakeLists.txt | 6 +-----
 2 files changed, 9 insertions(+), 6 deletions(-)

diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index 4c89c93668..1304e9f7d7 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -1,6 +1,11 @@
 cmake_minimum_required(VERSION 2.8.3)
 project(hip)
 
+#############################
+# Options
+#############################
+option(BUILD_HIPIFY_CLANG "Enable building the CUDA->HIP converter" OFF)
+
 #############################
 # Setup config generation
 #############################
@@ -142,7 +147,9 @@ add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER)
 # Build steps
 #############################
 # Build clang hipify if enabled
-add_subdirectory(hipify-clang)
+if (BUILD_HIPIFY_CLANG)
+    add_subdirectory(hipify-clang)
+endif()
 
 # Build hip_hcc if platform is hcc
 if(HIP_PLATFORM STREQUAL "hcc")
diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt
index cb8354157b..7b43f2beb7 100644
--- a/projects/hip/hipify-clang/CMakeLists.txt
+++ b/projects/hip/hipify-clang/CMakeLists.txt
@@ -1,8 +1,6 @@
 cmake_minimum_required(VERSION 2.8.8)
 project(hipify-clang)
 
-set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "")
-
 if (HIPIFY_CLANG_LLVM_DIR)
     find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED NO_DEFAULT_PATH)
 else()
@@ -10,7 +8,7 @@ else()
     return()
 endif()
 
-option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON)
+option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" OFF)
 
 # Disable the tests if `lit` or `FileCheck` is not installed.
 find_program(LIT_COMMAND lit)
@@ -102,5 +100,3 @@ if (HIPIFY_CLANG_TESTS)
     add_dependencies(test-hipify-clang test-hipify)
     set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests")
 endif()
-
-set(BUILD_HIPIFY_CLANG 1 CACHE INTERNAL "")

From b61c4a241fd860ea9da69a016dc7ca9b82261f9c Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Tue, 24 Oct 2017 20:59:30 +0100
Subject: [PATCH 09/28] Use add_dependencies to avoid duplication of
 pkg_hip_base

[ROCm/hip commit: 56b422204397c1858b17ca396d80be5c4ea25f93]
---
 projects/hip/CMakeLists.txt | 30 +++++++++++++-----------------
 1 file changed, 13 insertions(+), 17 deletions(-)

diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index 1304e9f7d7..dd8d604dd2 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -294,23 +294,19 @@ set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_base)
 configure_file(packaging/hip_base.txt ${BUILD_DIR}/CMakeLists.txt @ONLY)
 configure_file(packaging/hip_base.postinst ${BUILD_DIR}/postinst @ONLY)
 configure_file(packaging/hip_base.prerm ${BUILD_DIR}/prerm @ONLY)
-if(NOT BUILD_HIPIFY_CLANG)
-    add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} .
-        COMMAND rm -rf *.deb *.rpm *.tar.gz
-        COMMAND make package
-        COMMAND cp *.deb ${PROJECT_BINARY_DIR}
-        COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
-        COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
-        WORKING_DIRECTORY ${BUILD_DIR})
-else()
-    add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} .
-        COMMAND rm -rf *.deb *.rpm *.tar.gz
-        COMMAND make package
-        COMMAND cp *.deb ${PROJECT_BINARY_DIR}
-        COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
-        COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
-        WORKING_DIRECTORY ${BUILD_DIR}
-        DEPENDS hipify-clang)
+
+add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} .
+    COMMAND rm -rf *.deb *.rpm *.tar.gz
+    COMMAND make package
+    COMMAND cp *.deb ${PROJECT_BINARY_DIR}
+    COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
+    COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
+    WORKING_DIRECTORY ${BUILD_DIR}
+)
+
+# Packaging needs to wait for hipify-clang to build if it's enabled...
+if (BUILD_HIPIFY_CLANG)
+    add_dependencies(pkg_hip_base hipify-clang)
 endif()
 
 # Package: hip_hcc

From d92f3c8d764629401d361ef51724ee520a52ee41 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Tue, 24 Oct 2017 21:30:34 +0100
Subject: [PATCH 10/28] Don't attempt to find test dependencies if tests are
 disabled

And while we're at it, introduce a handy program-finder macro


[ROCm/hip commit: 921ff4c8a3dd8994bdb31de8af75bc90c121791c]
---
 projects/hip/hipify-clang/CMakeLists.txt | 20 +++++++++++---------
 1 file changed, 11 insertions(+), 9 deletions(-)

diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt
index 7b43f2beb7..13abd287c8 100644
--- a/projects/hip/hipify-clang/CMakeLists.txt
+++ b/projects/hip/hipify-clang/CMakeLists.txt
@@ -10,15 +10,6 @@ endif()
 
 option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" OFF)
 
-# Disable the tests if `lit` or `FileCheck` is not installed.
-find_program(LIT_COMMAND lit)
-find_program(FILECHECK_COMMAND FileCheck)
-find_program(SOCAT_COMMAND socat)
-if (NOT LIT_COMMAND OR NOT FILECHECK_COMMAND OR NOT SOCAT_COMMAND)
-    set(HIPIFY_CLANG_TESTS OFF CACHE INTERNAL "")
-    message(STATUS "hipify-clang's tests are not being built because `lit`,`FileCheck` or `socat` could not be found.")
-endif()
-
 list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR})
 include(AddLLVM)
 
@@ -79,6 +70,17 @@ install(TARGETS hipify-clang DESTINATION bin)
 if (HIPIFY_CLANG_TESTS)
     find_package(PythonInterp 2.7 REQUIRED EXACT)
 
+    function (require_program PROGRAM_NAME)
+        find_program(FOUND_PROGRAM ${PROGRAM_NAME})
+        if (NOT FOUND_PROGRAM)
+            message(FATAL_ERROR "Can't find ${PROGRAM_NAME}. Either set HIPIFY_CLANG_TESTS to OFF to disable hipify tests, or install the missing program.")
+        endif()
+    endfunction()
+
+    require_program(lit)
+    require_program(FileCheck)
+    require_program(socat)
+
     # Populates CUDA_TOOLKIT_ROOT_DIR, which is then applied to the lit config to give the
     # value of --cuda-path for the test runs.
     find_package(CUDA REQUIRED)

From cc1bc495a0ef2a7372f940f5ed2a8e3b63282d93 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Tue, 24 Oct 2017 21:31:24 +0100
Subject: [PATCH 11/28] We no longer rely on HIPIFY_CLANG_LLVM_DIR to disable
 hipify-clang

Since there's now an option for toggling hipify-clang, omitting the
path is no longer something we need to check for. We'll still
abort if LLVM isn't found, due to `REQUIRED`.


[ROCm/hip commit: c60c8d417ee19b901df125778ca80b1950e00f28]
---
 projects/hip/hipify-clang/CMakeLists.txt | 7 +------
 1 file changed, 1 insertion(+), 6 deletions(-)

diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt
index 13abd287c8..c8d205f0f4 100644
--- a/projects/hip/hipify-clang/CMakeLists.txt
+++ b/projects/hip/hipify-clang/CMakeLists.txt
@@ -1,12 +1,7 @@
 cmake_minimum_required(VERSION 2.8.8)
 project(hipify-clang)
 
-if (HIPIFY_CLANG_LLVM_DIR)
-    find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED NO_DEFAULT_PATH)
-else()
-    message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or higher using HIPIFY_CLANG_LLVM_DIR")
-    return()
-endif()
+find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED NO_DEFAULT_PATH)
 
 option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" OFF)
 

From 1461c69757dade16bcfd4d45160d1c47a8966b1c Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Tue, 24 Oct 2017 21:32:51 +0100
Subject: [PATCH 12/28] Move the "LLVM found" print adjacent to the
 find_package call

Very surprising that LLVM's finder module doesn't print this
itself like _literally every other finder module_. Blarg.


[ROCm/hip commit: 92c90a70685f17ef21eabe528251a12c66b9272b]
---
 projects/hip/hipify-clang/CMakeLists.txt | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt
index c8d205f0f4..e9709a6f6e 100644
--- a/projects/hip/hipify-clang/CMakeLists.txt
+++ b/projects/hip/hipify-clang/CMakeLists.txt
@@ -2,14 +2,13 @@ cmake_minimum_required(VERSION 2.8.8)
 project(hipify-clang)
 
 find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED NO_DEFAULT_PATH)
+message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}")
 
 option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" OFF)
 
 list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR})
 include(AddLLVM)
 
-message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}")
-
 include_directories(${LLVM_INCLUDE_DIRS})
 link_directories(${LLVM_LIBRARY_DIRS})
 add_definitions(${LLVM_DEFINITIONS})

From fd656a0fbbbe0e65a31c305dd0bc00c0cb1e8ac7 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Tue, 24 Oct 2017 21:36:51 +0100
Subject: [PATCH 13/28] Use cmake's builtin mechanism for handling library
 locations

See [the documentation](https://cmake.org/cmake/help/v3.0/command/find_package.html)
for exactly how the search procedure works. If you want to use an
LLVM from a specific location, use CMAKE_PREFIX_PATH as normal.

No longer do we have a nonstandard HIPIFY_CLANG_LLVM_DIR variable
for people to learn about.


[ROCm/hip commit: 8fefc6a2b789739dfad064cd32d1e57bf9ef03ff]
---
 projects/hip/CMakeLists.txt              | 7 -------
 projects/hip/hipify-clang/CMakeLists.txt | 2 +-
 2 files changed, 1 insertion(+), 8 deletions(-)

diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index dd8d604dd2..1c8f640afb 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -125,13 +125,6 @@ else()
     message(FATAL_ERROR "Don't know where to install HIP. Please specify absolute path using -DCMAKE_INSTALL_PREFIX")
 endif()
 
-# Check if we need to build hipify-clang
-if(NOT DEFINED HIPIFY_CLANG_LLVM_DIR)
-    if(DEFINED ENV{HIPIFY_CLANG_LLVM_DIR})
-        set(HIPIFY_CLANG_LLVM_DIR $ENV{HIPIFY_CLANG_LLVM_DIR})
-    endif()
-endif()
-
 # Check if we need to enable ATP marker
 if(NOT DEFINED COMPILE_HIP_ATP_MARKER)
     if(NOT DEFINED ENV{COMPILE_HIP_ATP_MARKER})
diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt
index e9709a6f6e..0dc1affec5 100644
--- a/projects/hip/hipify-clang/CMakeLists.txt
+++ b/projects/hip/hipify-clang/CMakeLists.txt
@@ -1,7 +1,7 @@
 cmake_minimum_required(VERSION 2.8.8)
 project(hipify-clang)
 
-find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED NO_DEFAULT_PATH)
+find_package(LLVM REQUIRED)
 message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}")
 
 option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" OFF)

From e3338fe5a4646f93184d740b656fa5eeca6c8f7d Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Tue, 24 Oct 2017 21:39:38 +0100
Subject: [PATCH 14/28] hipify does not add the `hipLaunchParm` option any more

This was removed a while ago - seems like it uses a different
variant of the launch kernel function now, so this is redundant.


[ROCm/hip commit: b412802c66b25e6dbc83cfc222c7f23f6f6c1861]
---
 projects/hip/hipify-clang/README.md | 1 -
 1 file changed, 1 deletion(-)

diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md
index 20456f3bff..bb2ffe55b5 100644
--- a/projects/hip/hipify-clang/README.md
+++ b/projects/hip/hipify-clang/README.md
@@ -12,7 +12,6 @@
 ## Using hipify-clang
 
 `hipify-clang` is a clang-based tool which can automate the translation of CUDA source code into portable HIP C++.
-The tool can automatically add extra HIP arguments (notably the "hipLaunchParm" required at the beginning of every HIP kernel call).
 `hipify-clang` has some additional dependencies explained below and can be built as a separate make step. The instructions below are specifically for **Ubuntu 14.04** and **Ubuntu 16.04**.
 
 ### Build and install

From 7a0eedabbebdba5d33fb409518050a4526e75625 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 25 Oct 2017 18:45:03 +0100
Subject: [PATCH 15/28] Update hipify-clang readme for simplified build process

[ROCm/hip commit: 82d05ee6f4236dea2c1bc1ac889f8b1f5b1f5fd4]
---
 projects/hip/hipify-clang/README.md | 88 +++++++++++++++--------------
 1 file changed, 45 insertions(+), 43 deletions(-)

diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md
index bb2ffe55b5..563440e709 100644
--- a/projects/hip/hipify-clang/README.md
+++ b/projects/hip/hipify-clang/README.md
@@ -1,3 +1,7 @@
+# hipify-clang
+
+`hipify-clang` is a clang-based tool to automatically translate CUDA source code into portable HIP C++.
+
 ## Table of Contents
 
 <!-- toc -->
@@ -9,67 +13,65 @@
 
 <!-- tocstop -->
 
-## Using hipify-clang
+## Build and install
 
-`hipify-clang` is a clang-based tool which can automate the translation of CUDA source code into portable HIP C++.
-`hipify-clang` has some additional dependencies explained below and can be built as a separate make step. The instructions below are specifically for **Ubuntu 14.04** and **Ubuntu 16.04**.
+### Dependencies
 
-### Build and install
+`hipify-clang` requires clang+llvm of at least version 3.8.
 
-- Download and unpack clang+llvm 3.8 binary package preqrequisite.
+In most cases, you can get a suitable version of clang+llvm with your package manager.
+
+Failing that, you can [download a release archive](http://releases.llvm.org/), extract it somewhere, and set
+[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.0/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it.
+
+### Build
+
+Assuming this repository is at `./HIP`:
 
-**Ubuntu 14.04**:
 ```shell
-wget http://llvm.org/releases/3.8.0/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04.tar.xz
-tar xvfJ clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-14.04.tar.xz
-```
-**Ubuntu 16.04**:
-```shell
-wget http://llvm.org/releases/3.8.0/clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-16.04.tar.xz
-tar xvfJ clang+llvm-3.8.0-x86_64-linux-gnu-ubuntu-16.04.tar.xz
-```
+mkdir build inst
 
-- Enable build of hipify-clang and specify path to LLVM.
-
-Note HIPIFY_CLANG_LLVM_DIR must be a full absolute path to the location extracted above. Here's an example assuming we extract the clang 3.8 package into ~/HIP/clang+llvm-3.8.0/
-```shell
-cd HIP
-mkdir build
 cd build
-cmake -DHIPIFY_CLANG_LLVM_DIR=~/HIP/clang+llvm-3.8.0/ -DCMAKE_BUILD_TYPE=Release ..
-make
-make install
+cmake \
+ -DCMAKE_INSTALL_PREFIX=../inst \
+ -DCMAKE_BUILD_TYPE=Release \
+ -DBUILD_HIPIFY_CLANG=ON \
+ ../HIP
+
+make -j install
 ```
 
-### Running and using hipify-clang
+The binary can then be found at `./inst/bin/hipify-clang`.
 
-`hipify-clang` performs an initial compile of the CUDA source code into a "symbol tree", and thus needs access to the appropriate header files.
+### Test
 
-In the case when `hipify-clang` doesn't find cuda headers, it reports various errors about unknown keywords (e.g. '\__global\__'), API function names (e.g. 'cudaMalloc'), syntax (e.g. 'foo<<<1,n>>>(...)'), etc.
+`hipify-clang` has unit tests using LLVM [`lit`](https://llvm.org/docs/CommandGuide/lit.html)/[`FileCheck`](https://llvm.org/docs/CommandGuide/FileCheck.html).
 
-To install CUDA headers, download the "deb(network)" variant of the target installer.
+To run it:
+1. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM.
+2. Ensure `socat` is installed - your distro almost certainly has a package for this.
+3. Build with the `HIPIFY_CLANG_TESTS` option turned on.
+4. `make test-hipify`
 
-**Ubuntu 14.04**:
-```shell
-wget http://developer.download.nvidia.com/compute/cuda/repos/ubuntu1404/x86_64/cuda-repo-ubuntu1404_7.5-18_amd64.deb
-sudo dpkg -i cuda-repo-ubuntu1404_7.5-18_amd64.deb
-sudo apt-get update && sudo apt-get install cuda-minimal-build-7-5 cuda-curand-dev-7-5
-```
-**Ubuntu 16.04**:
-```shell
-wget http://archive.ubuntu.com/ubuntu/pool/multiverse/n/nvidia-cuda-toolkit/nvidia-cuda-toolkit_7.5.18-0ubuntu1_amd64.deb
-sudo dpkg -i nvidia-cuda-toolkit_7.5.18-0ubuntu1_amd64.deb
-sudo apt-get update && sudo apt-get install cuda-minimal-build-7-5 cuda-curand-dev-7-5
-```
-To set additional options like Language Selection (only "-x cuda" is supported), Preprocessor Definition (-D), Include Path (-I), etc., options delimiter "--" should be used before them, for instance:
+## Running and using hipify-clang
+
+To process a file, `hipify-clang` needs access to the same headers that would be needed to compile it with clang.
+
+For example:
 
 ```shell
-./hipify-clang -print-stats sort_kernel.cu -- -x cuda -I/srv/git/HIP/include -I/usr/local/cuda-7.5/include -DX=1
+hipify-clang square.cu -- \
+  -x cuda \
+  --cuda-path=/opt/cuda \
+  --cuda-gpu-arch=sm_30 \
+  -isystem /opt/cuda/samples/common/inc
 ```
 
-Delimiter "--" is used to separate hipify-clang options (before the delimiter) from clang options (after the delimiter). It is strongly recommended to always specify the delimiter, even if there are no clang specific options at all, in order to avoid possible errors regarding compilation database; in such case delimeter should be the last option in hipify-clang's command line.
+`hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you
+were compiling the input file. The [Clang manual for compiling CUDA](https://llvm.org/docs/CompileCudaWithLLVM.html#compiling-cuda-code)
+may be useful.
 
-Option "-x cuda" is also worth specifying in order to convert source CUDA files with extensions other than standard extensions (*.cu, *.cuh).
+For a list of `hipify-clang` options, run `hipify-clang --help`.
 
 ## Disclaimer
 

From 22854820754f6678e966261b22e6e3d55ccbd61e Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Fri, 27 Oct 2017 16:17:24 +0100
Subject: [PATCH 16/28] Describe the LLVM we found

[ROCm/hip commit: 69e67fe25a8475ebc9dbc0fe30d23769352135f8]
---
 projects/hip/hipify-clang/CMakeLists.txt | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt
index 0dc1affec5..da6eaeaa99 100644
--- a/projects/hip/hipify-clang/CMakeLists.txt
+++ b/projects/hip/hipify-clang/CMakeLists.txt
@@ -2,7 +2,10 @@ cmake_minimum_required(VERSION 2.8.8)
 project(hipify-clang)
 
 find_package(LLVM REQUIRED)
-message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}")
+message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}:")
+message(STATUS "   - CMake module path: ${LLVM_CMAKE_DIR}")
+message(STATUS "   - Include path     : ${LLVM_INCLUDE_DIRS}")
+message(STATUS "   - Binary path      : ${LLVM_TOOLS_BINARY_DIR}")
 
 option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" OFF)
 

From 199d75adc0deb284005905d1babb1c92bdebf17f Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Mon, 23 Oct 2017 17:06:04 +0100
Subject: [PATCH 17/28] Make `unsupported` actually be a bool...

[ROCm/hip commit: 2f376c9b25a8654b508613ddd57851c015f14858]
---
 projects/hip/hipify-clang/src/CUDA2HipMap.h | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.h b/projects/hip/hipify-clang/src/CUDA2HipMap.h
index 95f9054576..bd172d6eb5 100644
--- a/projects/hip/hipify-clang/src/CUDA2HipMap.h
+++ b/projects/hip/hipify-clang/src/CUDA2HipMap.h
@@ -11,10 +11,10 @@ struct hipCounter {
     llvm::StringRef hipName;
     ConvTypes countType;
     ApiTypes countApiType;
-    int unsupported;
+    bool unsupported;
 };
 
-#define HIP_UNSUPPORTED -1
+#define HIP_UNSUPPORTED true
 
 /// Macros to ignore.
 extern const std::set<llvm::StringRef> CUDA_EXCLUDES;

From 48e7403762abb0fa706fe0e8fd4f2572a13526fc Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Mon, 23 Oct 2017 17:06:24 +0100
Subject: [PATCH 18/28] Remove CUDA_EXCLUDES

An artefact from a now-defunct hack to avoid corrupting programs


[ROCm/hip commit: c6707ef33c8978158832e80677fc0b34d1e82d7a]
---
 projects/hip/hipify-clang/src/CUDA2HipMap.cpp | 2 --
 projects/hip/hipify-clang/src/CUDA2HipMap.h   | 3 ---
 2 files changed, 5 deletions(-)

diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp
index fb7e920f8c..de6ddb2d74 100644
--- a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp
+++ b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp
@@ -1,7 +1,5 @@
 #include "CUDA2HipMap.h"
 
-const std::set<llvm::StringRef> CUDA_EXCLUDES{"CHECK_CUDA_ERROR", "CUDA_SAFE_CALL"};
-
 /// Maps the names of CUDA types to the corresponding hip types.
 const std::map<llvm::StringRef, hipCounter> CUDA_TYPE_NAME_MAP{
     // Error codes and return types
diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.h b/projects/hip/hipify-clang/src/CUDA2HipMap.h
index bd172d6eb5..ecef9ce2b5 100644
--- a/projects/hip/hipify-clang/src/CUDA2HipMap.h
+++ b/projects/hip/hipify-clang/src/CUDA2HipMap.h
@@ -16,9 +16,6 @@ struct hipCounter {
 
 #define HIP_UNSUPPORTED true
 
-/// Macros to ignore.
-extern const std::set<llvm::StringRef> CUDA_EXCLUDES;
-
 /// Maps cuda header names to hip header names.
 extern const std::map<llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP;
 

From 389fa2e68e95910e8bf064138e97fbfdb9c8896c Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Mon, 23 Oct 2017 17:11:51 +0100
Subject: [PATCH 19/28] Remove unused field

[ROCm/hip commit: 0c09bdf5230a7bbda04d2802bf8cbd6b53efcda3]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index 5c422bc6ff..1c2fe5195d 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -422,7 +422,6 @@ public:
 
   void EndOfMainFile() override {}
 
-  bool SeenEnd = false;
   void setSourceManager(SourceManager *sm) { _sm = sm; }
   void setPreprocessor(Preprocessor *pp) { _pp = pp; }
   void setMatch(Cuda2HipCallback *match) { Match = match; }

From 85fd2e6f5141a81e5d49d4881ef6b2239352135c Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 25 Oct 2017 19:17:28 +0100
Subject: [PATCH 20/28] Extract LLVM compatibility code into its own
 translation unit

[ROCm/hip commit: 1bd837b4b12bf9f09209fc1510c6c86fbd243852]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp   | 59 +++++---------------
 projects/hip/hipify-clang/src/LLVMCompat.cpp | 43 ++++++++++++++
 projects/hip/hipify-clang/src/LLVMCompat.h   | 49 ++++++++++++++++
 3 files changed, 105 insertions(+), 46 deletions(-)
 create mode 100644 projects/hip/hipify-clang/src/LLVMCompat.cpp
 create mode 100644 projects/hip/hipify-clang/src/LLVMCompat.h

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index 1c2fe5195d..270dec68b3 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -54,6 +54,7 @@ THE SOFTWARE.
 
 #include "CUDA2HipMap.h"
 #include "Types.h"
+#include "LLVMCompat.h"
 
 using namespace clang;
 using namespace clang::ast_matchers;
@@ -139,7 +140,7 @@ void removePrefixIfPresent(std::string& s, std::string prefix) {
 
 class Cuda2Hip {
 public:
-  Cuda2Hip(Replacements *R, const std::string &srcFileName) :
+  Cuda2Hip(Replacements& R, const std::string &srcFileName) :
     Replace(R), mainFileName(srcFileName) {}
   uint64_t countReps[CONV_LAST] = { 0 };
   uint64_t countApiReps[API_LAST] = { 0 };
@@ -163,23 +164,17 @@ public:
   }
 
 protected:
-  Replacements *Replace;
+  Replacements& Replace;
   std::string mainFileName;
 
   virtual void insertReplacement(const Replacement &rep, const FullSourceLoc &fullSL) {
-#if LLVM_VERSION_MAJOR > 3
-    // New clang added error checking to Replacements, and *insists* that you explicitly check it.
-    llvm::Error e = Replace->add(rep);
-#else
-    // In older versions, it's literally an std::set<Replacement>
-    Replace->insert(rep);
-#endif
+    llcompat::insertReplacement(Replace, rep);
     if (PrintStats) {
       LOCs.insert(fullSL.getExpansionLineNumber());
     }
   }
   void insertHipHeaders(Cuda2Hip *owner, const SourceManager &SM) {
-    if (owner->countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && Replace->size() > 0) {
+    if (owner->countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && Replace.size() > 0) {
       std::string repName = "#include <hip/hip_runtime.h>";
       hipCounter counter = { repName, CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME };
       updateCounters(counter, repName);
@@ -266,7 +261,7 @@ class Cuda2HipCallback;
 
 class HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks, public Cuda2Hip {
 public:
-  HipifyPPCallbacks(Replacements *R, const std::string &mainFileName)
+  HipifyPPCallbacks(Replacements& R, const std::string &mainFileName)
     : Cuda2Hip(R, mainFileName) {}
 
   virtual bool handleBeginSource(CompilerInstance &CI
@@ -385,14 +380,6 @@ public:
     // Is the macro itself a CUDA identifier? If so, rewrite it
     RewriteToken(MacroNameTok);
 
-    // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics
-    // remain unchanged.
-#if LLVM_VERSION_MAJOR > 4
-    #define GET_NUM_ARGS() getNumParams()
-#else
-    #define GET_NUM_ARGS() getNumArgs()
-#endif
-
     // If it's a macro with arguments, rewrite all the arguments as hip, too.
     for (unsigned int i = 0; Args && i < MD.getMacroInfo()->GET_NUM_ARGS(); i++) {
       std::vector<Token> toks;
@@ -400,11 +387,8 @@ public:
       // to workaround the 'const' MacroArgs passed into this hook.
       const Token *start = Args->getUnexpArgument(i);
       size_t len = Args->getArgLength(start) + 1;
-#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
-      _pp->EnterTokenStream(start, len, false, false);
-#else
-      _pp->EnterTokenStream(ArrayRef<Token>(start, len), false);
-#endif
+      llcompat::EnterPreprocessorTokenStream(*_pp, start, len, false);
+
       do {
         toks.push_back(Token());
         Token &tk = toks.back();
@@ -733,7 +717,7 @@ private:
   }
 
 public:
-  Cuda2HipCallback(Replacements *Replace, ast_matchers::MatchFinder *parent, HipifyPPCallbacks *PPCallbacks, const std::string &mainFileName)
+  Cuda2HipCallback(Replacements& Replace, ast_matchers::MatchFinder *parent, HipifyPPCallbacks *PPCallbacks, const std::string &mainFileName)
     : Cuda2Hip(Replace, mainFileName), owner(parent), PP(PPCallbacks) {
     PP->setMatch(this);
   }
@@ -1092,13 +1076,7 @@ int main(int argc, const char **argv) {
   auto start = std::chrono::steady_clock::now();
   auto begin = start;
 
-  // The signature of PrintStackTraceOnErrorSignal changed in llvm 3.9. We don't support
-  // anything older than 3.8, so let's specifically detect the one old version we support.
-#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
-  llvm::sys::PrintStackTraceOnErrorSignal();
-#else
-  llvm::sys::PrintStackTraceOnErrorSignal(StringRef());
-#endif
+  llcompat::PrintStackTraceOnErrorSignal();
 
   CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore);
   std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
@@ -1161,15 +1139,7 @@ int main(int argc, const char **argv) {
     ast_matchers::MatchFinder Finder;
 
     // The Replacements to apply to the file `src`.
-    Replacements* replacementsToUse;
-#if LLVM_VERSION_MAJOR > 3
-    // getReplacements() now returns a map from filename to Replacements - so create an entry
-    // for this source file and return a pointer to it.
-    replacementsToUse = &(Tool.getReplacements()[tmpFile]);
-#else
-    replacementsToUse = &Tool.getReplacements();
-#endif
-
+    Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile);
     HipifyPPCallbacks* PPCallbacks = new HipifyPPCallbacks(replacementsToUse, tmpFile);
     Cuda2HipCallback Callback(replacementsToUse, &Finder, PPCallbacks, tmpFile);
 
@@ -1199,11 +1169,8 @@ int main(int argc, const char **argv) {
     SourceManager SM(Diagnostics, Tool.getFiles());
     if (PrintStats) {
       DEBUG(dbgs() << "Replacements collected by the tool:\n");
-#if LLVM_VERSION_MAJOR > 3
-        Replacements& replacements = Tool.getReplacements().begin()->second;
-#else
-        Replacements& replacements = Tool.getReplacements();
-#endif
+
+      Replacements& replacements = llcompat::getReplacements(Tool, tmpFile);
       for (const auto &replacement : replacements) {
         DEBUG(dbgs() << replacement.toString() << "\n");
         repBytes += replacement.getLength();
diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp
new file mode 100644
index 0000000000..ef9bdc193c
--- /dev/null
+++ b/projects/hip/hipify-clang/src/LLVMCompat.cpp
@@ -0,0 +1,43 @@
+#include "LLVMCompat.h"
+
+namespace llcompat {
+
+void PrintStackTraceOnErrorSignal() {
+    // The signature of PrintStackTraceOnErrorSignal changed in llvm 3.9. We don't support
+    // anything older than 3.8, so let's specifically detect the one old version we support.
+#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
+    llvm::sys::PrintStackTraceOnErrorSignal();
+#else
+    llvm::sys::PrintStackTraceOnErrorSignal(clang::StringRef());
+#endif
+}
+
+ct::Replacements& getReplacements(ct::RefactoringTool& Tool, clang::StringRef file) {
+#if LLVM_VERSION_MAJOR > 3
+    // getReplacements() now returns a map from filename to Replacements - so create an entry
+    // for this source file and return a reference to it.
+    return Tool.getReplacements()[file];
+#else
+    return Tool.getReplacements();
+#endif
+}
+
+void insertReplacement(ct::Replacements& replacements, const ct::Replacement& rep) {
+#if LLVM_VERSION_MAJOR > 3
+    // New clang added error checking to Replacements, and *insists* that you explicitly check it.
+    llvm::Error e = replacements.add(rep);
+#else
+    // In older versions, it's literally an std::set<Replacement>
+    replacements.insert(rep);
+#endif
+}
+
+void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *start, size_t len, bool DisableMacroExpansion) {
+#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
+    _pp.EnterTokenStream(start, len, false, DisableMacroExpansion;
+#else
+    _pp.EnterTokenStream(clang::ArrayRef<clang::Token>{start, len}, DisableMacroExpansion);
+#endif
+}
+
+} // namespace llcompat
diff --git a/projects/hip/hipify-clang/src/LLVMCompat.h b/projects/hip/hipify-clang/src/LLVMCompat.h
new file mode 100644
index 0000000000..3e2fe1aebb
--- /dev/null
+++ b/projects/hip/hipify-clang/src/LLVMCompat.h
@@ -0,0 +1,49 @@
+#pragma once
+
+#include <clang/Tooling/Core/Replacement.h>
+#include <clang/Tooling/Refactoring.h>
+#include <llvm/Support/Signals.h>
+#include <clang/Lex/Token.h>
+#include <clang/Lex/Preprocessor.h>
+
+namespace ct = clang::tooling;
+
+// Things for papering over the differences between different LLVM versions.
+
+namespace llcompat {
+
+
+/**
+ * The getNumArgs function on macros was rather unhelpfully renamed in clang 4.0. Its semantics
+ * remain unchanged, so let's be slightly ugly about it here. :D
+ */
+#if LLVM_VERSION_MAJOR > 4
+    #define GET_NUM_ARGS() getNumParams()
+#else
+    #define GET_NUM_ARGS() getNumArgs()
+#endif
+
+void PrintStackTraceOnErrorSignal();
+
+/**
+ * Get the replacement map for a given filename in a RefactoringTool.
+ *
+ * Older LLVM versions don't actually support multiple filenames, so everything all gets
+ * smushed together. It is the caller's responsibility to cope with this.
+ */
+ct::Replacements& getReplacements(ct::RefactoringTool& Tool, clang::StringRef file);
+
+/**
+ * Add a Replacement to a Replacements.
+ */
+void insertReplacement(ct::Replacements& replacements, const ct::Replacement& rep);
+
+/**
+ * Version-agnostic version of Preprocessor::EnterTokenStream().
+ */
+void EnterPreprocessorTokenStream(clang::Preprocessor& _pp,
+                                  const clang::Token *start,
+                                  size_t len,
+                                  bool DisableMacroExpansion);
+
+} // namespace llcompat

From 25517e41fdf5448727927ea952d7cf30f268550f Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 25 Oct 2017 19:23:32 +0100
Subject: [PATCH 21/28] Move string utility functions into their own
 translation unit

[ROCm/hip commit: ee8e11a72094456697021a74bff05b5dbf2e4e11]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp    | 18 +-----------------
 projects/hip/hipify-clang/src/StringUtils.cpp | 17 +++++++++++++++++
 projects/hip/hipify-clang/src/StringUtils.h   | 14 ++++++++++++++
 3 files changed, 32 insertions(+), 17 deletions(-)
 create mode 100644 projects/hip/hipify-clang/src/StringUtils.cpp
 create mode 100644 projects/hip/hipify-clang/src/StringUtils.h

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index 270dec68b3..dabcfe05ed 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -55,6 +55,7 @@ THE SOFTWARE.
 #include "CUDA2HipMap.h"
 #include "Types.h"
 #include "LLVMCompat.h"
+#include "StringUtils.h"
 
 using namespace clang;
 using namespace clang::ast_matchers;
@@ -121,23 +122,6 @@ uint64_t countApiRepsTotalUnsupported[API_LAST] = { 0 };
 std::map<std::string, uint64_t> cuda2hipConvertedTotal;
 std::map<std::string, uint64_t> cuda2hipUnconvertedTotal;
 
-StringRef unquoteStr(StringRef s) {
-  if (s.size() > 1 && s.front() == '"' && s.back() == '"')
-    return s.substr(1, s.size() - 2);
-  return s;
-}
-
-/**
- * If `s` starts with `prefix`, remove it. Otherwise, does nothing.
- */
-void removePrefixIfPresent(std::string& s, std::string prefix) {
-  if (s.find(prefix) != 0) {
-    return;
-  }
-
-  s.erase(0, prefix.size());
-}
-
 class Cuda2Hip {
 public:
   Cuda2Hip(Replacements& R, const std::string &srcFileName) :
diff --git a/projects/hip/hipify-clang/src/StringUtils.cpp b/projects/hip/hipify-clang/src/StringUtils.cpp
new file mode 100644
index 0000000000..ad55333bc8
--- /dev/null
+++ b/projects/hip/hipify-clang/src/StringUtils.cpp
@@ -0,0 +1,17 @@
+#include "StringUtils.h"
+
+llvm::StringRef unquoteStr(llvm::StringRef s) {
+    if (s.size() > 1 && s.front() == '"' && s.back() == '"') {
+        return s.substr(1, s.size() - 2);
+    }
+
+    return s;
+}
+
+void removePrefixIfPresent(std::string &s, std::string prefix) {
+    if (s.find(prefix) != 0) {
+        return;
+    }
+
+    s.erase(0, prefix.size());
+}
diff --git a/projects/hip/hipify-clang/src/StringUtils.h b/projects/hip/hipify-clang/src/StringUtils.h
new file mode 100644
index 0000000000..66a9be780f
--- /dev/null
+++ b/projects/hip/hipify-clang/src/StringUtils.h
@@ -0,0 +1,14 @@
+#pragma once
+
+#include <string>
+#include "llvm/ADT/StringRef.h"
+
+/**
+ * Remove double-quotes from the start/end of a string, if present.
+ */
+llvm::StringRef unquoteStr(llvm::StringRef s);
+
+/**
+ * If `s` starts with `prefix`, remove it. Otherwise, does nothing.
+ */
+void removePrefixIfPresent(std::string &s, std::string prefix);

From 51df5b20c9522b0c571427a3d8acffbf2e763439 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 25 Oct 2017 19:41:18 +0100
Subject: [PATCH 22/28] Prefer references to pointers in updateCountersExt()

[ROCm/hip commit: 00bb447e551b3a5b669d4bf6d92e6571e521403d]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 20 ++++++++++----------
 1 file changed, 10 insertions(+), 10 deletions(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index dabcfe05ed..bc4bdb4587 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -175,21 +175,21 @@ protected:
   }
 
   void updateCountersExt(const hipCounter &counter, const std::string &cudaName) {
-    std::map<std::string, uint64_t> *map = &cuda2hipConverted;
-    std::map<std::string, uint64_t> *mapTotal = &cuda2hipConvertedTotal;
+    std::map<std::string, uint64_t>& map = cuda2hipConverted;
+    std::map<std::string, uint64_t>& mapTotal = cuda2hipConvertedTotal;
     if (counter.unsupported) {
-      map = &cuda2hipUnconverted;
-      mapTotal = &cuda2hipUnconvertedTotal;
+      map = cuda2hipUnconverted;
+      mapTotal = cuda2hipUnconvertedTotal;
     }
-    auto found = map->find(cudaName);
-    if (found == map->end()) {
-      map->insert(std::pair<std::string, uint64_t>(cudaName, 1));
+    auto found = map.find(cudaName);
+    if (found == map.end()) {
+      map.insert(std::pair<std::string, uint64_t>(cudaName, 1));
     } else {
       found->second++;
     }
-    auto foundT = mapTotal->find(cudaName);
-    if (foundT == mapTotal->end()) {
-      mapTotal->insert(std::pair<std::string, uint64_t>(cudaName, 1));
+    auto foundT = mapTotal.find(cudaName);
+    if (foundT == mapTotal.end()) {
+      mapTotal.insert(std::pair<std::string, uint64_t>(cudaName, 1));
     } else {
       foundT->second++;
     }

From a3c1d30745d40d07a4f69ff618c2e65d837464a6 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 25 Oct 2017 19:42:46 +0100
Subject: [PATCH 23/28] 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: d8beee891855dfe23c97e7d1e0a4179185fd040e]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 15 +++------------
 1 file changed, 3 insertions(+), 12 deletions(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index bc4bdb4587..af44f84ba2 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -181,18 +181,9 @@ protected:
       map = cuda2hipUnconverted;
       mapTotal = cuda2hipUnconvertedTotal;
     }
-    auto found = map.find(cudaName);
-    if (found == map.end()) {
-      map.insert(std::pair<std::string, uint64_t>(cudaName, 1));
-    } else {
-      found->second++;
-    }
-    auto foundT = mapTotal.find(cudaName);
-    if (foundT == mapTotal.end()) {
-      mapTotal.insert(std::pair<std::string, uint64_t>(cudaName, 1));
-    } else {
-      foundT->second++;
-    }
+
+    map[cudaName]++;
+    mapTotal[cudaName]++;
   }
 
   virtual void updateCounters(const hipCounter &counter, const std::string &cudaName) {

From e4569bc84e3f4ba6c8ec768aef9860e60fd965af Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 25 Oct 2017 19:44:21 +0100
Subject: [PATCH 24/28] Inline updateCountersExt

[ROCm/hip commit: 50448aec3ba6ea7852308f9cd9be35d76920e58f]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 21 ++++++++-------------
 1 file changed, 8 insertions(+), 13 deletions(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index af44f84ba2..5c3e41d58d 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -174,24 +174,16 @@ protected:
     llvm::errs() << "[HIPIFY] " << getMsgType(msgType) << ": " << mainFileName << ":" << fullSL.getExpansionLineNumber() << ":" << fullSL.getExpansionColumnNumber() << ": " << message << "\n";
   }
 
-  void updateCountersExt(const hipCounter &counter, const std::string &cudaName) {
+  virtual void updateCounters(const hipCounter &counter, const std::string &cudaName) {
+    if (!PrintStats) {
+      return;
+    }
+
     std::map<std::string, uint64_t>& map = cuda2hipConverted;
     std::map<std::string, uint64_t>& mapTotal = cuda2hipConvertedTotal;
     if (counter.unsupported) {
       map = cuda2hipUnconverted;
       mapTotal = cuda2hipUnconvertedTotal;
-    }
-
-    map[cudaName]++;
-    mapTotal[cudaName]++;
-  }
-
-  virtual void updateCounters(const hipCounter &counter, const std::string &cudaName) {
-    if (!PrintStats) {
-      return;
-    }
-    updateCountersExt(counter, cudaName);
-    if (counter.unsupported) {
       countRepsUnsupported[counter.countType]++;
       countRepsTotalUnsupported[counter.countType]++;
       countApiRepsUnsupported[counter.countApiType]++;
@@ -202,6 +194,9 @@ protected:
       countApiReps[counter.countApiType]++;
       countApiRepsTotal[counter.countApiType]++;
     }
+
+    map[cudaName]++;
+    mapTotal[cudaName]++;
   }
 
   void processString(StringRef s, SourceManager &SM, SourceLocation start) {

From 1454bf96512249d7cd87e01c5b383a80fbb24969 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Wed, 25 Oct 2017 20:02:59 +0100
Subject: [PATCH 25/28] Copy-paste less in the statistics printing code

[ROCm/hip commit: 5699c18adc6f0be22b6d9257845599c1090ef0fe]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 162 +++++++++------------
 1 file changed, 65 insertions(+), 97 deletions(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index 5c3e41d58d..df3bfaf040 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -787,6 +787,15 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac
   );
 }
 
+/**
+ * Print a named stat value to both the terminal and the CSV file.
+ */
+template<typename T>
+void printStat(std::ofstream &csv, const std::string& name, T value) {
+  llvm::outs() << "  " << name << ": " << value << "\n";
+  csv << name << ";" << value << "\n";
+}
+
 int64_t printStats(const std::string &csvFile, const std::string &srcFile,
                    HipifyPPCallbacks &PPCallbacks, Cuda2HipCallback &Callback,
                    uint64_t replacedBytes, uint64_t totalBytes, unsigned totalLines,
@@ -802,53 +811,37 @@ int64_t printStats(const std::string &csvFile, const std::string &srcFile,
   for (int i = 0; i < CONV_LAST; i++) {
     sum_unsupported += Callback.countRepsUnsupported[i] + PPCallbacks.countRepsUnsupported[i];
   }
+
   if (sum > 0 || sum_unsupported > 0) {
     str = "file \'" + srcFile + "\' statistics:\n";
     llvm::outs() << "\n" << hipify_info << str;
     csv << "\n" << str;
-    str = "CONVERTED refs count";
-    llvm::outs() << "  " << str << ": " << sum << "\n";
-    csv << "\n" << str << separator << sum << "\n";
-    str = "UNCONVERTED refs count";
-    llvm::outs() << "  " << str << ": " << sum_unsupported << "\n";
-    csv << str << separator << sum_unsupported << "\n";
-    str = "CONVERSION %";
-    long conv = 100 - std::lround(double(sum_unsupported*100)/double(sum + sum_unsupported));
-    llvm::outs() << "  " << str << ": " << conv << "%\n";
-    csv << str << separator << conv << "%\n";
-    str = "REPLACED bytes";
-    llvm::outs() << "  " << str << ": " << replacedBytes << "\n";
-    csv << str << separator << replacedBytes << "\n";
-    str = "TOTAL bytes";
-    llvm::outs() << "  " << str << ": " << totalBytes << "\n";
-    csv << str << separator << totalBytes << "\n";
-    str = "CHANGED lines of code";
-    unsigned changedLines = Callback.LOCs.size() + PPCallbacks.LOCs.size();
-    llvm::outs() << "  " << str << ": " << changedLines << "\n";
-    csv << str << separator << changedLines << "\n";
-    str = "TOTAL lines of code";
-    llvm::outs() << "  " << str << ": " << totalLines << "\n";
-    csv << str << separator << totalLines << "\n";
+
+    size_t changedLines = Callback.LOCs.size() + PPCallbacks.LOCs.size();
+
+    printStat(csv, "CONVERTED refs count", sum);
+    printStat(csv, "UNCONVERTED refs count", sum_unsupported);
+    printStat(csv, "CONVERSION %", 100 - std::lround(double(sum_unsupported * 100) / double(sum + sum_unsupported)));
+    printStat(csv, "REPLACED bytes", replacedBytes);
+    printStat(csv, "TOTAL bytes", totalBytes);
+    printStat(csv, "CHANGED lines of code", changedLines);
+    printStat(csv, "TOTAL lines of code", totalLines);
+
     if (totalBytes > 0) {
-      str = "CODE CHANGED (in bytes) %";
-      conv = std::lround(double(replacedBytes * 100) / double(totalBytes));
-      llvm::outs() << "  " << str << ": " << conv << "%\n";
-      csv << str << separator << conv << "%\n";
+      printStat(csv, "CODE CHANGED (in bytes) %", std::lround(double(replacedBytes * 100) / double(totalBytes)));
     }
+
     if (totalLines > 0) {
-      str = "CODE CHANGED (in lines) %";
-      conv = std::lround(double(changedLines * 100) / double(totalLines));
-      llvm::outs() << "  " << str << ": " << conv << "%\n";
-      csv << str << separator << conv << "%\n";
+      printStat(csv, "CODE CHANGED (in lines) %", std::lround(double(changedLines * 100) / double(totalLines)));
     }
+
     typedef std::chrono::duration<double, std::milli> duration;
     duration elapsed = std::chrono::steady_clock::now() - start;
-    str = "TIME ELAPSED s";
     std::stringstream stream;
     stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000;
-    llvm::outs() << "  " << str << ": " << stream.str() << "\n";
-    csv << str << separator << stream.str() << "\n";
+    printStat(csv, "TIME ELAPSED s", stream.str());
   }
+
   if (sum > 0) {
     llvm::outs() << hipify_info << "CONVERTED refs by type:\n";
     csv << "\nCUDA ref type" << separator << "Count\n";
@@ -857,14 +850,12 @@ int64_t printStats(const std::string &csvFile, const std::string &srcFile,
       if (0 == sum_interm) {
         continue;
       }
-      llvm::outs() << "  " << counterNames[i] << ": " << sum_interm << "\n";
-      csv << counterNames[i] << separator << sum_interm << "\n";
+      printStat(csv, counterNames[i], sum_interm);
     }
     llvm::outs() << hipify_info << "CONVERTED refs by API:\n";
     csv << "\nCUDA API" << separator << "Count\n";
     for (int i = 0; i < API_LAST; i++) {
-      llvm::outs() << "  " << apiNames[i] << ": " << Callback.countApiReps[i] + PPCallbacks.countApiReps[i] << "\n";
-      csv << apiNames[i] << separator << Callback.countApiReps[i] + PPCallbacks.countApiReps[i] << "\n";
+      printStat(csv, apiNames[i], Callback.countApiReps[i] + PPCallbacks.countApiReps[i]);
     }
     for (const auto & it : PPCallbacks.cuda2hipConverted) {
       const auto found = Callback.cuda2hipConverted.find(it.first);
@@ -877,10 +868,10 @@ int64_t printStats(const std::string &csvFile, const std::string &srcFile,
     llvm::outs() << hipify_info << "CONVERTED refs by names:\n";
     csv << "\nCUDA ref name" << separator << "Count\n";
     for (const auto & it : Callback.cuda2hipConverted) {
-      llvm::outs() << "  " << it.first << ": " << it.second << "\n";
-      csv << it.first << separator << it.second << "\n";
+      printStat(csv, it.first, it.second);
     }
   }
+
   if (sum_unsupported > 0) {
     str = "UNCONVERTED refs by type:";
     llvm::outs() << hipify_info << str << "\n";
@@ -890,14 +881,12 @@ int64_t printStats(const std::string &csvFile, const std::string &srcFile,
       if (0 == sum_interm) {
         continue;
       }
-      llvm::outs() << "  " << counterNames[i] << ": " << sum_interm << "\n";
-      csv << counterNames[i] << separator << sum_interm << "\n";
+      printStat(csv, counterNames[i], sum_interm);
     }
     llvm::outs() << hipify_info << "UNCONVERTED refs by API:\n";
     csv << "\nUNCONVERTED CUDA API" << separator << "Count\n";
     for (int i = 0; i < API_LAST; i++) {
-      llvm::outs() << "  " << apiNames[i] << ": " << Callback.countApiRepsUnsupported[i] + PPCallbacks.countApiRepsUnsupported[i] << "\n";
-      csv << apiNames[i] << separator << Callback.countApiRepsUnsupported[i] + PPCallbacks.countApiRepsUnsupported[i] << "\n";
+      printStat(csv, apiNames[i], Callback.countApiRepsUnsupported[i] + PPCallbacks.countApiRepsUnsupported[i]);
     }
     for (const auto & it : PPCallbacks.cuda2hipUnconverted) {
       const auto found = Callback.cuda2hipUnconverted.find(it.first);
@@ -907,13 +896,14 @@ int64_t printStats(const std::string &csvFile, const std::string &srcFile,
         found->second += it.second;
       }
     }
+
     llvm::outs() << hipify_info << "UNCONVERTED refs by names:\n";
     csv << "\nUNCONVERTED CUDA ref name" << separator << "Count\n";
     for (const auto & it : Callback.cuda2hipUnconverted) {
-      llvm::outs() << "  " << it.first << ": " << it.second << "\n";
-      csv << it.first << separator << it.second << "\n";
+      printStat(csv, it.first, it.second);
     }
   }
+
   csv.close();
   return sum;
 }
@@ -932,58 +922,34 @@ void printAllStats(const std::string &csvFile, int64_t totalFiles, int64_t conve
   for (int i = 0; i < CONV_LAST; i++) {
     sum_unsupported += countRepsTotalUnsupported[i];
   }
+
   if (sum > 0 || sum_unsupported > 0) {
     str = "TOTAL statistics:\n";
     llvm::outs() << "\n" << hipify_info << str;
     csv << "\n" << str;
-    str = "CONVERTED files";
-    llvm::outs() << "  " << str << ": " << convertedFiles << "\n";
-    csv << "\n" << str << separator << convertedFiles << "\n";
-    str = "PROCESSED files";
-    llvm::outs() << "  " << str << ": " << totalFiles << "\n";
-    csv << str << separator << totalFiles << "\n";
-    str = "CONVERTED refs count";
-    llvm::outs() << "  " << str << ": " << sum << "\n";
-    csv << str << separator << sum << "\n";
-    str = "UNCONVERTED refs count";
-    llvm::outs() << "  " << str << ": " << sum_unsupported << "\n";
-    csv << str << separator << sum_unsupported << "\n";
-    str = "CONVERSION %";
-    long conv = 100 - std::lround(double(sum_unsupported * 100) / double(sum + sum_unsupported));
-    llvm::outs() << "  " << str << ": " << conv << "%\n";
-    csv << str << separator << conv << "%\n";
-    str = "REPLACED bytes";
-    llvm::outs() << "  " << str << ": " << replacedBytes << "\n";
-    csv << str << separator << replacedBytes << "\n";
-    str = "TOTAL bytes";
-    llvm::outs() << "  " << str << ": " << totalBytes << "\n";
-    csv << str << separator << totalBytes << "\n";
-    str = "CHANGED lines of code";
-    llvm::outs() << "  " << str << ": " << changedLines << "\n";
-    csv << str << separator << changedLines << "\n";
-    str = "TOTAL lines of code";
-    llvm::outs() << "  " << str << ": " << totalLines << "\n";
-    csv << str << separator << totalLines << "\n";
+    printStat(csv, "CONVERTED files", convertedFiles);
+    printStat(csv, "PROCESSED files", totalFiles);
+    printStat(csv, "CONVERTED refs count", sum);
+    printStat(csv, "UNCONVERTED refs count", sum_unsupported);
+    printStat(csv, "CONVERSION %", 100 - std::lround(double(sum_unsupported * 100) / double(sum + sum_unsupported)));
+    printStat(csv, "REPLACED bytes", replacedBytes);
+    printStat(csv, "TOTAL bytes", totalBytes);
+    printStat(csv, "CHANGED lines of code", changedLines);
+    printStat(csv, "TOTAL lines of code", totalLines);
     if (totalBytes > 0) {
-      str = "CODE CHANGED (in bytes) %";
-      conv = std::lround(double(replacedBytes * 100) / double(totalBytes));
-      llvm::outs() << "  " << str << ": " << conv << "%\n";
-      csv << str << separator << conv << "%\n";
+      printStat(csv, "CODE CHANGED (in bytes) %", std::lround(double(replacedBytes * 100) / double(totalBytes)));
     }
     if (totalLines > 0) {
-      str = "CODE CHANGED (in lines) %";
-      conv = std::lround(double(changedLines * 100) / double(totalLines));
-      llvm::outs() << "  " << str << ": " << conv << "%\n";
-      csv << str << separator << conv << "%\n";
+      printStat(csv, "CODE CHANGED (in lines) %", std::lround(double(changedLines * 100) / double(totalLines)));
     }
+
     typedef std::chrono::duration<double, std::milli> duration;
     duration elapsed = std::chrono::steady_clock::now() - start;
-    str = "TIME ELAPSED s";
     std::stringstream stream;
     stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000;
-    llvm::outs() << "  " << str << ": " << stream.str() << "\n";
-    csv << str << separator << stream.str() << "\n";
+    printStat(csv, "TIME ELAPSED s", stream.str());
   }
+
   if (sum > 0) {
     llvm::outs() << hipify_info << "CONVERTED refs by type:\n";
     csv << "\nCUDA ref type" << separator << "Count\n";
@@ -992,22 +958,23 @@ void printAllStats(const std::string &csvFile, int64_t totalFiles, int64_t conve
       if (0 == sum_interm) {
         continue;
       }
-      llvm::outs() << "  " << counterNames[i] << ": " << sum_interm << "\n";
-      csv << counterNames[i] << separator << sum_interm << "\n";
+
+      printStat(csv, counterNames[i], sum_interm);
     }
+
     llvm::outs() << hipify_info << "CONVERTED refs by API:\n";
     csv << "\nCUDA API" << separator << "Count\n";
     for (int i = 0; i < API_LAST; i++) {
-      llvm::outs() << "  " << apiNames[i] << ": " << countApiRepsTotal[i] << "\n";
-      csv << apiNames[i] << separator << countApiRepsTotal[i] << "\n";
+      printStat(csv, apiNames[i], countApiRepsTotal[i]);
     }
+
     llvm::outs() << hipify_info << "CONVERTED refs by names:\n";
     csv << "\nCUDA ref name" << separator << "Count\n";
     for (const auto & it : cuda2hipConvertedTotal) {
-      llvm::outs() << "  " << it.first << ": " << it.second << "\n";
-      csv << it.first << separator << it.second << "\n";
+      printStat(csv, it.first, it.second);
     }
   }
+
   if (sum_unsupported > 0) {
     str = "UNCONVERTED refs by type:";
     llvm::outs() << hipify_info << str << "\n";
@@ -1017,22 +984,23 @@ void printAllStats(const std::string &csvFile, int64_t totalFiles, int64_t conve
       if (0 == sum_interm) {
         continue;
       }
-      llvm::outs() << "  " << counterNames[i] << ": " << sum_interm << "\n";
-      csv << counterNames[i] << separator << sum_interm << "\n";
+
+      printStat(csv, counterNames[i], sum_interm);
     }
+
     llvm::outs() << hipify_info << "UNCONVERTED refs by API:\n";
     csv << "\nUNCONVERTED CUDA API" << separator << "Count\n";
     for (int i = 0; i < API_LAST; i++) {
-      llvm::outs() << "  " << apiNames[i] << ": " << countApiRepsTotalUnsupported[i] << "\n";
-      csv << apiNames[i] << separator << countApiRepsTotalUnsupported[i] << "\n";
+      printStat(csv, apiNames[i], countApiRepsTotalUnsupported[i]);
     }
+
     llvm::outs() << hipify_info << "UNCONVERTED refs by names:\n";
     csv << "\nUNCONVERTED CUDA ref name" << separator << "Count\n";
     for (const auto & it : cuda2hipUnconvertedTotal) {
-      llvm::outs() << "  " << it.first << ": " << it.second << "\n";
-      csv << it.first << separator << it.second << "\n";
+      printStat(csv, it.first, it.second);
     }
   }
+
   csv.close();
 }
 

From 05699e37792f04e00adde3a77cd9c07a49617fba Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Thu, 26 Oct 2017 04:25:05 +0100
Subject: [PATCH 26/28] 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: b303ffe53e3665bbd1322b7333de311a828b2790]
---
 projects/hip/hipify-clang/src/CUDA2HipMap.h  |  10 +-
 projects/hip/hipify-clang/src/Cuda2Hip.cpp   | 369 ++-----------------
 projects/hip/hipify-clang/src/Statistics.cpp | 227 ++++++++++++
 projects/hip/hipify-clang/src/Statistics.h   | 174 +++++++++
 projects/hip/hipify-clang/src/Types.h        |  47 ---
 5 files changed, 440 insertions(+), 387 deletions(-)
 create mode 100644 projects/hip/hipify-clang/src/Statistics.cpp
 create mode 100644 projects/hip/hipify-clang/src/Statistics.h
 delete mode 100644 projects/hip/hipify-clang/src/Types.h

diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.h b/projects/hip/hipify-clang/src/CUDA2HipMap.h
index ecef9ce2b5..605acf7aac 100644
--- a/projects/hip/hipify-clang/src/CUDA2HipMap.h
+++ b/projects/hip/hipify-clang/src/CUDA2HipMap.h
@@ -4,15 +4,7 @@
 #include <set>
 #include <map>
 
-#include "Types.h"
-
-// TODO: This shouldn't really be here. More restructuring needed...
-struct hipCounter {
-    llvm::StringRef hipName;
-    ConvTypes countType;
-    ApiTypes countApiType;
-    bool unsupported;
-};
+#include "Statistics.h"
 
 #define HIP_UNSUPPORTED true
 
diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index df3bfaf040..ae83d7f794 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -53,7 +53,6 @@ THE SOFTWARE.
 #include <sstream>
 
 #include "CUDA2HipMap.h"
-#include "Types.h"
 #include "LLVMCompat.h"
 #include "StringUtils.h"
 
@@ -64,16 +63,6 @@ using namespace llvm;
 
 #define DEBUG_TYPE "cuda2hip"
 
-const char *counterNames[CONV_LAST] = {
-    "version",      "init",     "device",  "mem",       "kern",        "coord_func", "math_func",
-    "special_func", "stream",   "event",   "occupancy", "ctx",         "peer",       "module",
-    "cache",        "exec",     "err",     "def",       "tex",         "gl",         "graphics",
-    "surface",      "jit",      "d3d9",    "d3d10",     "d3d11",       "vdpau",      "egl",
-    "thread",       "other",    "include", "include_cuda_main_header", "type",       "literal",
-    "numeric_literal"};
-
-const char *apiNames[API_LAST] = {
-    "CUDA Driver API", "CUDA RT API", "CUBLAS API"};
 
 // Set up the command line options
 static cl::OptionCategory ToolTemplateCategory("CUDA to HIP source translator options");
@@ -115,24 +104,10 @@ static cl::opt<bool> Examine("examine",
 
 static cl::extrahelp CommonHelp(CommonOptionsParser::HelpMessage);
 
-uint64_t countRepsTotal[CONV_LAST] = { 0 };
-uint64_t countApiRepsTotal[API_LAST] = { 0 };
-uint64_t countRepsTotalUnsupported[CONV_LAST] = { 0 };
-uint64_t countApiRepsTotalUnsupported[API_LAST] = { 0 };
-std::map<std::string, uint64_t> cuda2hipConvertedTotal;
-std::map<std::string, uint64_t> cuda2hipUnconvertedTotal;
-
 class Cuda2Hip {
 public:
   Cuda2Hip(Replacements& R, const std::string &srcFileName) :
     Replace(R), mainFileName(srcFileName) {}
-  uint64_t countReps[CONV_LAST] = { 0 };
-  uint64_t countApiReps[API_LAST] = { 0 };
-  uint64_t countRepsUnsupported[CONV_LAST] = { 0 };
-  uint64_t countApiRepsUnsupported[API_LAST] = { 0 };
-  std::map<std::string, uint64_t> cuda2hipConverted;
-  std::map<std::string, uint64_t> cuda2hipUnconverted;
-  std::set<unsigned> LOCs;
 
   enum msgTypes {
     HIPIFY_ERROR = 0,
@@ -154,14 +129,15 @@ protected:
   virtual void insertReplacement(const Replacement &rep, const FullSourceLoc &fullSL) {
     llcompat::insertReplacement(Replace, rep);
     if (PrintStats) {
-      LOCs.insert(fullSL.getExpansionLineNumber());
+      rep.getLength();
+      Statistics::current().lineTouched(fullSL.getExpansionLineNumber());
+      Statistics::current().bytesChanged(rep.getLength());
     }
   }
   void insertHipHeaders(Cuda2Hip *owner, const SourceManager &SM) {
-    if (owner->countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && Replace.size() > 0) {
+    if (Replace.size() > 0) {
       std::string repName = "#include <hip/hip_runtime.h>";
-      hipCounter counter = { repName, CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME };
-      updateCounters(counter, repName);
+      Statistics::current().incrementCounter({repName, ConvTypes::CONV_INCLUDE_CUDA_MAIN_H, ApiTypes::API_RUNTIME}, "#include <cuda>");
       SourceLocation sl = SM.getLocForStartOfFile(SM.getMainFileID());
       FullSourceLoc fullSL(sl, SM);
       Replacement Rep(SM, sl, 0, repName + "\n");
@@ -174,31 +150,6 @@ protected:
     llvm::errs() << "[HIPIFY] " << getMsgType(msgType) << ": " << mainFileName << ":" << fullSL.getExpansionLineNumber() << ":" << fullSL.getExpansionColumnNumber() << ": " << message << "\n";
   }
 
-  virtual void updateCounters(const hipCounter &counter, const std::string &cudaName) {
-    if (!PrintStats) {
-      return;
-    }
-
-    std::map<std::string, uint64_t>& map = cuda2hipConverted;
-    std::map<std::string, uint64_t>& mapTotal = cuda2hipConvertedTotal;
-    if (counter.unsupported) {
-      map = cuda2hipUnconverted;
-      mapTotal = cuda2hipUnconvertedTotal;
-      countRepsUnsupported[counter.countType]++;
-      countRepsTotalUnsupported[counter.countType]++;
-      countApiRepsUnsupported[counter.countApiType]++;
-      countApiRepsTotalUnsupported[counter.countApiType]++;
-    } else {
-      countReps[counter.countType]++;
-      countRepsTotal[counter.countType]++;
-      countApiReps[counter.countApiType]++;
-      countApiRepsTotal[counter.countApiType]++;
-    }
-
-    map[cudaName]++;
-    mapTotal[cudaName]++;
-  }
-
   void processString(StringRef s, SourceManager &SM, SourceLocation start) {
     size_t begin = 0;
     while ((begin = s.find("cu", begin)) != StringRef::npos) {
@@ -207,8 +158,8 @@ protected:
       const auto found = CUDA_RENAMES_MAP().find(name);
       if (found != CUDA_RENAMES_MAP().end()) {
         StringRef repName = found->second.hipName;
-        hipCounter counter = {"", CONV_LITERAL, API_RUNTIME, found->second.unsupported};
-        updateCounters(counter, name.str());
+        hipCounter counter = {"[string literal]", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.unsupported};
+        Statistics::current().incrementCounter(counter, name.str());
         if (!counter.unsupported) {
           SourceLocation sl = start.getLocWithOffset(begin + 1);
           Replacement Rep(SM, sl, name.size(), repName);
@@ -266,7 +217,7 @@ public:
       return;
     }
 
-    updateCounters(found->second, file_name.str());
+    Statistics::current().incrementCounter(found->second, file_name.str());
     if (found->second.unsupported) {
       // An unsupported CUDA header? Oh dear. Print a warning.
       printHipifyMessage(*_sm, hash_loc, "Unsupported CUDA header used: " + file_name.str());
@@ -312,7 +263,7 @@ public:
       // So it's an identifier, but not CUDA? Boring.
       return;
     }
-    updateCounters(found->second, name.str());
+    Statistics::current().incrementCounter(found->second, name.str());
 
     SourceLocation sl = t.getLocation();
     if (found->second.unsupported) {
@@ -408,7 +359,7 @@ private:
     }
 
     const hipCounter& hipCtr = found->second;
-    updateCounters(found->second, name);
+    Statistics::current().incrementCounter(hipCtr, name);
 
     if (hipCtr.unsupported) {
       return true; // Silently fail when you find an unsupported member.
@@ -416,7 +367,6 @@ private:
     }
 
     size_t length = name.size();
-    updateCounters(found->second, name);
     Replacement Rep(*SM, sl, length, hipCtr.hipName);
     FullSourceLoc fullSL(sl, *SM);
     insertReplacement(Rep, fullSL);
@@ -517,8 +467,8 @@ private:
       Replacement Rep(*SM, launchStart, length, OS.str());
       FullSourceLoc fullSL(launchStart, *SM);
       insertReplacement(Rep, fullSL);
-      hipCounter counter = {"hipLaunchKernelGGL", CONV_KERN, API_RUNTIME};
-      updateCounters(counter, refName.str());
+      hipCounter counter = {"hipLaunchKernelGGL", ConvTypes::CONV_KERN, ApiTypes::API_RUNTIME};
+      Statistics::current().incrementCounter(counter, refName.str());
       return true;
     }
     return false;
@@ -542,7 +492,7 @@ private:
           // TODO: Make a lookup table just for builtins to improve performance.
           const auto found = CUDA_IDENTIFIER_MAP.find(name);
           if (found != CUDA_IDENTIFIER_MAP.end()) {
-            updateCounters(found->second, name.str());
+            Statistics::current().incrementCounter(found->second, name.str());
             if (!found->second.unsupported) {
               StringRef repName = found->second.hipName;
               Replacement Rep(*SM, sl, name.size(), repName);
@@ -569,7 +519,7 @@ private:
       // TODO: Make a lookup table just for enum values to improve performance.
       const auto found = CUDA_IDENTIFIER_MAP.find(name);
       if (found != CUDA_IDENTIFIER_MAP.end()) {
-        updateCounters(found->second, name.str());
+        Statistics::current().incrementCounter(found->second, name.str());
         if (!found->second.unsupported) {
           StringRef repName = found->second.hipName;
           Replacement Rep(*SM, sl, name.size(), repName);
@@ -665,8 +615,8 @@ private:
           Replacement Rep(*SM, slStart, repLength, repName);
           FullSourceLoc fullSL(slStart, *SM);
           insertReplacement(Rep, fullSL);
-          hipCounter counter = { "HIP_DYNAMIC_SHARED", CONV_MEM, API_RUNTIME };
-          updateCounters(counter, refName.str());
+          hipCounter counter = { "HIP_DYNAMIC_SHARED", ConvTypes::CONV_MEM, ApiTypes::API_RUNTIME };
+          Statistics::current().incrementCounter(counter, refName.str());
         }
       }
       return true;
@@ -787,223 +737,6 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac
   );
 }
 
-/**
- * Print a named stat value to both the terminal and the CSV file.
- */
-template<typename T>
-void printStat(std::ofstream &csv, const std::string& name, T value) {
-  llvm::outs() << "  " << name << ": " << value << "\n";
-  csv << name << ";" << value << "\n";
-}
-
-int64_t printStats(const std::string &csvFile, const std::string &srcFile,
-                   HipifyPPCallbacks &PPCallbacks, Cuda2HipCallback &Callback,
-                   uint64_t replacedBytes, uint64_t totalBytes, unsigned totalLines,
-                   const std::chrono::steady_clock::time_point &start) {
-  std::ofstream csv(csvFile, std::ios::app);
-  int64_t sum = 0, sum_interm = 0;
-  std::string str;
-  const std::string hipify_info = "[HIPIFY] info: ", separator = ";";
-  for (int i = 0; i < CONV_LAST; i++) {
-    sum += Callback.countReps[i] + PPCallbacks.countReps[i];
-  }
-  int64_t sum_unsupported = 0;
-  for (int i = 0; i < CONV_LAST; i++) {
-    sum_unsupported += Callback.countRepsUnsupported[i] + PPCallbacks.countRepsUnsupported[i];
-  }
-
-  if (sum > 0 || sum_unsupported > 0) {
-    str = "file \'" + srcFile + "\' statistics:\n";
-    llvm::outs() << "\n" << hipify_info << str;
-    csv << "\n" << str;
-
-    size_t changedLines = Callback.LOCs.size() + PPCallbacks.LOCs.size();
-
-    printStat(csv, "CONVERTED refs count", sum);
-    printStat(csv, "UNCONVERTED refs count", sum_unsupported);
-    printStat(csv, "CONVERSION %", 100 - std::lround(double(sum_unsupported * 100) / double(sum + sum_unsupported)));
-    printStat(csv, "REPLACED bytes", replacedBytes);
-    printStat(csv, "TOTAL bytes", totalBytes);
-    printStat(csv, "CHANGED lines of code", changedLines);
-    printStat(csv, "TOTAL lines of code", totalLines);
-
-    if (totalBytes > 0) {
-      printStat(csv, "CODE CHANGED (in bytes) %", std::lround(double(replacedBytes * 100) / double(totalBytes)));
-    }
-
-    if (totalLines > 0) {
-      printStat(csv, "CODE CHANGED (in lines) %", std::lround(double(changedLines * 100) / double(totalLines)));
-    }
-
-    typedef std::chrono::duration<double, std::milli> duration;
-    duration elapsed = std::chrono::steady_clock::now() - start;
-    std::stringstream stream;
-    stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000;
-    printStat(csv, "TIME ELAPSED s", stream.str());
-  }
-
-  if (sum > 0) {
-    llvm::outs() << hipify_info << "CONVERTED refs by type:\n";
-    csv << "\nCUDA ref type" << separator << "Count\n";
-    for (int i = 0; i < CONV_LAST; i++) {
-      sum_interm = Callback.countReps[i] + PPCallbacks.countReps[i];
-      if (0 == sum_interm) {
-        continue;
-      }
-      printStat(csv, counterNames[i], sum_interm);
-    }
-    llvm::outs() << hipify_info << "CONVERTED refs by API:\n";
-    csv << "\nCUDA API" << separator << "Count\n";
-    for (int i = 0; i < API_LAST; i++) {
-      printStat(csv, apiNames[i], Callback.countApiReps[i] + PPCallbacks.countApiReps[i]);
-    }
-    for (const auto & it : PPCallbacks.cuda2hipConverted) {
-      const auto found = Callback.cuda2hipConverted.find(it.first);
-      if (found == Callback.cuda2hipConverted.end()) {
-        Callback.cuda2hipConverted.insert(std::pair<std::string, uint64_t>(it.first, 1));
-      } else {
-        found->second += it.second;
-      }
-    }
-    llvm::outs() << hipify_info << "CONVERTED refs by names:\n";
-    csv << "\nCUDA ref name" << separator << "Count\n";
-    for (const auto & it : Callback.cuda2hipConverted) {
-      printStat(csv, it.first, it.second);
-    }
-  }
-
-  if (sum_unsupported > 0) {
-    str = "UNCONVERTED refs by type:";
-    llvm::outs() << hipify_info << str << "\n";
-    csv << "\nUNCONVERTED CUDA ref type" << separator << "Count\n";
-    for (int i = 0; i < CONV_LAST; i++) {
-      sum_interm = Callback.countRepsUnsupported[i] + PPCallbacks.countRepsUnsupported[i];
-      if (0 == sum_interm) {
-        continue;
-      }
-      printStat(csv, counterNames[i], sum_interm);
-    }
-    llvm::outs() << hipify_info << "UNCONVERTED refs by API:\n";
-    csv << "\nUNCONVERTED CUDA API" << separator << "Count\n";
-    for (int i = 0; i < API_LAST; i++) {
-      printStat(csv, apiNames[i], Callback.countApiRepsUnsupported[i] + PPCallbacks.countApiRepsUnsupported[i]);
-    }
-    for (const auto & it : PPCallbacks.cuda2hipUnconverted) {
-      const auto found = Callback.cuda2hipUnconverted.find(it.first);
-      if (found == Callback.cuda2hipUnconverted.end()) {
-        Callback.cuda2hipUnconverted.insert(std::pair<std::string, uint64_t>(it.first, 1));
-      } else {
-        found->second += it.second;
-      }
-    }
-
-    llvm::outs() << hipify_info << "UNCONVERTED refs by names:\n";
-    csv << "\nUNCONVERTED CUDA ref name" << separator << "Count\n";
-    for (const auto & it : Callback.cuda2hipUnconverted) {
-      printStat(csv, it.first, it.second);
-    }
-  }
-
-  csv.close();
-  return sum;
-}
-
-void printAllStats(const std::string &csvFile, int64_t totalFiles, int64_t convertedFiles,
-                   uint64_t replacedBytes, uint64_t totalBytes, unsigned changedLines, unsigned totalLines,
-                   const std::chrono::steady_clock::time_point &start) {
-  std::ofstream csv(csvFile, std::ios::app);
-  int64_t sum = 0, sum_interm = 0;
-  std::string str;
-  const std::string hipify_info = "[HIPIFY] info: ", separator = ";";
-  for (int i = 0; i < CONV_LAST; i++) {
-    sum += countRepsTotal[i];
-  }
-  int64_t sum_unsupported = 0;
-  for (int i = 0; i < CONV_LAST; i++) {
-    sum_unsupported += countRepsTotalUnsupported[i];
-  }
-
-  if (sum > 0 || sum_unsupported > 0) {
-    str = "TOTAL statistics:\n";
-    llvm::outs() << "\n" << hipify_info << str;
-    csv << "\n" << str;
-    printStat(csv, "CONVERTED files", convertedFiles);
-    printStat(csv, "PROCESSED files", totalFiles);
-    printStat(csv, "CONVERTED refs count", sum);
-    printStat(csv, "UNCONVERTED refs count", sum_unsupported);
-    printStat(csv, "CONVERSION %", 100 - std::lround(double(sum_unsupported * 100) / double(sum + sum_unsupported)));
-    printStat(csv, "REPLACED bytes", replacedBytes);
-    printStat(csv, "TOTAL bytes", totalBytes);
-    printStat(csv, "CHANGED lines of code", changedLines);
-    printStat(csv, "TOTAL lines of code", totalLines);
-    if (totalBytes > 0) {
-      printStat(csv, "CODE CHANGED (in bytes) %", std::lround(double(replacedBytes * 100) / double(totalBytes)));
-    }
-    if (totalLines > 0) {
-      printStat(csv, "CODE CHANGED (in lines) %", std::lround(double(changedLines * 100) / double(totalLines)));
-    }
-
-    typedef std::chrono::duration<double, std::milli> duration;
-    duration elapsed = std::chrono::steady_clock::now() - start;
-    std::stringstream stream;
-    stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000;
-    printStat(csv, "TIME ELAPSED s", stream.str());
-  }
-
-  if (sum > 0) {
-    llvm::outs() << hipify_info << "CONVERTED refs by type:\n";
-    csv << "\nCUDA ref type" << separator << "Count\n";
-    for (int i = 0; i < CONV_LAST; i++) {
-      sum_interm = countRepsTotal[i];
-      if (0 == sum_interm) {
-        continue;
-      }
-
-      printStat(csv, counterNames[i], sum_interm);
-    }
-
-    llvm::outs() << hipify_info << "CONVERTED refs by API:\n";
-    csv << "\nCUDA API" << separator << "Count\n";
-    for (int i = 0; i < API_LAST; i++) {
-      printStat(csv, apiNames[i], countApiRepsTotal[i]);
-    }
-
-    llvm::outs() << hipify_info << "CONVERTED refs by names:\n";
-    csv << "\nCUDA ref name" << separator << "Count\n";
-    for (const auto & it : cuda2hipConvertedTotal) {
-      printStat(csv, it.first, it.second);
-    }
-  }
-
-  if (sum_unsupported > 0) {
-    str = "UNCONVERTED refs by type:";
-    llvm::outs() << hipify_info << str << "\n";
-    csv << "\nUNCONVERTED CUDA ref type" << separator << "Count\n";
-    for (int i = 0; i < CONV_LAST; i++) {
-      sum_interm = countRepsTotalUnsupported[i];
-      if (0 == sum_interm) {
-        continue;
-      }
-
-      printStat(csv, counterNames[i], sum_interm);
-    }
-
-    llvm::outs() << hipify_info << "UNCONVERTED refs by API:\n";
-    csv << "\nUNCONVERTED CUDA API" << separator << "Count\n";
-    for (int i = 0; i < API_LAST; i++) {
-      printStat(csv, apiNames[i], countApiRepsTotalUnsupported[i]);
-    }
-
-    llvm::outs() << hipify_info << "UNCONVERTED refs by names:\n";
-    csv << "\nUNCONVERTED CUDA ref name" << separator << "Count\n";
-    for (const auto & it : cuda2hipUnconvertedTotal) {
-      printStat(csv, it.first, it.second);
-    }
-  }
-
-  csv.close();
-}
-
 void copyFile(const std::string& src, const std::string& dst) {
   std::ifstream source(src, std::ios::binary);
   std::ofstream dest(dst, std::ios::binary);
@@ -1011,9 +744,6 @@ void copyFile(const std::string& src, const std::string& dst) {
 }
 
 int main(int argc, const char **argv) {
-  auto start = std::chrono::steady_clock::now();
-  auto begin = start;
-
   llcompat::PrintStackTraceOnErrorSignal();
 
   CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore);
@@ -1023,6 +753,7 @@ int main(int argc, const char **argv) {
     llvm::errs() << "[HIPIFY] conflict: -o and multiple source files are specified.\n";
     return 1;
   }
+
   if (NoOutput) {
     if (Inplace) {
       llvm::errs() << "[HIPIFY] conflict: both -no-output and -inplace options are specified.\n";
@@ -1033,24 +764,23 @@ int main(int argc, const char **argv) {
       return 1;
     }
   }
+
   if (Examine) {
     NoOutput = PrintStats = true;
   }
+
   int Result = 0;
-  std::string csv;
+
+  // Arguments for the Statistics print routines.
+  std::unique_ptr<std::ostream> csv = nullptr;
+  llvm::raw_ostream* statPrint = nullptr;
   if (!OutputStatsFilename.empty()) {
-    csv = OutputStatsFilename;
-  } else {
-    csv = "hipify_stats.csv";
+    csv = std::unique_ptr<std::ostream>(new std::ofstream(OutputStatsFilename, std::ios_base::trunc));
   }
-  size_t filesTranslated = fileSources.size();
-  uint64_t repBytesTotal = 0;
-  uint64_t bytesTotal = 0;
-  unsigned changedLinesTotal = 0;
-  unsigned linesTotal = 0;
-  if (PrintStats && filesTranslated > 1) {
-    std::remove(csv.c_str());
+  if (PrintStats) {
+    statPrint = &llvm::errs();
   }
+
   for (const auto & src : fileSources) {
     if (dst.empty()) {
       if (Inplace) {
@@ -1070,6 +800,9 @@ int main(int argc, const char **argv) {
     // Should we fail for some reason, we'll just leak this file and not corrupt the input.
     copyFile(src, tmpFile);
 
+    // Initialise the statistics counters for this file.
+    Statistics::setActive(src);
+
     // RefactoringTool operates on the file in-place. Giving it the output path is no good,
     // because that'll break relative includes, and we don't want to overwrite the input file.
     // So what we do is operate on a copy, which we then move to the output.
@@ -1101,24 +834,8 @@ int main(int argc, const char **argv) {
     TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts);
     DiagnosticsEngine Diagnostics(IntrusiveRefCntPtr<DiagnosticIDs>(new DiagnosticIDs()), &*DiagOpts, &DiagnosticPrinter, false);
 
-    uint64_t repBytes = 0;
-    uint64_t bytes = 0;
-    unsigned lines = 0;
     SourceManager SM(Diagnostics, Tool.getFiles());
-    if (PrintStats) {
-      DEBUG(dbgs() << "Replacements collected by the tool:\n");
 
-      Replacements& replacements = llcompat::getReplacements(Tool, tmpFile);
-      for (const auto &replacement : replacements) {
-        DEBUG(dbgs() << replacement.toString() << "\n");
-        repBytes += replacement.getLength();
-      }
-      std::ifstream src_file(dst, std::ios::binary | std::ios::ate);
-      src_file.clear();
-      src_file.seekg(0);
-      lines = std::count(std::istreambuf_iterator<char>(src_file), std::istreambuf_iterator<char>(), '\n');
-      bytes = src_file.tellg();
-    }
     Rewriter Rewrite(SM, DefaultLangOptions);
     if (!Tool.applyAllReplacements(Rewrite)) {
       DEBUG(dbgs() << "Skipped some replacements.\n");
@@ -1131,26 +848,16 @@ int main(int argc, const char **argv) {
     } else {
       remove(tmpFile.c_str());
     }
-    if (PrintStats) {
-      if (fileSources.size() == 1) {
-        if (OutputStatsFilename.empty()) {
-          csv = dst + ".csv";
-        }
-        std::remove(csv.c_str());
-      }
-      if (0 == printStats(csv, src, *PPCallbacks, Callback, repBytes, bytes, lines, start)) {
-        filesTranslated--;
-      }
-      start = std::chrono::steady_clock::now();
-      repBytesTotal += repBytes;
-      bytesTotal += bytes;
-      changedLinesTotal += PPCallbacks->LOCs.size() + Callback.LOCs.size();
-      linesTotal += lines;
-    }
+
+    Statistics::current().markCompletion();
+    Statistics::current().print(csv.get(), statPrint);
+
     dst.clear();
   }
-  if (PrintStats && fileSources.size() > 1) {
-    printAllStats(csv, fileSources.size(), filesTranslated, repBytesTotal, bytesTotal, changedLinesTotal, linesTotal, begin);
+
+  if (fileSources.size() > 1) {
+    Statistics::printAggregate(csv.get(), statPrint);
   }
+
   return Result;
 }
diff --git a/projects/hip/hipify-clang/src/Statistics.cpp b/projects/hip/hipify-clang/src/Statistics.cpp
new file mode 100644
index 0000000000..00b8a66f11
--- /dev/null
+++ b/projects/hip/hipify-clang/src/Statistics.cpp
@@ -0,0 +1,227 @@
+#include "Statistics.h"
+#include <assert.h>
+#include <sstream>
+#include <iomanip>
+
+
+const char *counterNames[NUM_CONV_TYPES] = {
+    "version", "init", "device", "mem", "kern", "coord_func", "math_func",
+    "special_func", "stream", "event", "occupancy", "ctx", "peer", "module",
+    "cache", "exec", "err", "def", "tex", "gl", "graphics",
+    "surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl",
+    "thread", "other", "include", "include_cuda_main_header", "type", "literal",
+    "numeric_literal"
+};
+
+const char *apiNames[NUM_API_TYPES] = {
+    "CUDA Driver API", "CUDA RT API", "CUBLAS API"
+};
+
+namespace {
+
+template<typename ST, typename ST2>
+void conditionalPrint(ST *stream1,
+                      ST2* stream2,
+                      const std::string& s1,
+                      const std::string& s2) {
+    if (stream1) {
+        *stream1 << s1;
+    }
+
+    if (stream2) {
+        *stream2 << s2;
+    }
+}
+
+
+/**
+ * Print a named stat value to both the terminal and the CSV file.
+ */
+template<typename T>
+void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string &name, T value) {
+    if (printOut) {
+        *printOut << "  " << name << ": " << value << "\n";
+    }
+
+    if (csv) {
+        *csv << name << ";" << value << "\n";
+    }
+}
+
+
+} // Anonymous namespace
+
+void StatCounter::incrementCounter(const hipCounter& counter, std::string name) {
+    counters[name]++;
+    apiCounters[(int) counter.countApiType]++;
+    convTypeCounters[(int) counter.countType]++;
+}
+
+void StatCounter::add(const StatCounter& other) {
+    for (const auto& p : other.counters) {
+        counters[p.first] += p.second;
+    }
+
+    for (int i = 0; i < NUM_API_TYPES; i++) {
+        apiCounters[i] += other.apiCounters[i];
+    }
+
+    for (int i = 0; i < NUM_CONV_TYPES; i++) {
+        convTypeCounters[i] += other.convTypeCounters[i];
+    }
+}
+
+int StatCounter::getConvSum() {
+    int acc = 0;
+    for (const int& i : convTypeCounters) {
+        acc += i;
+    }
+
+    return acc;
+}
+
+void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix) {
+    conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n", "[HIPIFY] info: " + prefix + " refs by type:\n");
+    for (int i = 0; i < NUM_CONV_TYPES; i++) {
+        if (convTypeCounters[i] > 0) {
+            printStat(csv, printOut, counterNames[i], convTypeCounters[i]);
+        }
+    }
+
+    conditionalPrint(csv, printOut, "\nCUDA API;Count\n", "[HIPIFY] info: " + prefix + " refs by API:\n");
+    for (int i = 0; i < NUM_API_TYPES; i++) {
+        printStat(csv, printOut, apiNames[i], apiCounters[i]);
+    }
+
+    conditionalPrint(csv, printOut, "\nCUDA ref name;Count\n", "[HIPIFY] info: " + prefix + " refs by names:\n");
+    for (const auto &it : counters) {
+        printStat(csv, printOut, it.first, it.second);
+    }
+}
+
+
+Statistics::Statistics(std::string name): fileName(name) {
+    // Compute the total bytes/lines in the input file.
+    std::ifstream src_file(name, std::ios::binary | std::ios::ate);
+    src_file.clear();
+    src_file.seekg(0);
+    totalLines = (int) std::count(std::istreambuf_iterator<char>(src_file), std::istreambuf_iterator<char>(), '\n');
+    totalBytes = (int) src_file.tellg();
+
+    // Mark the start time...
+    startTime = chr::steady_clock::now();
+};
+
+
+///////// Counter update routines //////////
+
+void Statistics::incrementCounter(const hipCounter &counter, std::string name) {
+    if (counter.unsupported) {
+        unsupported.incrementCounter(counter, name);
+    } else {
+        supported.incrementCounter(counter, name);
+    }
+}
+
+void Statistics::add(const Statistics &other) {
+    supported.add(other.supported);
+    unsupported.add(other.unsupported);
+    totalBytes += other.totalBytes;
+    totalLines += other.totalLines;
+    touchedBytes += other.touchedBytes;
+}
+
+void Statistics::lineTouched(int lineNumber) {
+    touchedLines.insert(lineNumber);
+}
+void Statistics::bytesChanged(int bytes) {
+    touchedBytes += bytes;
+}
+void Statistics::markCompletion() {
+    completionTime = chr::steady_clock::now();
+}
+
+
+///////// Output functions //////////
+
+void Statistics::print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader) {
+   if (!skipHeader) {
+       std::string str = "file \'" + fileName + "\' statistics:\n";
+       conditionalPrint(csv, printOut, "\n" + str, "\n[HIPIFY] info: " + str);
+   }
+
+   size_t changedLines = touchedLines.size();
+
+   // Total number of (un)supported refs that were converted.
+   int supportedSum = supported.getConvSum();
+   int unsupportedSum = unsupported.getConvSum();
+
+   printStat(csv, printOut, "CONVERTED refs count", supportedSum);
+   printStat(csv, printOut, "UNCONVERTED refs count", unsupportedSum);
+   printStat(csv, printOut, "CONVERSION %", 100 - std::lround(double(unsupportedSum * 100) / double(supportedSum + unsupportedSum)));
+   printStat(csv, printOut, "REPLACED bytes", touchedBytes);
+   printStat(csv, printOut, "TOTAL bytes", totalBytes);
+   printStat(csv, printOut, "CHANGED lines of code", changedLines);
+   printStat(csv, printOut, "TOTAL lines of code", totalLines);
+
+   if (totalBytes > 0) {
+       printStat(csv, printOut, "CODE CHANGED (in bytes) %", std::lround(double(touchedBytes * 100) / double(totalBytes)));
+   }
+
+   if (totalLines > 0) {
+       printStat(csv, printOut, "CODE CHANGED (in lines) %", std::lround(double(changedLines * 100) / double(totalLines)));
+   }
+
+   typedef std::chrono::duration<double, std::milli> duration;
+   duration elapsed = completionTime - startTime;
+   std::stringstream stream;
+   stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000;
+   printStat(csv, printOut, "TIME ELAPSED s", stream.str());
+
+   supported.print(csv, printOut, "CONVERTED");
+   unsupported.print(csv, printOut, "UNCONVERTED");
+}
+
+void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut) {
+    Statistics globalStats = getAggregate();
+
+    conditionalPrint(csv, printOut, "\nTOTAL statistics:\n", "\n[HIPIFY] info: TOTAL statistics:\n");
+
+    // A file is considered "converted" if we made any changes to it.
+    int convertedFiles = 0;
+    for (const auto& p : stats) {
+        if (!p.second.touchedLines.empty()) {
+            convertedFiles++;
+        }
+    }
+
+    printStat(csv, printOut, "CONVERTED files", convertedFiles);
+    printStat(csv, printOut, "PROCESSED files", stats.size());
+
+    globalStats.print(csv, printOut);
+}
+
+//// Static state management ////
+
+Statistics Statistics::getAggregate() {
+    Statistics globalStats("global");
+
+    for (const auto& p : stats) {
+        globalStats.add(p.second);
+    }
+
+    return globalStats;
+}
+
+Statistics& Statistics::current() {
+    assert(Statistics::currentStatistics);
+    return *Statistics::currentStatistics;
+}
+
+void Statistics::setActive(std::string name) {
+    stats.emplace(std::make_pair(name, Statistics{name}));
+    Statistics::currentStatistics = &stats.at(name);
+}
+
+std::map<std::string, Statistics> Statistics::stats = {};
+Statistics* Statistics::currentStatistics = nullptr;
diff --git a/projects/hip/hipify-clang/src/Statistics.h b/projects/hip/hipify-clang/src/Statistics.h
new file mode 100644
index 0000000000..da4e296db0
--- /dev/null
+++ b/projects/hip/hipify-clang/src/Statistics.h
@@ -0,0 +1,174 @@
+#pragma once
+
+#include <chrono>
+#include <string>
+#include <fstream>
+#include <llvm/ADT/StringRef.h>
+#include <map>
+#include <set>
+#include <llvm/Support/raw_ostream.h>
+
+namespace chr = std::chrono;
+
+enum ConvTypes {
+    CONV_VERSION = 0,
+    CONV_INIT,
+    CONV_DEVICE,
+    CONV_MEM,
+    CONV_KERN,
+    CONV_COORD_FUNC,
+    CONV_MATH_FUNC,
+    CONV_SPECIAL_FUNC,
+    CONV_STREAM,
+    CONV_EVENT,
+    CONV_OCCUPANCY,
+    CONV_CONTEXT,
+    CONV_PEER,
+    CONV_MODULE,
+    CONV_CACHE,
+    CONV_EXEC,
+    CONV_ERROR,
+    CONV_DEF,
+    CONV_TEX,
+    CONV_GL,
+    CONV_GRAPHICS,
+    CONV_SURFACE,
+    CONV_JIT,
+    CONV_D3D9,
+    CONV_D3D10,
+    CONV_D3D11,
+    CONV_VDPAU,
+    CONV_EGL,
+    CONV_THREAD,
+    CONV_OTHER,
+    CONV_INCLUDE,
+    CONV_INCLUDE_CUDA_MAIN_H,
+    CONV_TYPE,
+    CONV_LITERAL,
+    CONV_NUMERIC_LITERAL,
+    CONV_LAST
+};
+constexpr int NUM_CONV_TYPES = (int) ConvTypes::CONV_LAST;
+
+enum ApiTypes {
+    API_DRIVER = 0,
+    API_RUNTIME,
+    API_BLAS,
+    API_LAST
+};
+constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST;
+
+// The names of various fields in in the statistics reports.
+extern const char *counterNames[NUM_CONV_TYPES];
+extern const char *apiNames[NUM_API_TYPES];
+
+
+struct hipCounter {
+    llvm::StringRef hipName;
+    ConvTypes countType;
+    ApiTypes countApiType;
+    bool unsupported;
+};
+
+
+/**
+ * Tracks a set of named counters, as well as counters for each of the type enums defined above.
+ */
+class StatCounter {
+private:
+    // Each thing we track is either "supported" or "unsupported"...
+    std::map<std::string, int> counters;
+
+    int apiCounters[NUM_API_TYPES] = {};
+    int convTypeCounters[NUM_CONV_TYPES] = {};
+
+public:
+    void incrementCounter(const hipCounter& counter, std::string name);
+
+    /**
+     * Add the counters from `other` onto the counters of this object.
+     */
+    void add(const StatCounter& other);
+
+    int getConvSum();
+
+    void print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix);
+};
+
+/**
+ * Tracks the statistics for a single input file.
+ */
+class Statistics {
+    StatCounter supported;
+    StatCounter unsupported;
+
+    std::string fileName;
+
+    std::set<int> touchedLines = {};
+    int touchedBytes = 0;
+
+    int totalLines = 0;
+    int totalBytes = 0;
+
+    chr::steady_clock::time_point startTime;
+    chr::steady_clock::time_point completionTime;
+
+public:
+    Statistics(std::string name);
+
+    void incrementCounter(const hipCounter &counter, std::string name);
+
+    /**
+     * Add the counters from `other` onto the counters of this object.
+     */
+    void add(const Statistics &other);
+
+    void lineTouched(int lineNumber);
+    void bytesChanged(int bytes);
+
+    /**
+     * Set the completion timestamp to now.
+     */
+    void markCompletion();
+
+    /////// Output functions ///////
+
+public:
+   /**
+     * Pretty-print the statistics stored in this object.
+     *
+     * @param csv Pointer to an output stream for the CSV to write. If null, no CSV is written
+     * @param printOut Pointer to an output stream to print human-readable textual stats to. If null, no
+     *                 such stats are produced.
+     */
+    void print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader = false);
+
+    /// Print aggregated statistics for all registered counters.
+    static void printAggregate(std::ostream *csv, llvm::raw_ostream* printOut);
+
+    /////// Static nonsense ///////
+
+    // The Statistics for each input file.
+    static std::map<std::string, Statistics> stats;
+
+    // The Statistics objects for the currently-being-processed input file.
+    static Statistics* currentStatistics;
+
+    /**
+     * Aggregate statistics over all entries in `stats` and return the resulting Statistics object.
+     */
+    static Statistics getAggregate();
+
+    /**
+     * Convenient global entry point for updating the "active" Statistics. Since we operate single-threadedly
+     * processing one file at a time, this allows us to simply expose the stats for the current file globally,
+     * simplifying things.
+     */
+    static Statistics& current();
+
+    /**
+     * Set the active Statistics object to the named one, creating it if necessary, and write the completion
+     * timestamp into the currently active one.
+     */
+    static void setActive(std::string name);
+};
diff --git a/projects/hip/hipify-clang/src/Types.h b/projects/hip/hipify-clang/src/Types.h
deleted file mode 100644
index f30e4895ca..0000000000
--- a/projects/hip/hipify-clang/src/Types.h
+++ /dev/null
@@ -1,47 +0,0 @@
-#pragma once
-
-enum ConvTypes {
-    CONV_VERSION = 0,
-    CONV_INIT,
-    CONV_DEVICE,
-    CONV_MEM,
-    CONV_KERN,
-    CONV_COORD_FUNC,
-    CONV_MATH_FUNC,
-    CONV_SPECIAL_FUNC,
-    CONV_STREAM,
-    CONV_EVENT,
-    CONV_OCCUPANCY,
-    CONV_CONTEXT,
-    CONV_PEER,
-    CONV_MODULE,
-    CONV_CACHE,
-    CONV_EXEC,
-    CONV_ERROR,
-    CONV_DEF,
-    CONV_TEX,
-    CONV_GL,
-    CONV_GRAPHICS,
-    CONV_SURFACE,
-    CONV_JIT,
-    CONV_D3D9,
-    CONV_D3D10,
-    CONV_D3D11,
-    CONV_VDPAU,
-    CONV_EGL,
-    CONV_THREAD,
-    CONV_OTHER,
-    CONV_INCLUDE,
-    CONV_INCLUDE_CUDA_MAIN_H,
-    CONV_TYPE,
-    CONV_LITERAL,
-    CONV_NUMERIC_LITERAL,
-    CONV_LAST
-};
-
-enum ApiTypes {
-    API_DRIVER = 0,
-    API_RUNTIME,
-    API_BLAS,
-    API_LAST
-};

From 54f786583b6dfa9170724f92fcbd219ec59967e1 Mon Sep 17 00:00:00 2001
From: Chris Kitching <chriskitching@linux.com>
Date: Thu, 26 Oct 2017 04:30:38 +0100
Subject: [PATCH 27/28] 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: 20871a3a0744e9221b83a8def50232fe822bd316]
---
 projects/hip/hipify-clang/src/Cuda2Hip.cpp | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/projects/hip/hipify-clang/src/Cuda2Hip.cpp b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
index ae83d7f794..a1cf80fde9 100644
--- a/projects/hip/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/hip/hipify-clang/src/Cuda2Hip.cpp
@@ -166,10 +166,8 @@ protected:
           FullSourceLoc fullSL(sl, SM);
           insertReplacement(Rep, fullSL);
         }
-      } else {
-        // std::string msg = "the following reference is not handled: '" + name.str() + "' [string literal].";
-        // printHipifyMessage(SM, start, msg);
       }
+
       if (end == StringRef::npos) {
         break;
       }

From 27c5e94c81eb2ee7a9d6331bb56b9fc49a51f566 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Fri, 27 Oct 2017 23:31:43 +0300
Subject: [PATCH 28/28] [HIPIFY] fix typo - missing `)`

[ROCm/hip commit: 44c74b651193abce3dc659a05a9fcb0c6d5efa06]
---
 projects/hip/hipify-clang/src/LLVMCompat.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp
index ef9bdc193c..474ba2a7dd 100644
--- a/projects/hip/hipify-clang/src/LLVMCompat.cpp
+++ b/projects/hip/hipify-clang/src/LLVMCompat.cpp
@@ -34,7 +34,7 @@ void insertReplacement(ct::Replacements& replacements, const ct::Replacement& re
 
 void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *start, size_t len, bool DisableMacroExpansion) {
 #if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
-    _pp.EnterTokenStream(start, len, false, DisableMacroExpansion;
+    _pp.EnterTokenStream(start, len, false, DisableMacroExpansion);
 #else
     _pp.EnterTokenStream(clang::ArrayRef<clang::Token>{start, len}, DisableMacroExpansion);
 #endif