diff --git a/CMakeLists.txt b/CMakeLists.txt
index 4c89c93668..1c8f640afb 100644
--- a/CMakeLists.txt
+++ b/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
 #############################
@@ -120,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})
@@ -142,7 +140,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")
@@ -287,23 +287,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
diff --git a/hipify-clang/CMakeLists.txt b/hipify-clang/CMakeLists.txt
index cb8354157b..da6eaeaa99 100644
--- a/hipify-clang/CMakeLists.txt
+++ b/hipify-clang/CMakeLists.txt
@@ -1,31 +1,17 @@
 cmake_minimum_required(VERSION 2.8.8)
 project(hipify-clang)
 
-set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "")
+find_package(LLVM REQUIRED)
+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}")
 
-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()
-
-option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON)
-
-# 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()
+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})
@@ -81,6 +67,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)
@@ -102,5 +99,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 "")
diff --git a/hipify-clang/README.md b/hipify-clang/README.md
index 20456f3bff..563440e709 100644
--- a/hipify-clang/README.md
+++ b/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,68 +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++.
-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**.
+### 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
 
diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp
index fb7e920f8c..de6ddb2d74 100644
--- a/hipify-clang/src/CUDA2HipMap.cpp
+++ b/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/hipify-clang/src/CUDA2HipMap.h b/hipify-clang/src/CUDA2HipMap.h
index 95f9054576..605acf7aac 100644
--- a/hipify-clang/src/CUDA2HipMap.h
+++ b/hipify-clang/src/CUDA2HipMap.h
@@ -4,20 +4,9 @@
 #include <set>
 #include <map>
 
-#include "Types.h"
+#include "Statistics.h"
 
-// TODO: This shouldn't really be here. More restructuring needed...
-struct hipCounter {
-    llvm::StringRef hipName;
-    ConvTypes countType;
-    ApiTypes countApiType;
-    int unsupported;
-};
-
-#define HIP_UNSUPPORTED -1
-
-/// Macros to ignore.
-extern const std::set<llvm::StringRef> CUDA_EXCLUDES;
+#define HIP_UNSUPPORTED true
 
 /// Maps cuda header names to hip header names.
 extern const std::map<llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP;
diff --git a/hipify-clang/src/Cuda2Hip.cpp b/hipify-clang/src/Cuda2Hip.cpp
index 4bffcdaed9..a1cf80fde9 100644
--- a/hipify-clang/src/Cuda2Hip.cpp
+++ b/hipify-clang/src/Cuda2Hip.cpp
@@ -53,7 +53,8 @@ THE SOFTWARE.
 #include <sstream>
 
 #include "CUDA2HipMap.h"
-#include "Types.h"
+#include "LLVMCompat.h"
+#include "StringUtils.h"
 
 using namespace clang;
 using namespace clang::ast_matchers;
@@ -62,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");
@@ -113,41 +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;
-
-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) :
+  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,
@@ -163,26 +123,21 @@ 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());
+      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");
@@ -195,45 +150,6 @@ protected:
     llvm::errs() << "[HIPIFY] " << getMsgType(msgType) << ": " << mainFileName << ":" << fullSL.getExpansionLineNumber() << ":" << fullSL.getExpansionColumnNumber() << ": " << message << "\n";
   }
 
-  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;
-    if (counter.unsupported) {
-      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++;
-    }
-  }
-
-  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]++;
-      countApiRepsTotalUnsupported[counter.countApiType]++;
-    } else {
-      countReps[counter.countType]++;
-      countRepsTotal[counter.countType]++;
-      countApiReps[counter.countApiType]++;
-      countApiRepsTotal[counter.countApiType]++;
-    }
-  }
-
   void processString(StringRef s, SourceManager &SM, SourceLocation start) {
     size_t begin = 0;
     while ((begin = s.find("cu", begin)) != StringRef::npos) {
@@ -242,18 +158,16 @@ 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);
           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;
       }
@@ -266,7 +180,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
@@ -291,156 +205,126 @@ 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;
+    }
+
+    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());
+      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);
+  }
+
+  /**
+   * 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;
+    }
+    Statistics::current().incrementCounter(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()) &&
-        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()) {
+      RewriteToken(T);
     }
   }
 
   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 LLVM_VERSION_MAJOR > 4
-        #define GET_NUM_ARGS() getNumParams()
-#else
-        #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;
-#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
-        _pp->EnterTokenStream(start, len, false, false);
-#else
-        _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();
-            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";
-            }
-          } 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());
-              }
-            }
-          }
-        }
+
+    if (!_sm->isWrittenInMainFile(MacroNameTok.getLocation())) {
+      return; // Macros in headers are not our concern.
+    }
+
+    // Is the macro itself a CUDA identifier? If so, rewrite it
+    RewriteToken(MacroNameTok);
+
+    // 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'
+      // to workaround the 'const' MacroArgs passed into this hook.
+      const Token *start = Args->getUnexpArgument(i);
+      size_t len = Args->getArgLength(start) + 1;
+      llcompat::EnterPreprocessorTokenStream(*_pp, start, len, false);
+
+      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) {
+        RewriteToken(tok);
       }
     }
   }
 
   void EndOfMainFile() override {}
 
-  bool SeenEnd = false;
   void setSourceManager(SourceManager *sm) { _sm = sm; }
   void setPreprocessor(Preprocessor *pp) { _pp = pp; }
   void setMatch(Cuda2HipCallback *match) { Match = match; }
@@ -453,28 +337,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) {
@@ -495,7 +357,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.
@@ -503,108 +365,108 @@ 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);
-    }
+    Replacement Rep(*SM, sl, length, hipCtr.hipName);
+    FullSourceLoc fullSL(sl, *SM);
+    insertReplacement(Rep, fullSL);
 
     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)) {
       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 << ",";
-      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)) {
-          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,";
-        }
+      SourceManager *SM = Result.SourceManager;
+
+      const Expr& calleeExpr = *(launchKernel->getCallee());
+      OS << "hipLaunchKernelGGL(" << readSourceText(*SM, calleeExpr.getSourceRange()) << ", ";
+
+      // Next up are the four kernel configuration parameters, the last two of which are optional and default to zero.
+      const CallExpr& config = *(launchKernel->getConfig());
+
+      // 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());
+      hipCounter counter = {"hipLaunchKernelGGL", ConvTypes::CONV_KERN, ApiTypes::API_RUNTIME};
+      Statistics::current().incrementCounter(counter, refName.str());
       return true;
     }
     return false;
@@ -628,7 +490,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);
@@ -655,7 +517,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);
@@ -751,8 +613,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;
@@ -760,15 +622,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) {
@@ -782,7 +635,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);
   }
@@ -795,7 +648,6 @@ public:
     if (cudaLaunchKernel(Result)) return;
     if (cudaSharedIncompleteArrayVar(Result)) return;
     if (stringLiteral(Result)) return;
-    if (unresolvedTemplateName(Result)) return;
   }
 
 private:
@@ -814,7 +666,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"),
@@ -875,255 +735,6 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac
   );
 }
 
-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;
-    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";
-    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";
-    }
-    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";
-    }
-    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";
-  }
-  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;
-      }
-      llvm::outs() << "  " << counterNames[i] << ": " << sum_interm << "\n";
-      csv << counterNames[i] << separator << sum_interm << "\n";
-    }
-    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";
-    }
-    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) {
-      llvm::outs() << "  " << it.first << ": " << it.second << "\n";
-      csv << it.first << separator << it.second << "\n";
-    }
-  }
-  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;
-      }
-      llvm::outs() << "  " << counterNames[i] << ": " << sum_interm << "\n";
-      csv << counterNames[i] << separator << sum_interm << "\n";
-    }
-    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";
-    }
-    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) {
-      llvm::outs() << "  " << it.first << ": " << it.second << "\n";
-      csv << it.first << separator << it.second << "\n";
-    }
-  }
-  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;
-    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";
-    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";
-    }
-    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";
-    }
-    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";
-  }
-  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;
-      }
-      llvm::outs() << "  " << counterNames[i] << ": " << sum_interm << "\n";
-      csv << counterNames[i] << separator << sum_interm << "\n";
-    }
-    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";
-    }
-    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";
-    }
-  }
-  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;
-      }
-      llvm::outs() << "  " << counterNames[i] << ": " << sum_interm << "\n";
-      csv << counterNames[i] << separator << sum_interm << "\n";
-    }
-    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";
-    }
-    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";
-    }
-  }
-  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);
@@ -1131,16 +742,7 @@ 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;
-
-  // 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();
@@ -1149,6 +751,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";
@@ -1159,24 +762,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) {
@@ -1196,6 +798,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.
@@ -1203,15 +808,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);
 
@@ -1235,27 +832,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");
-#if LLVM_VERSION_MAJOR > 3
-        Replacements& replacements = Tool.getReplacements().begin()->second;
-#else
-        Replacements& replacements = Tool.getReplacements();
-#endif
-      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");
@@ -1268,26 +846,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/hipify-clang/src/LLVMCompat.cpp b/hipify-clang/src/LLVMCompat.cpp
new file mode 100644
index 0000000000..474ba2a7dd
--- /dev/null
+++ b/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/hipify-clang/src/LLVMCompat.h b/hipify-clang/src/LLVMCompat.h
new file mode 100644
index 0000000000..3e2fe1aebb
--- /dev/null
+++ b/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
diff --git a/hipify-clang/src/Statistics.cpp b/hipify-clang/src/Statistics.cpp
new file mode 100644
index 0000000000..00b8a66f11
--- /dev/null
+++ b/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/hipify-clang/src/Statistics.h b/hipify-clang/src/Statistics.h
new file mode 100644
index 0000000000..da4e296db0
--- /dev/null
+++ b/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/hipify-clang/src/StringUtils.cpp b/hipify-clang/src/StringUtils.cpp
new file mode 100644
index 0000000000..ad55333bc8
--- /dev/null
+++ b/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/hipify-clang/src/StringUtils.h b/hipify-clang/src/StringUtils.h
new file mode 100644
index 0000000000..66a9be780f
--- /dev/null
+++ b/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);
diff --git a/hipify-clang/src/Types.h b/hipify-clang/src/Types.h
deleted file mode 100644
index f30e4895ca..0000000000
--- a/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
-};
diff --git a/tests/hipify-clang/axpy.cu b/tests/hipify-clang/axpy.cu
index 8c6b0e0d8d..2fd62ac344 100644
--- a/tests/hipify-clang/axpy.cu
+++ b/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();