From 4c6174616277cefba5caeabedd8b3f8195d9fe54 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Mon, 12 Dec 2016 10:16:58 -0600
Subject: [PATCH 01/16] Changed threadfences to match target parsing by hcc

Change-Id: I28fcabbaacd13495b707f263fd09afaead0665fa


[ROCm/clr commit: 812cf130b64de3a2cbf18cecc8b4995bc6459393]
---
 projects/clr/hipamd/src/hip_ir.ll | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_ir.ll b/projects/clr/hipamd/src/hip_ir.ll
index d0e2a879a3..21123dd7c0 100644
--- a/projects/clr/hipamd/src/hip_ir.ll
+++ b/projects/clr/hipamd/src/hip_ir.ll
@@ -2,12 +2,12 @@ target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:
 target triple = "amdgcn--amdhsa"
 
 
-define void @__threadfence() #1 {
+define linkonce_odr spir_func void @__threadfence() #1 {
     fence syncscope(2) seq_cst
     ret void
 }
 
-define void @__threadfence_block()  #1 {
+define linkonce_odr spir_func void @__threadfence_block()  #1 {
     fence syncscope(3) seq_cst
     ret void
 }

From c9aa0841d82d94b5fb8f579786c23e55218fc93b Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Mon, 12 Dec 2016 19:57:19 +0300
Subject: [PATCH 02/16] [HIPIFY] Multiple source files support.

[ROCm/clr commit: d5eb7f494f35d6d28963f9d67e83f575170b8e08]
---
 .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp  | 141 ++++++++++--------
 1 file changed, 75 insertions(+), 66 deletions(-)

diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp
index c66a2b1b87..31da2537cb 100644
--- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp
@@ -2106,11 +2106,12 @@ void printStats(std::string fileSource, HipifyPPCallbacks &PPCallbacks, Cuda2Hip
 
 int main(int argc, const char **argv) {
   llvm::sys::PrintStackTraceOnErrorSignal();
-  CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::Required);
+  CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore);
   std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
   std::string dst = OutputFilename;
-  if (N) {
-    NoOutput = PrintStats = true;
+  if (!dst.empty() && fileSources.size() > 1) {
+    llvm::errs() << "Conflict: -o and multiple source files are specified.\n";
+    return 1;
   }
   if (NoOutput) {
     if (Inplace) {
@@ -2122,84 +2123,92 @@ int main(int argc, const char **argv) {
       return 1;
     }
   }
-  if (dst.empty()) {
-    dst = fileSources[0];
-    if (!Inplace) {
-      size_t pos = dst.rfind(".");
-      if (pos != std::string::npos && pos+1 < dst.size()) {
-        dst = dst.substr(0, pos) + ".hip." + dst.substr(pos+1, dst.size()-pos-1);
-      } else {
-        dst += ".hip.cu";
+  if (N) {
+    NoOutput = PrintStats = true;
+  }
+  int Result = 0;
+  for (const auto & src : fileSources) {
+    if (dst.empty()) {
+      dst = src;
+      if (!Inplace) {
+        size_t pos = dst.rfind(".");
+        if (pos != std::string::npos && pos + 1 < dst.size()) {
+          dst = dst.substr(0, pos) + ".hip." + dst.substr(pos + 1, dst.size() - pos - 1);
+        }
+        else {
+          dst += ".hip.cu";
+        }
       }
     }
-  } else {
-    if (Inplace) {
-      llvm::errs() << "Conflict: both -o and -inplace options are specified.\n";
-      return 1;
+    else {
+      if (Inplace) {
+        llvm::errs() << "Conflict: both -o and -inplace options are specified.\n";
+        return 1;
+      }
+      dst += ".hip";
+    }
+    // backup source file since tooling may change "inplace"
+    if (!NoBackup || !Inplace) {
+      std::ifstream source(src, std::ios::binary);
+      std::ofstream dest(Inplace ? dst + ".prehip" : dst, std::ios::binary);
+      dest << source.rdbuf();
+      source.close();
+      dest.close();
     }
-    dst += ".hip";
-  }
-  // backup source file since tooling may change "inplace"
-  if (!NoBackup || !Inplace) {
-    std::ifstream source(fileSources[0], std::ios::binary);
-    std::ofstream dest(Inplace ? dst + ".prehip" : dst, std::ios::binary);
-    dest << source.rdbuf();
-    source.close();
-    dest.close();
-  }
 
-  RefactoringTool Tool(OptionsParser.getCompilations(), dst);
-  ast_matchers::MatchFinder Finder;
-  HipifyPPCallbacks PPCallbacks(&Tool.getReplacements());
-  Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder, &PPCallbacks);
+    RefactoringTool Tool(OptionsParser.getCompilations(), dst);
+    ast_matchers::MatchFinder Finder;
+    HipifyPPCallbacks PPCallbacks(&Tool.getReplacements());
+    Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder, &PPCallbacks);
 
-  addAllMatchers(Finder, &Callback);
+    addAllMatchers(Finder, &Callback);
 
-  auto action = newFrontendActionFactory(&Finder, &PPCallbacks);
-  std::vector<const char*> compilationStages;
-  compilationStages.push_back("--cuda-host-only");
-  Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster(compilationStages[0], ArgumentInsertPosition::BEGIN));
-  Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-std=c++11"));
+    auto action = newFrontendActionFactory(&Finder, &PPCallbacks);
+    std::vector<const char*> compilationStages;
+    compilationStages.push_back("--cuda-host-only");
+    Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster(compilationStages[0], ArgumentInsertPosition::BEGIN));
+    Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-std=c++11"));
 #if defined(HIPIFY_CLANG_RES)
-  Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
+    Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
 #endif
-  Tool.appendArgumentsAdjuster(getClangSyntaxOnlyAdjuster());
-  int Result = Tool.run(action.get());
-  Tool.clearArgumentsAdjusters();
+    Tool.appendArgumentsAdjuster(getClangSyntaxOnlyAdjuster());
+    Result += Tool.run(action.get());
+    Tool.clearArgumentsAdjusters();
 
-  LangOptions DefaultLangOptions;
-  IntrusiveRefCntPtr<DiagnosticOptions> DiagOpts = new DiagnosticOptions();
-  TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts);
-  DiagnosticsEngine Diagnostics(
+    LangOptions DefaultLangOptions;
+    IntrusiveRefCntPtr<DiagnosticOptions> DiagOpts = new DiagnosticOptions();
+    TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts);
+    DiagnosticsEngine Diagnostics(
       IntrusiveRefCntPtr<DiagnosticIDs>(new DiagnosticIDs()), &*DiagOpts,
       &DiagnosticPrinter, false);
 
-  DEBUG(dbgs() << "Replacements collected by the tool:\n");
-  for (const auto &r : Tool.getReplacements()) {
-    DEBUG(dbgs() << r.toString() << "\n");
-  }
+    DEBUG(dbgs() << "Replacements collected by the tool:\n");
+    for (const auto &r : Tool.getReplacements()) {
+      DEBUG(dbgs() << r.toString() << "\n");
+    }
 
-  SourceManager Sources(Diagnostics, Tool.getFiles());
-  Rewriter Rewrite(Sources, DefaultLangOptions);
+    SourceManager Sources(Diagnostics, Tool.getFiles());
+    Rewriter Rewrite(Sources, DefaultLangOptions);
 
-  if (!Tool.applyAllReplacements(Rewrite)) {
-    DEBUG(dbgs() << "Skipped some replacements.\n");
-  }
-  if (!NoOutput) {
-    Result = Rewrite.overwriteChangedFiles();
-  }
-  if (!Inplace && !NoOutput) {
-    size_t pos = dst.rfind(".");
-    if (pos != std::string::npos) {
-      rename(dst.c_str(), dst.substr(0, pos).c_str());
+    if (!Tool.applyAllReplacements(Rewrite)) {
+      DEBUG(dbgs() << "Skipped some replacements.\n");
+    }
+    if (!NoOutput) {
+      Result += Rewrite.overwriteChangedFiles();
+    }
+    if (!Inplace && !NoOutput) {
+      size_t pos = dst.rfind(".");
+      if (pos != std::string::npos) {
+        rename(dst.c_str(), dst.substr(0, pos).c_str());
+      }
+    }
+    if (NoOutput) {
+      remove(dst.c_str());
+    }
+    dst.clear();
+    if (PrintStats) {
+      printStats(src, PPCallbacks, Callback);
     }
   }
-  if (NoOutput) {
-    remove(dst.c_str());
-  }
-  if (PrintStats) {
-    printStats(fileSources[0], PPCallbacks, Callback);
-  }
-
   return Result;
 }

From 1da43c7216f63c425078b7559afefe8014dfe4c0 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Mon, 12 Dec 2016 20:03:01 +0300
Subject: [PATCH 03/16] [HIPIFY] Rename -n to -examine

[ROCm/clr commit: b73ffa4e7c48f0df2e3d4c7b0fedd40c0fbfec58]
---
 projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp
index 31da2537cb..b4e364c52f 100644
--- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp
@@ -2019,7 +2019,7 @@ static cl::opt<bool> PrintStats("print-stats",
        cl::value_desc("print-stats"),
        cl::cat(ToolTemplateCategory));
 
-static cl::opt<bool> N("n",
+static cl::opt<bool> Examine("examine",
        cl::desc("Combines -no-output and -print-stats options"),
        cl::value_desc("n"),
        cl::cat(ToolTemplateCategory));
@@ -2123,7 +2123,7 @@ int main(int argc, const char **argv) {
       return 1;
     }
   }
-  if (N) {
+  if (Examine) {
     NoOutput = PrintStats = true;
   }
   int Result = 0;

From e22ed56bd79d489e1fc6f165270bbc3bc6f19735 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Tue, 13 Dec 2016 18:01:08 +0300
Subject: [PATCH 04/16] [HIPIFY] Code refactoring and total stat collecting

+ Single base class for Preprocessor and MatchFinder classes.
+ Total Stats for multiple files is added.


[ROCm/clr commit: 6bcacf83a41f96617eb30ad4b068583e3eced935]
---
 .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp  | 187 ++++++++++--------
 1 file changed, 103 insertions(+), 84 deletions(-)

diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp
index b4e364c52f..4f37bb9365 100644
--- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp
+++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp
@@ -85,6 +85,7 @@ const char *counterNames[CONV_LAST] = {
     "special_func", "stream",  "event", "ctx",   "module",     "cache",
     "err",          "def",     "tex",   "other", "include",    "include_cuda_main_header",
     "type",         "literal", "numeric_literal"};
+
 enum ApiTypes {
   API_DRIVER = 0,
   API_RUNTIME,
@@ -97,6 +98,15 @@ const char *apiNames[API_LAST] = {
 
 namespace {
 
+int64_t countRepsTotal[CONV_LAST] = { 0 };
+int64_t countApiRepsTotal[API_LAST] = { 0 };
+
+struct hipCounter {
+  StringRef hipName;
+  ConvTypes countType;
+  ApiTypes countApiType;
+};
+
 struct cuda2hipMap {
   cuda2hipMap() {
 
@@ -1317,13 +1327,7 @@ struct cuda2hipMap {
     //cuda2hipRename["cublasDrotmg_v2"] = {"hipblasDrotmg", CONV_MATH_FUNC, API_BLAS};
   }
 
-  struct HipNames {
-    StringRef hipName;
-    ConvTypes countType;
-    ApiTypes countApiType;
-  };
-
-  SmallDenseMap<StringRef, HipNames> cuda2hipRename;
+  SmallDenseMap<StringRef, hipCounter> cuda2hipRename;
   std::set<StringRef> cudaExcludes;
 };
 
@@ -1333,36 +1337,52 @@ StringRef unquoteStr(StringRef s) {
   return s;
 }
 
-static void processString(StringRef s, const cuda2hipMap &map,
-                          Replacements *Replace, SourceManager &SM,
-                          SourceLocation start,
-                          int64_t countReps[CONV_LAST],
-                          int64_t countApiReps[API_LAST]) {
-  size_t begin = 0;
-  while ((begin = s.find("cu", begin)) != StringRef::npos) {
-    const size_t end = s.find_first_of(" ", begin + 4);
-    StringRef name = s.slice(begin, end);
-    const auto found = map.cuda2hipRename.find(name);
-    if (found != map.cuda2hipRename.end()) {
-      StringRef repName = found->second.hipName;
-      countReps[CONV_LITERAL]++;
-      countApiReps[API_RUNTIME]++;
-      SourceLocation sl = start.getLocWithOffset(begin + 1);
-      Replacement Rep(SM, sl, name.size(), repName);
-      Replace->insert(Rep);
-    }
-    if (end == StringRef::npos)
-      break;
-    begin = end + 1;
+class Cuda2Hip {
+public:
+  Cuda2Hip(Replacements *R): Replace(R) {}
+
+  int64_t countReps[CONV_LAST] = { 0 };
+  int64_t countApiReps[API_LAST] = { 0 };
+
+protected:
+  struct cuda2hipMap N;
+  Replacements *Replace;
+
+  virtual void updateCounters(const hipCounter & counter) {
+    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) {
+      const size_t end = s.find_first_of(" ", begin + 4);
+      StringRef name = s.slice(begin, end);
+      const auto found = N.cuda2hipRename.find(name);
+      if (found != N.cuda2hipRename.end()) {
+        StringRef repName = found->second.hipName;
+        hipCounter counter = { "", CONV_LITERAL, API_RUNTIME };
+        updateCounters(counter);
+        SourceLocation sl = start.getLocWithOffset(begin + 1);
+        Replacement Rep(SM, sl, name.size(), repName);
+        Replace->insert(Rep);
+      }
+      if (end == StringRef::npos) {
+        break;
+      }
+      begin = end + 1;
+    }
+  }
+};
 
 class Cuda2HipCallback;
 
-class HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks {
+class HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks, public Cuda2Hip {
 public:
   HipifyPPCallbacks(Replacements *R)
-      : SeenEnd(false), _sm(nullptr), _pp(nullptr), Replace(R) {}
+    : Cuda2Hip(R), SeenEnd(false), _sm(nullptr), _pp(nullptr) {}
 
   virtual bool handleBeginSource(CompilerInstance &CI,
                                  StringRef Filename) override {
@@ -1389,8 +1409,7 @@ public:
         const auto found = N.cuda2hipRename.find(file_name);
         if (found != N.cuda2hipRename.end()) {
           StringRef repName = found->second.hipName;
-          countReps[found->second.countType]++;
-          countApiReps[found->second.countApiType]++;
+          updateCounters(found->second);
           DEBUG(dbgs() << "Include file found: " << file_name << "\n"
                        << "SourceLocation:"
                        << filename_range.getBegin().printToString(*_sm) << "\n"
@@ -1418,8 +1437,7 @@ public:
           const auto found = N.cuda2hipRename.find(name);
           if (found != N.cuda2hipRename.end()) {
             StringRef repName = found->second.hipName;
-            countReps[found->second.countType]++;
-            countApiReps[found->second.countApiType]++;
+            updateCounters(found->second);
             SourceLocation sl = T.getLocation();
             DEBUG(dbgs() << "Identifier " << name
                          << " found in definition of macro "
@@ -1465,8 +1483,7 @@ public:
               const auto found = N.cuda2hipRename.find(name);
               if (found != N.cuda2hipRename.end()) {
                 StringRef repName = found->second.hipName;
-                countReps[found->second.countType]++;
-                countApiReps[found->second.countApiType]++;
+                updateCounters(found->second);
                 DEBUG(dbgs()
                       << "Identifier " << name
                       << " found as an actual argument in expansion of macro "
@@ -1488,16 +1505,14 @@ public:
                 if (found != N.cuda2hipRename.end()) {
                   sl = sl_macro;
                   StringRef repName = found->second.hipName;
-                  countReps[found->second.countType]++;
-                  countApiReps[found->second.countApiType]++;
+                  updateCounters(found->second);
                   Replacement Rep(*_sm, sl, length, repName);
                   Replace->insert(Rep);
                 }
               } else {
                 if (tok.is(tok::string_literal)) {
                   StringRef s(tok.getLiteralData(), tok.getLength());
-                  processString(unquoteStr(s), N, Replace, *_sm, tok.getLocation(),
-                    countReps, countApiReps);
+                  processString(unquoteStr(s), *_sm, tok.getLocation());
                 }
               }
             }
@@ -1513,18 +1528,14 @@ public:
   void setSourceManager(SourceManager *sm) { _sm = sm; }
   void setPreprocessor(Preprocessor *pp) { _pp = pp; }
   void setMatch(Cuda2HipCallback *match) { Match = match; }
-  int64_t countReps[CONV_LAST]  = { 0 };
-  int64_t countApiReps[API_LAST] = { 0 };
 
 private:
   SourceManager *_sm;
   Preprocessor *_pp;
   Cuda2HipCallback *Match;
-  Replacements *Replace;
-  struct cuda2hipMap N;
 };
 
-class Cuda2HipCallback : public MatchFinder::MatchCallback {
+class Cuda2HipCallback : public MatchFinder::MatchCallback, public Cuda2Hip {
 private:
   void convertKernelDecl(const FunctionDecl *kernelDecl, const MatchFinder::MatchResult &Result) {
     SourceManager *SM = Result.SourceManager;
@@ -1579,8 +1590,7 @@ private:
           }
         }
         if (bReplace) {
-          countReps[found->second.countType]++;
-          countApiReps[found->second.countApiType]++;
+          updateCounters(found->second);
           Replacement Rep(*SM, sl, length, repName);
           Replace->insert(Rep);
         }
@@ -1653,8 +1663,8 @@ private:
                         SM->getCharacterData(launchKernel->getLocStart());
       Replacement Rep(*SM, launchKernel->getLocStart(), length, OS.str());
       Replace->insert(Rep);
-      countReps[CONV_KERN]++;
-      countApiReps[API_RUNTIME]++;
+      hipCounter counter = { "", CONV_KERN, API_RUNTIME };
+      updateCounters(counter);
       return true;
     }
     return false;
@@ -1675,8 +1685,7 @@ private:
           const auto found = N.cuda2hipRename.find(name);
           if (found != N.cuda2hipRename.end()) {
             StringRef repName = found->second.hipName;
-            countReps[found->second.countType]++;
-            countApiReps[found->second.countApiType]++;
+            updateCounters(found->second);
             SourceLocation sl = threadIdx->getLocStart();
             SourceManager *SM = Result.SourceManager;
             Replacement Rep(*SM, sl, name.size(), repName);
@@ -1695,8 +1704,7 @@ private:
       const auto found = N.cuda2hipRename.find(name);
       if (found != N.cuda2hipRename.end()) {
         StringRef repName = found->second.hipName;
-        countReps[found->second.countType]++;
-        countApiReps[found->second.countApiType]++;
+        updateCounters(found->second);
         SourceLocation sl = enumConstantRef->getLocStart();
         SourceManager *SM = Result.SourceManager;
         Replacement Rep(*SM, sl, name.size(), repName);
@@ -1719,8 +1727,7 @@ private:
       const auto found = N.cuda2hipRename.find(name);
       if (found != N.cuda2hipRename.end()) {
         StringRef repName = found->second.hipName;
-        countReps[found->second.countType]++;
-        countApiReps[found->second.countApiType]++;
+        updateCounters(found->second);
         SourceLocation sl = enumConstantDecl->getLocStart();
         SourceManager *SM = Result.SourceManager;
         Replacement Rep(*SM, sl, name.size(), repName);
@@ -1742,8 +1749,7 @@ private:
       const auto found = N.cuda2hipRename.find(name);
       if (found != N.cuda2hipRename.end()) {
         StringRef repName = found->second.hipName;
-        countReps[found->second.countType]++;
-        countApiReps[found->second.countApiType]++;
+        updateCounters(found->second);
         SourceLocation sl = typedefVar->getLocStart();
         SourceManager *SM = Result.SourceManager;
         Replacement Rep(*SM, sl, name.size(), repName);
@@ -1763,8 +1769,7 @@ private:
       const auto found = N.cuda2hipRename.find(name);
       if (found != N.cuda2hipRename.end()) {
         StringRef repName = found->second.hipName;
-        countReps[found->second.countType]++;
-        countApiReps[found->second.countApiType]++;
+        updateCounters(found->second);
         TypeLoc TL = structVar->getTypeSourceInfo()->getTypeLoc();
         SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
         SourceManager *SM = Result.SourceManager;
@@ -1784,8 +1789,7 @@ private:
         const auto found = N.cuda2hipRename.find(name);
         if (found != N.cuda2hipRename.end()) {
           StringRef repName = found->second.hipName;
-          countReps[found->second.countType]++;
-          countApiReps[found->second.countApiType]++;
+          updateCounters(found->second);
           TypeLoc TL = structVarPtr->getTypeSourceInfo()->getTypeLoc();
           SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
           SourceManager *SM = Result.SourceManager;
@@ -1807,8 +1811,7 @@ private:
       const auto found = N.cuda2hipRename.find(name);
       if (found != N.cuda2hipRename.end()) {
         StringRef repName = found->second.hipName;
-        countReps[found->second.countType]++;
-        countApiReps[found->second.countApiType]++;
+        updateCounters(found->second);
         TypeLoc TL = typeInfo->getTypeLoc();
         SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
         SourceManager *SM = Result.SourceManager;
@@ -1852,8 +1855,8 @@ private:
           StringRef repName = Twine("HIP_DYNAMIC_SHARED(" + typeName + ", " + varName + ")").toStringRef(tmpData);
           Replacement Rep(*SM, slStart, repLength, repName);
           Replace->insert(Rep);
-          countReps[CONV_MEM]++;
-          countApiReps[API_RUNTIME]++;
+          hipCounter counter = { "", CONV_MEM, API_RUNTIME };
+          updateCounters(counter);
         }
       }
       return true;
@@ -1872,8 +1875,7 @@ private:
       const auto found = N.cuda2hipRename.find(name);
       if (found != N.cuda2hipRename.end()) {
         StringRef repName = found->second.hipName;
-        countReps[found->second.countType]++;
-        countApiReps[found->second.countApiType]++;
+        updateCounters(found->second);
         TypeLoc TL = paramDecl->getTypeSourceInfo()->getTypeLoc();
         SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
         SourceManager *SM = Result.SourceManager;
@@ -1897,8 +1899,7 @@ private:
         const auto found = N.cuda2hipRename.find(name);
         if (found != N.cuda2hipRename.end()) {
           StringRef repName = found->second.hipName;
-          countReps[found->second.countType]++;
-          countApiReps[found->second.countApiType]++;
+          updateCounters(found->second);
           TypeLoc TL = paramDeclPtr->getTypeSourceInfo()->getTypeLoc();
           SourceLocation sl = TL.getUnqualifiedLoc().getLocStart();
           SourceManager *SM = Result.SourceManager;
@@ -1925,7 +1926,7 @@ private:
       if (sLiteral->getCharByteWidth() == 1) {
         StringRef s = sLiteral->getString();
         SourceManager *SM = Result.SourceManager;
-        processString(s, N, Replace, *SM, sLiteral->getLocStart(), countReps, countApiReps);
+        processString(s, *SM, sLiteral->getLocStart());
       }
       return true;
     }
@@ -1934,7 +1935,7 @@ private:
 
 public:
   Cuda2HipCallback(Replacements *Replace, ast_matchers::MatchFinder *parent, HipifyPPCallbacks *PPCallbacks)
-    : Replace(Replace), owner(parent), PP(PPCallbacks) {
+    : Cuda2Hip(Replace), owner(parent), PP(PPCallbacks) {
     PP->setMatch(this);
   }
 
@@ -1962,19 +1963,14 @@ public:
       SourceManager *SM = Result.SourceManager;
       Replacement Rep(*SM, SM->getLocForStartOfFile(SM->getMainFileID()), 0, repName);
       Replace->insert(Rep);
-      countReps[CONV_INCLUDE_CUDA_MAIN_H]++;
-      countApiReps[API_RUNTIME]++;
+      hipCounter counter = { "", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME };
+      updateCounters(counter);
     }
   }
 
-  int64_t countReps[CONV_LAST]  = { 0 };
-  int64_t countApiReps[API_LAST] = { 0 };
-
 private:
-  Replacements *Replace;
   ast_matchers::MatchFinder *owner;
   HipifyPPCallbacks *PP;
-  struct cuda2hipMap N;
 };
 
 void HipifyPPCallbacks::handleEndSource() {
@@ -1983,8 +1979,8 @@ void HipifyPPCallbacks::handleEndSource() {
     StringRef repName = "#include <hip/hip_runtime.h>\n";
     Replacement Rep(*_sm, _sm->getLocForStartOfFile(_sm->getMainFileID()), 0, repName);
     Replace->insert(Rep);
-    countReps[CONV_INCLUDE_CUDA_MAIN_H]++;
-    countApiReps[API_RUNTIME]++;
+    hipCounter counter = { "", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME };
+    updateCounters(counter);
   }
 }
 
@@ -2088,12 +2084,12 @@ void addAllMatchers(ast_matchers::MatchFinder &Finder, Cuda2HipCallback *Callbac
                             Callback);
 }
 
-void printStats(std::string fileSource, HipifyPPCallbacks &PPCallbacks, Cuda2HipCallback &Callback) {
+int64_t printStats(std::string fileSource, HipifyPPCallbacks &PPCallbacks, Cuda2HipCallback &Callback) {
   int64_t sum = 0;
   for (int i = 0; i < CONV_LAST; i++) {
     sum += Callback.countReps[i] + PPCallbacks.countReps[i];
   }
-  llvm::outs() << "info: converted " << sum << " CUDA->HIP refs ( ";
+  llvm::outs() << "Info: converted " << sum << " CUDA->HIP refs ( ";
   for (int i = 0; i < CONV_LAST; i++) {
     llvm::outs() << counterNames[i] << ':' << Callback.countReps[i] + PPCallbacks.countReps[i] << ' ';
   }
@@ -2102,6 +2098,23 @@ void printStats(std::string fileSource, HipifyPPCallbacks &PPCallbacks, Cuda2Hip
     llvm::outs() << apiNames[i] << ':' << Callback.countApiReps[i] + PPCallbacks.countApiReps[i] << ' ';
   }
   llvm::outs() << ") in \'" << fileSource << "\'\n";
+  return sum;
+}
+
+void printAllStats(int64_t totalFiles, int64_t convertedFiles) {
+  int64_t sum = 0;
+  for (int i = 0; i < CONV_LAST; i++) {
+    sum += countRepsTotal[i];
+  }
+  llvm::outs() << "Info: totally converted " << sum << " CUDA->HIP refs ( ";
+  for (int i = 0; i < CONV_LAST; i++) {
+    llvm::outs() << counterNames[i] << ':' << countRepsTotal[i] << ' ';
+  }
+  llvm::outs() << "), by APIs ( ";
+  for (int i = 0; i < API_LAST; i++) {
+    llvm::outs() << apiNames[i] << ':' << countApiRepsTotal[i] << ' ';
+  }
+  llvm::outs() << ") in " << convertedFiles << " files of " << totalFiles << " processed files.\n";
 }
 
 int main(int argc, const char **argv) {
@@ -2127,6 +2140,7 @@ int main(int argc, const char **argv) {
     NoOutput = PrintStats = true;
   }
   int Result = 0;
+  size_t filesTransleted = fileSources.size();
   for (const auto & src : fileSources) {
     if (dst.empty()) {
       dst = src;
@@ -2207,8 +2221,13 @@ int main(int argc, const char **argv) {
     }
     dst.clear();
     if (PrintStats) {
-      printStats(src, PPCallbacks, Callback);
+      if (0 == printStats(src, PPCallbacks, Callback)) {
+        filesTransleted--;
+      }
     }
   }
+  if (PrintStats && fileSources.size() > 1) {
+    printAllStats(fileSources.size(), filesTransleted);
+  }
   return Result;
 }

From c865d937926b831e58d156e0c9e0edc72cfafdfe Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 13 Dec 2016 09:18:34 -0600
Subject: [PATCH 05/16] added half math addition ISA support

Change-Id: I293b771f695b499b795d7e53f600c9e4fe2a2071


[ROCm/clr commit: 02eab122c5eae36891b1337393986f6eb5823667]
---
 .../clr/hipamd/include/hip/hcc_detail/hip_fp16.h | 16 ++++++++++++++++
 projects/clr/hipamd/src/hip_fp16.cpp             |  5 +++++
 projects/clr/hipamd/src/hip_ir.ll                |  5 +++++
 3 files changed, 26 insertions(+)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
index bcf2605f28..445df78eb4 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -25,6 +25,10 @@ THE SOFTWARE.
 
 #include "hip/hip_runtime.h"
 
+#define __CLANG_VERSION__ __clang_major__ * 10 + __clang_minor__
+
+#if __CLANG_VERSION__ == 35
+
 typedef struct{
   unsigned x: 16;
 } __half;
@@ -175,3 +179,15 @@ __device__ __half2 __lowhigh2highlow(const __half2 a);
 __device__ __half2 __low2half2(const __half2 a, const __half2 b);
 
 #endif
+
+#if __CLANG_VERSION__ == 40
+
+typedef __fp16 __half;
+extern "C" __half __hip_hadd_clang40_gfx803(__half a, __half b);
+
+__device__ inline __half __hadd(__half a, __half b){
+  return __hip_hadd_clang40_gfx803(a, b);
+}
+
+#endif
+#endif
diff --git a/projects/clr/hipamd/src/hip_fp16.cpp b/projects/clr/hipamd/src/hip_fp16.cpp
index 1a9d04474f..5d01f73cf7 100644
--- a/projects/clr/hipamd/src/hip_fp16.cpp
+++ b/projects/clr/hipamd/src/hip_fp16.cpp
@@ -22,6 +22,8 @@ THE SOFTWARE.
 
 #include"hip/hip_fp16.h"
 
+#if __CLANG_VERSION__ == 35
+
 static const unsigned sign_val = 0x8000;
 static const __half __half_value_one_float = {0x3C00};
 static const __half __half_value_zero_float = {0x0};
@@ -372,3 +374,6 @@ __device__ __half2 __lowhigh2highlow(const __half2 a){
 __device__ __half2 __low2half2(const __half2 a, const __half2 b){
   return {a.q, b.q};
 }
+
+#endif
+
diff --git a/projects/clr/hipamd/src/hip_ir.ll b/projects/clr/hipamd/src/hip_ir.ll
index 21123dd7c0..078dc3eed5 100644
--- a/projects/clr/hipamd/src/hip_ir.ll
+++ b/projects/clr/hipamd/src/hip_ir.ll
@@ -34,4 +34,9 @@ define linkonce_odr spir_func i32 @__rocm_hadd(i32 %in1, i32 %in2) {
   ret i32 %val
 }
 
+define linkonce_odr spir_func half @__hip_hadd_clang40_gfx803(half %a, half %b) {
+  %val = tail call half asm "v_add_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
+  ret half %val
+}
+
 attributes #1 = { alwaysinline nounwind }

From 9691299b0ddcc0292234aa22c8e13765ff9a913e Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 13 Dec 2016 14:41:36 -0600
Subject: [PATCH 06/16] added few type reinterpret cast device functions

1. __int_as_float
2. __hiloint2double

Change-Id: Id247c196887b24a12090f0521bf91e13afeec733


[ROCm/clr commit: 7a712aa76b44896d43898dfdc6b981784ef4d9de]
---
 projects/clr/hipamd/CMakeLists.txt            |  3 +-
 .../clr/hipamd/include/hip/device_functions.h | 33 +++++++++++
 .../include/hip/hcc_detail/device_functions.h | 32 ++++++++++
 .../include/hip/hcc_detail/hip_vector_types.h |  3 -
 projects/clr/hipamd/src/device_functions.cpp  | 58 +++++++++++++++++++
 5 files changed, 125 insertions(+), 4 deletions(-)
 create mode 100644 projects/clr/hipamd/include/hip/device_functions.h
 create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
 create mode 100644 projects/clr/hipamd/src/device_functions.cpp

diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt
index 6f9a819c4d..c76410f06d 100644
--- a/projects/clr/hipamd/CMakeLists.txt
+++ b/projects/clr/hipamd/CMakeLists.txt
@@ -177,7 +177,8 @@ if(HIP_PLATFORM STREQUAL "hcc")
     set(SOURCE_FILES_DEVICE
         src/device_util.cpp
         src/hip_ldg.cpp
-        src/hip_fp16.cpp)
+        src/hip_fp16.cpp
+        src/device_functions.cpp)
 
     set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic")
     add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME})
diff --git a/projects/clr/hipamd/include/hip/device_functions.h b/projects/clr/hipamd/include/hip/device_functions.h
new file mode 100644
index 0000000000..838bad8f0c
--- /dev/null
+++ b/projects/clr/hipamd/include/hip/device_functions.h
@@ -0,0 +1,33 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_DEVICE_FUNCTIONS_H
+#define HIP_DEVICE_FUNCTIONS_H
+
+#include<hip/hip_common.h>
+
+#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
+#include <hip/hcc_detail/device_functions.h>
+#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
+#include <device_functions.h>
+#else 
+#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
+#endif 
+
+#endif
diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
new file mode 100644
index 0000000000..8fa870664f
--- /dev/null
+++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
@@ -0,0 +1,32 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
+#define HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
+
+#include "hip_runtime.h"
+
+__device__ float __int_as_float (int x);
+
+__device__ double __hiloint2double (int hi, int lo);
+
+extern __HIP_DEVICE__ double  __longlong_as_double(long long int x);
+extern __HIP_DEVICE__ long long int __double_as_longlong(double x);
+
+#endif
diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h
index ffe15a27a4..7c48985996 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h
@@ -343,9 +343,6 @@ __HIP_DEVICE__ double2 make_double2(double, double );
 __HIP_DEVICE__ double3 make_double3(double, double, double );
 __HIP_DEVICE__ double4 make_double4(double, double, double, double );
 
-extern __HIP_DEVICE__ double  __longlong_as_double(long long int x);
-extern __HIP_DEVICE__ long long int __double_as_longlong(double x);
-
 
 /*
 ///---
diff --git a/projects/clr/hipamd/src/device_functions.cpp b/projects/clr/hipamd/src/device_functions.cpp
new file mode 100644
index 0000000000..30a09e6e02
--- /dev/null
+++ b/projects/clr/hipamd/src/device_functions.cpp
@@ -0,0 +1,58 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip/device_functions.h>
+
+extern "C" float __hip_int_as_float(int);
+
+typedef struct {
+  signed int hi;
+  signed int lo;
+} __hip_signed_2;
+
+typedef struct {
+union {
+  double d;
+  long long int lli;
+  __hip_signed_2 s2;
+} ;
+} __hip_64bit_struct;
+
+typedef struct {
+union {
+  float f;
+  unsigned int ui;
+  signed int si;
+};
+} __hip_32bit_struct;
+
+__device__ float __int_as_float (int x) {
+  __hip_32bit_struct s;
+  s.si = x;
+  return s.f;
+}
+
+__device__ double __hiloint2double (int hi, int lo) {
+  __hip_64bit_struct s;
+  s.s2.hi = hi;
+  s.s2.lo = lo;
+  return s.d;
+}
+
+

From 1f638c1da61b73f0122a9ce48658e22d42f40414 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Tue, 13 Dec 2016 15:07:04 -0600
Subject: [PATCH 07/16] Add USE_IPC to disable use of IPC APIs.  Set to 0.

[ROCm/clr commit: 5eed20be123db9e575efed800512328f21807716]
---
 projects/clr/hipamd/src/hip_hcc.h      |  3 +++
 projects/clr/hipamd/src/hip_memory.cpp | 13 ++++++++++++-
 2 files changed, 15 insertions(+), 1 deletion(-)

diff --git a/projects/clr/hipamd/src/hip_hcc.h b/projects/clr/hipamd/src/hip_hcc.h
index 82290bc489..b5417080e0 100644
--- a/projects/clr/hipamd/src/hip_hcc.h
+++ b/projects/clr/hipamd/src/hip_hcc.h
@@ -34,6 +34,7 @@ THE SOFTWARE.
 #endif
 
 #define USE_DISPATCH_HSA_KERNEL 1
+#define USE_IPC 0
 //
 
 
@@ -375,7 +376,9 @@ struct LockedBase {
 class ihipIpcMemHandle_t
 {
 public:
+#if USE_IPC
     hsa_amd_ipc_memory_t ipc_handle; ///< ipc memory handle on ROCr
+#endif
     char reserved[HIP_IPC_HANDLE_SIZE];
     size_t psize;
 };
diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
index f2ab6d19a0..7aa6ef4942 100644
--- a/projects/clr/hipamd/src/hip_memory.cpp
+++ b/projects/clr/hipamd/src/hip_memory.cpp
@@ -1056,11 +1056,15 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
     // Save the size of the pointer to hipIpcMemHandle
     (*handle)->psize = psize;
 
+#if USE_IPC
     // Create HSA ipc memory
     hsa_status_t hsa_status =
         hsa_amd_ipc_memory_create(devPtr, psize, &(*handle)->ipc_handle);
     if(hsa_status!= HSA_STATUS_SUCCESS)
         hipStatus = hipErrorMemoryAllocation;
+#else
+    hipStatus = hipErrorRuntimeOther;
+#endif
 
     return hipStatus;
 }
@@ -1069,6 +1073,7 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
 // HIP_INIT_API ( devPtr, handle.handle , flags);
     hipError_t hipStatus = hipSuccess;
 
+#if USE_IPC
     // Get the current device agent.
     hc::accelerator acc;
     hsa_agent_t *agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
@@ -1080,7 +1085,9 @@ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned
         hsa_amd_ipc_memory_attach(&handle->ipc_handle, handle->psize, 1, agent, devPtr);
     if(hsa_status != HSA_STATUS_SUCCESS)
         hipStatus = hipErrorMapBufferObjectFailed;
-
+#else 
+    hipStatus = hipErrorRuntimeOther;
+#endif
     return hipStatus;
 }
 
@@ -1088,10 +1095,14 @@ hipError_t hipIpcCloseMemHandle(void *devPtr){
     HIP_INIT_API ( devPtr );
     hipError_t hipStatus = hipSuccess;
 
+#if USE_IPC
     hsa_status_t hsa_status =
         hsa_amd_ipc_memory_detach(devPtr);
     if(hsa_status != HSA_STATUS_SUCCESS)
         return hipErrorInvalidResourceHandle;
+#else
+    hipStatus = hipErrorRuntimeOther;
+#endif
     return hipStatus;
 }
 

From b26857a631db974b85d84637d4b20684b19ac47c Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 13 Dec 2016 20:06:56 -0600
Subject: [PATCH 08/16] disabled compiler flag hcc 4.0 for half support

Change-Id: I32175113f4c05d43310b3a05c2a14e12f6d48b09


[ROCm/clr commit: 26934a920cb4ea88c939620b87fece420ef36f16]
---
 .../hipamd/include/hip/hcc_detail/hip_fp16.h  | 22 +++++++++----------
 projects/clr/hipamd/src/hip_fp16.cpp          |  3 ---
 2 files changed, 10 insertions(+), 15 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
index 445df78eb4..735f915bd2 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -27,7 +27,16 @@ THE SOFTWARE.
 
 #define __CLANG_VERSION__ __clang_major__ * 10 + __clang_minor__
 
-#if __CLANG_VERSION__ == 35
+#ifdef HIP_HALF_HW_SUPPORT
+
+typedef __fp16 __half;
+extern "C" __half __hip_hadd_clang40_gfx803(__half a, __half b);
+
+__device__ inline __half __hadd(__half a, __half b){
+  return __hip_hadd_clang40_gfx803(a, b);
+}
+
+#else
 
 typedef struct{
   unsigned x: 16;
@@ -178,16 +187,5 @@ __device__ __half2 __lowhigh2highlow(const __half2 a);
 
 __device__ __half2 __low2half2(const __half2 a, const __half2 b);
 
-#endif
-
-#if __CLANG_VERSION__ == 40
-
-typedef __fp16 __half;
-extern "C" __half __hip_hadd_clang40_gfx803(__half a, __half b);
-
-__device__ inline __half __hadd(__half a, __half b){
-  return __hip_hadd_clang40_gfx803(a, b);
-}
-
 #endif
 #endif
diff --git a/projects/clr/hipamd/src/hip_fp16.cpp b/projects/clr/hipamd/src/hip_fp16.cpp
index 5d01f73cf7..63d91eb107 100644
--- a/projects/clr/hipamd/src/hip_fp16.cpp
+++ b/projects/clr/hipamd/src/hip_fp16.cpp
@@ -22,7 +22,6 @@ THE SOFTWARE.
 
 #include"hip/hip_fp16.h"
 
-#if __CLANG_VERSION__ == 35
 
 static const unsigned sign_val = 0x8000;
 static const __half __half_value_one_float = {0x3C00};
@@ -375,5 +374,3 @@ __device__ __half2 __low2half2(const __half2 a, const __half2 b){
   return {a.q, b.q};
 }
 
-#endif
-

From a86ef4e57726fd8fed73421020f54f08062c3077 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 13 Dec 2016 20:20:58 -0600
Subject: [PATCH 09/16] added simple half math ops

Change-Id: I10b1d1023a9e5f2ba63f28c4a2bbe60ee49a8aee


[ROCm/clr commit: 01ed8e91e9c1f340c2dd62023dcc8115db756c42]
---
 .../hipamd/include/hip/hcc_detail/hip_fp16.h  | 37 +++++++++++++++++--
 projects/clr/hipamd/src/hip_ir.ll             | 17 ++++++++-
 2 files changed, 50 insertions(+), 4 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
index 735f915bd2..5acf114518 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -30,10 +30,41 @@ THE SOFTWARE.
 #ifdef HIP_HALF_HW_SUPPORT
 
 typedef __fp16 __half;
-extern "C" __half __hip_hadd_clang40_gfx803(__half a, __half b);
+extern "C" __half __hip_hadd_gfx803(__half a, __half b);
+extern "C" __half __hip_hfma_gfx803(__half a, __half b);
+extern "C" __half __hip_hmul_gfx803(__half a, __half b);
+extern "C" __half __hip_hsub_gfx803(__half a, __half b);
 
-__device__ inline __half __hadd(__half a, __half b){
-  return __hip_hadd_clang40_gfx803(a, b);
+__device__ inline __half __hadd(__half a, __half b) {
+  return __hip_hadd_gfx803(a, b);
+}
+
+__device__ inline __half __hadd_sat(__half a, __half b) {
+  return __hip_add_gfx803(a, b);
+}
+
+__device__ inline __half __hfma(__half a, __half b) {
+  return __hip_hfma_gfx803(a, b);
+}
+
+__device__ inline __half __hfma_sat(__half a, __half b) {
+  return __hip_hfma_gfx803(a, b);
+}
+
+__device__ inline __half __hmul(__half a, __half b) {
+  return __hip_hmul_gfx803(a, b);
+}
+
+__device__ inline __half __hmul_sat(__half a, __half b) {
+  return __hip_hmul_gfx803(a, b);
+}
+
+__device__ inline __half __hsub(__half a, __half b) {
+  return __hip_hsub_gfx803(a, b);
+}
+
+__device__ inline __half __hsub_sat(__half a, __half b) {
+  return __hip_hsub_gfx803(a, b);
 }
 
 #else
diff --git a/projects/clr/hipamd/src/hip_ir.ll b/projects/clr/hipamd/src/hip_ir.ll
index 078dc3eed5..623be19084 100644
--- a/projects/clr/hipamd/src/hip_ir.ll
+++ b/projects/clr/hipamd/src/hip_ir.ll
@@ -34,9 +34,24 @@ define linkonce_odr spir_func i32 @__rocm_hadd(i32 %in1, i32 %in2) {
   ret i32 %val
 }
 
-define linkonce_odr spir_func half @__hip_hadd_clang40_gfx803(half %a, half %b) {
+define linkonce_odr spir_func half @__hip_hadd_gfx803(half %a, half %b) #1 {
   %val = tail call half asm "v_add_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
   ret half %val
 }
 
+define linkonce_odr spir_func half @__hip_hfma_gfx803(half %a, half %b, half %c) #1 {
+  %val = tail call half asm "v_fma_f16 $0, $1, $2, $3","=v,v,v,v"(half %a, half %b, half %c)
+  ret half %val
+}
+
+define linkonce_odr spir_func half @__hip_hmul_gfx803(half %a, half %b) #1 {
+  %val = tail call half asm "v_mul_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
+  ret half %val
+}
+
+define linkonce_odr spir_func half @__hip_hsub_gfx803(half %a, half %b) #1 {
+  %val = tail call half asm "v_sub_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
+  ret half %val
+}
+
 attributes #1 = { alwaysinline nounwind }

From 4d78e20eeda1236c8b77615e83e51ecf45784fff Mon Sep 17 00:00:00 2001
From: Sandeep Kumar <Sandeep.Kumar@amd.com>
Date: Wed, 14 Dec 2016 15:49:40 +0530
Subject: [PATCH 10/16] Fixes in Makefile of couple of samples

- modified Makefile for hipblas_saxpy to replaced hcblas.so with hipblas.so as part of HCSWAP-100
- Resolved missing separator issue in peer2peer cookbook Makefile

Change-Id: I678fea267eee1481f02da09379339ed78d3f95f2


[ROCm/clr commit: 8b0ce6e81bd092c9907c29092dd059faa9f093bc]
---
 .../hipamd/samples/2_Cookbook/8_peer2peer/Makefile  | 13 +++++++------
 .../samples/7_Advanced/hipblas_saxpy/Makefile       |  2 +-
 2 files changed, 8 insertions(+), 7 deletions(-)

diff --git a/projects/clr/hipamd/samples/2_Cookbook/8_peer2peer/Makefile b/projects/clr/hipamd/samples/2_Cookbook/8_peer2peer/Makefile
index 5cb7473921..0bf9e6f93e 100644
--- a/projects/clr/hipamd/samples/2_Cookbook/8_peer2peer/Makefile
+++ b/projects/clr/hipamd/samples/2_Cookbook/8_peer2peer/Makefile
@@ -1,6 +1,6 @@
 HIP_PATH?= $(wildcard /opt/rocm/hip)
 ifeq (,$(HIP_PATH))
-    HIP_PATH=../../..
+	HIP_PATH=../../..
 endif
 
 HIPCC=$(HIP_PATH)/bin/hipcc
@@ -22,14 +22,15 @@ CXX=$(HIPCC)
 
 
 $(EXECUTABLE): $(OBJECTS)
-    $(HIPCC) $(OBJECTS) -o $@
+	$(HIPCC) $(OBJECTS) -o $@
+
 
 test: $(EXECUTABLE)
-    $(EXECUTABLE)
+	$(EXECUTABLE)
 
 
 clean:
-rm -f $(EXECUTABLE)
-rm -f $(OBJECTS)
-rm -f $(HIP_PATH)/src/*.o
+	rm -f $(EXECUTABLE)
+	rm -f $(OBJECTS)
+	rm -f $(HIP_PATH)/src/*.o
 
diff --git a/projects/clr/hipamd/samples/7_Advanced/hipblas_saxpy/Makefile b/projects/clr/hipamd/samples/7_Advanced/hipblas_saxpy/Makefile
index ed88be2dd0..8586e75d25 100644
--- a/projects/clr/hipamd/samples/7_Advanced/hipblas_saxpy/Makefile
+++ b/projects/clr/hipamd/samples/7_Advanced/hipblas_saxpy/Makefile
@@ -12,7 +12,7 @@ endif
 ifeq (${HIP_PLATFORM}, hcc)
 	HCBLAS_ROOT?= $(wildcard /opt/rocm/hcblas)
 	HIPCC_FLAGS += -stdlib=libc++ -I$(HCBLAS_ROOT)/include
-	LIBS = -L$(HCBLAS_ROOT)/lib -lhcblas
+	LIBS = -L$(HCBLAS_ROOT)/lib -lhipblas -rpath $(HIP_PATH)/lib
 endif
 
 

From 657739b19f39bcf83c6d90d6be5a020a6532bf48 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Wed, 14 Dec 2016 14:18:48 -0600
Subject: [PATCH 11/16] added half2 support

Change-Id: I0f3b9b7037fed97e80ec99f5369c75a63f001aae


[ROCm/clr commit: c20a86d86612eed98aad1e74d615183aa2f6346c]
---
 .../hipamd/include/hip/hcc_detail/hip_fp16.h  | 21 +++++++++++++++++--
 projects/clr/hipamd/src/hip_fp16.cpp          |  3 ++-
 projects/clr/hipamd/src/hip_ir.ll             |  5 +++++
 3 files changed, 26 insertions(+), 3 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
index 5acf114518..3b03174708 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -27,20 +27,30 @@ THE SOFTWARE.
 
 #define __CLANG_VERSION__ __clang_major__ * 10 + __clang_minor__
 
-#ifdef HIP_HALF_HW_SUPPORT
+#if __CLANG_VERSION__ == 40
 
 typedef __fp16 __half;
+
+typedef struct __attribute__((aligned(4))){
+  int a;
+} __half2;
+
 extern "C" __half __hip_hadd_gfx803(__half a, __half b);
 extern "C" __half __hip_hfma_gfx803(__half a, __half b);
 extern "C" __half __hip_hmul_gfx803(__half a, __half b);
 extern "C" __half __hip_hsub_gfx803(__half a, __half b);
 
+extern "C" int __hip_hadd2_gfx803(int a, int b);
+extern "C" int __hip_hfma2_gfx803(int a, int b);
+extern "C" int __hip_hmul2_gfx803(int a, int b);
+extern "C" int __hip_hsub2_gfx803(int a, int b);
+
 __device__ inline __half __hadd(__half a, __half b) {
   return __hip_hadd_gfx803(a, b);
 }
 
 __device__ inline __half __hadd_sat(__half a, __half b) {
-  return __hip_add_gfx803(a, b);
+  return __hip_hadd_gfx803(a, b);
 }
 
 __device__ inline __half __hfma(__half a, __half b) {
@@ -67,6 +77,13 @@ __device__ inline __half __hsub_sat(__half a, __half b) {
   return __hip_hsub_gfx803(a, b);
 }
 
+
+__device__ inline __half2 __hadd2(__half2 a, __half2 b) {
+  __half2 ret;
+  ret.a = __hip_hadd2_gfx803(a.a, b.a);
+  return ret;
+}
+
 #else
 
 typedef struct{
diff --git a/projects/clr/hipamd/src/hip_fp16.cpp b/projects/clr/hipamd/src/hip_fp16.cpp
index 63d91eb107..3bf6bd395f 100644
--- a/projects/clr/hipamd/src/hip_fp16.cpp
+++ b/projects/clr/hipamd/src/hip_fp16.cpp
@@ -22,6 +22,7 @@ THE SOFTWARE.
 
 #include"hip/hip_fp16.h"
 
+#if __CLANG_VERSION__ == 35
 
 static const unsigned sign_val = 0x8000;
 static const __half __half_value_one_float = {0x3C00};
@@ -373,4 +374,4 @@ __device__ __half2 __lowhigh2highlow(const __half2 a){
 __device__ __half2 __low2half2(const __half2 a, const __half2 b){
   return {a.q, b.q};
 }
-
+#endif
diff --git a/projects/clr/hipamd/src/hip_ir.ll b/projects/clr/hipamd/src/hip_ir.ll
index 623be19084..831c4159f0 100644
--- a/projects/clr/hipamd/src/hip_ir.ll
+++ b/projects/clr/hipamd/src/hip_ir.ll
@@ -54,4 +54,9 @@ define linkonce_odr spir_func half @__hip_hsub_gfx803(half %a, half %b) #1 {
   ret half %val
 }
 
+define linkonce_odr spir_func i32 @__hip_hadd2_gfx803(i32 %a i32 %b) #1 {
+  %val = tail call i32 asm "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_0 dst_unused:UNUSED_PRESERVE src0_sel:WORD_0 src1_sel:WORD_0","=v,v,v"(i32 %a, i32 %b)
+  ret i32 %val
+}
+
 attributes #1 = { alwaysinline nounwind }

From 1844a95003c40bfe6985a511fa57241b66c0173c Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Wed, 14 Dec 2016 16:50:16 -0600
Subject: [PATCH 12/16] fixed compilation issues

Change-Id: I96692538736e2e4f2da9dba9c8c29a164aec4c0d


[ROCm/clr commit: a1d1fcfdacd53852628bc9173ebb2fa9056a6ea9]
---
 projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h | 2 +-
 projects/clr/hipamd/src/hip_ir.ll                     | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
index 3b03174708..fb6cd0a44c 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -27,7 +27,7 @@ THE SOFTWARE.
 
 #define __CLANG_VERSION__ __clang_major__ * 10 + __clang_minor__
 
-#if __CLANG_VERSION__ == 40
+#ifdef HIP_HALF_HW_SUPPORT
 
 typedef __fp16 __half;
 
diff --git a/projects/clr/hipamd/src/hip_ir.ll b/projects/clr/hipamd/src/hip_ir.ll
index 831c4159f0..4b7bd3b10e 100644
--- a/projects/clr/hipamd/src/hip_ir.ll
+++ b/projects/clr/hipamd/src/hip_ir.ll
@@ -54,7 +54,7 @@ define linkonce_odr spir_func half @__hip_hsub_gfx803(half %a, half %b) #1 {
   ret half %val
 }
 
-define linkonce_odr spir_func i32 @__hip_hadd2_gfx803(i32 %a i32 %b) #1 {
+define linkonce_odr spir_func i32 @__hip_hadd2_gfx803(i32 %a, i32 %b) #1 {
   %val = tail call i32 asm "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_0 dst_unused:UNUSED_PRESERVE src0_sel:WORD_0 src1_sel:WORD_0","=v,v,v"(i32 %a, i32 %b)
   ret i32 %val
 }

From 432f1c2b7a01a7e9b0fe7bde1761859e680636dd Mon Sep 17 00:00:00 2001
From: Peter Steinberger <steipete@gmail.com>
Date: Tue, 13 Dec 2016 23:48:04 +0100
Subject: [PATCH 13/16] Fixes a typo: perforamnce -> performance

Change-Id: I85e3b3d22c98c16556227283bfb33530e1bce2cf


[ROCm/clr commit: e2dc1cc27f6b5a14fceb2b94b7f04534bc1cc1d5]
---
 projects/clr/hipamd/README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/projects/clr/hipamd/README.md b/projects/clr/hipamd/README.md
index 83d041786e..802bed75b3 100644
--- a/projects/clr/hipamd/README.md
+++ b/projects/clr/hipamd/README.md
@@ -89,7 +89,7 @@ The HIP Runtime API code and compute kernel definition can exist in the same sou
 ## HIP Portability and Compiler Technology
 HIP C++ code can be compiled with either :
 - On the Nvidia CUDA platform, HIP provides header file which translate from the HIP runtime APIs to CUDA runtime APIs.  The header file contains mostly inlined 
-  functions and thus has very low overhead - developers coding in HIP should expect the same perforamnce as coding in native CUDA.  The code is then 
+  functions and thus has very low overhead - developers coding in HIP should expect the same performance as coding in native CUDA.  The code is then
   compiled with nvcc, the standard C++ compiler provided with the CUDA SDK.  Developers can use any tools supported by the CUDA SDK including the CUDA
   profiler and debugger.
 - On the AMD ROCm platform, HIP provides a header and runtime library built on top of hcc compiler.  The HIP runtime implements HIP streams, events, and memory APIs, 

From 716fac9a996d517bd79879ef9c4c31c35965a199 Mon Sep 17 00:00:00 2001
From: Martin Schleiss <schleissmartin@gmail.com>
Date: Wed, 14 Dec 2016 00:13:30 +0100
Subject: [PATCH 14/16] Fix various typos

Conflicts:
	README.md

Change-Id: Ie296d503d16121a62fed1a208352ec2b81c97fd9


[ROCm/clr commit: dc8db38f044d06b6e1bbcaf9867e6ffcbd3a34b9]
---
 projects/clr/hipamd/CONTRIBUTING.md                         | 6 +++---
 projects/clr/hipamd/README.md                               | 4 ++--
 projects/clr/hipamd/RELEASE.md                              | 4 ++--
 projects/clr/hipamd/docs/markdown/hip_kernel_language.md    | 2 +-
 projects/clr/hipamd/docs/markdown/hip_porting_driver_api.md | 6 +++---
 projects/clr/hipamd/docs/markdown/hip_porting_guide.md      | 2 +-
 projects/clr/hipamd/docs/markdown/hip_terms2.md             | 2 +-
 7 files changed, 13 insertions(+), 13 deletions(-)

diff --git a/projects/clr/hipamd/CONTRIBUTING.md b/projects/clr/hipamd/CONTRIBUTING.md
index f6ed47acef..81c4bc8c32 100644
--- a/projects/clr/hipamd/CONTRIBUTING.md
+++ b/projects/clr/hipamd/CONTRIBUTING.md
@@ -93,7 +93,7 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl
 ## Coding Guidelines (in brief)
 - Code Indentation:
     - Tabs should be expanded to spaces.
-    - Use 4 spaces indendation.
+    - Use 4 spaces indentation.
 - Capitalization and Naming
     - Prefer camelCase for HIP interfaces and internal symbols.  Note HCC uses _ for separator.  
       This guideline is not yet consistently followed in HIP code - eventual compliance is aspirational.
@@ -120,7 +120,7 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl
 
 - HIP_INIT_API() should be placed at the start of each top-level HIP API.  This function will make sure the HIP runtime is initialized,
   and also constructs an appropriate API string for tracing and CodeXL marker tracing.  The arguments to HIP_INIT_API should match
-  those of the parent fucntion.  
+  those of the parent function.  
 - ihipLogStatus should only be called from top-level HIP APIs,and should be called to log and return the error code.  The error code 
   is used by the GetLastError and PeekLastError functions - if a HIP API simply returns, then the error will not be logged correctly.
 
@@ -161,4 +161,4 @@ doxygen bug list.
 
 ##  Other Tips:
 ### Markdown Editing
-Recommended to use an offline Markown viewer to review documentation, such as Markdown Preview Plus extension in Chrome browser, or Remarkable.
+Recommended to use an offline Markdown viewer to review documentation, such as Markdown Preview Plus extension in Chrome browser, or Remarkable.
diff --git a/projects/clr/hipamd/README.md b/projects/clr/hipamd/README.md
index 802bed75b3..f61c3b106a 100644
--- a/projects/clr/hipamd/README.md
+++ b/projects/clr/hipamd/README.md
@@ -88,7 +88,7 @@ The HIP Runtime API code and compute kernel definition can exist in the same sou
 
 ## HIP Portability and Compiler Technology
 HIP C++ code can be compiled with either :
-- On the Nvidia CUDA platform, HIP provides header file which translate from the HIP runtime APIs to CUDA runtime APIs.  The header file contains mostly inlined 
+- On the NVIDIA CUDA platform, HIP provides header file which translate from the HIP runtime APIs to CUDA runtime APIs.  The header file contains mostly inlined 
   functions and thus has very low overhead - developers coding in HIP should expect the same performance as coding in native CUDA.  The code is then
   compiled with nvcc, the standard C++ compiler provided with the CUDA SDK.  Developers can use any tools supported by the CUDA SDK including the CUDA
   profiler and debugger.
@@ -120,7 +120,7 @@ make
 
  
 ## More Examples
-The GitHub repot [HIP-Examples](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP-Examples.git) contains a hipified vesion of the popular Rodinia benchmark suite.
+The GitHub repository [HIP-Examples](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP-Examples.git) contains a hipified version of the popular Rodinia benchmark suite.
 The README with the procedures and tips the team used during this porting effort is here: [Rodinia Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP-Examples/blob/master/rodinia_3.0/hip/README.hip_porting)
 
 ## Tour of the HIP Directories
diff --git a/projects/clr/hipamd/RELEASE.md b/projects/clr/hipamd/RELEASE.md
index a7c770f611..34eab60833 100644
--- a/projects/clr/hipamd/RELEASE.md
+++ b/projects/clr/hipamd/RELEASE.md
@@ -3,7 +3,7 @@
 We have attempted to document known bugs and limitations - in particular the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) document uses the phrase "Under Development", and the [HIP Runtime API bug list](http://gpuopen-professionalcompute-tools.github.io/HIP/bug.html) lists known bugs. 
 
 Upcoming:
-- Stability: Enforce perioidic host synchronization to reclaim resources if the application has launched a large
+- Stability: Enforce periodic host synchronization to reclaim resources if the application has launched a large
   number of commands (>1K) without synchronizing.  
 - Register keyword now silently ignored on HCC (previously would emit warning).
 - Doc updates: Add some more frequently asked questions to FAQ, fix TOC in some files, review.
@@ -73,7 +73,7 @@ Date: 2016.04.25
     - Create static library and link.
     - Set HIP_PATH to install.
 - Make hipDevice and hipStream thread-safe.
-    - Prefered hipStream usage is still to create new streams for each new thread, but it works even if you don;t.
+    - Preferred hipStream usage is still to create new streams for each new thread, but it works even if you don;t.
 - Improve automated platform detection: If AMD GPU is installed and detected by driver, default HIP_PLATFORM to hcc.
 - HIP_TRACE_API now prints arguments to the HIP function (in addition to name of function).
 - Deprecate hipDeviceGetProp (Replace with hipGetDeviceProp)
diff --git a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md
index f84868987c..3fd4ea9aca 100644
--- a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md
+++ b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md
@@ -105,7 +105,7 @@ HIP parses the `__noinline__` and `__forceinline__` keywords and converts them t
 
 
 ```
-// Example psuedocode introducing hipLaunchKernel:
+// Example pseudo code introducing hipLaunchKernel:
 __global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N)
 {
 ...
diff --git a/projects/clr/hipamd/docs/markdown/hip_porting_driver_api.md b/projects/clr/hipamd/docs/markdown/hip_porting_driver_api.md
index b0ac3ecf1d..5093291baa 100644
--- a/projects/clr/hipamd/docs/markdown/hip_porting_driver_api.md
+++ b/projects/clr/hipamd/docs/markdown/hip_porting_driver_api.md
@@ -31,8 +31,8 @@ accelerator language front-end.  Here, NVCC is not used.   Instead, the environm
 different kernel language or different compilation flow.
 Other environments have many kernels and do not want them to be all loaded automatically.
 The Module functions can be used to load the generated code objects and launch kernels.
-As we will see below, HIP defines a Module API which provides similar explict control over code
-object managemenet.
+As we will see below, HIP defines a Module API which provides similar explicit control over code
+object management.
 
 ### cuCtx API
 The Driver API defines "Context" and "Devices" as separate entities.  
@@ -144,7 +144,7 @@ table shows the type equivalence to enable this interaction.
 The hipModule interface does not support the hipModuleLoadEx function, which is used to control PTX compilaton options.
 HCC does not use PTX and does not support the same compilation options.
 In fact, HCC code objects always contain fully compiled ISA and do not require additional compilation as part of the load step.
-Code which requires this functionaly should use platform-specific coding, calling `cuModuleLoadEx` 
+Code which requires this functionally should use platform-specific coding, calling `cuModuleLoadEx` 
 on the NVCC path and hipModuleLoad on the hcc path.  For example:
 
 ```
diff --git a/projects/clr/hipamd/docs/markdown/hip_porting_guide.md b/projects/clr/hipamd/docs/markdown/hip_porting_guide.md
index c530df5098..e0e74c0f89 100644
--- a/projects/clr/hipamd/docs/markdown/hip_porting_guide.md
+++ b/projects/clr/hipamd/docs/markdown/hip_porting_guide.md
@@ -1,5 +1,5 @@
 # HIP Porting Guide 
-In addition to providing a portable C++ programmming environement for GPUs, HIP is designed to ease
+In addition to providing a portable C++ programming environment for GPUs, HIP is designed to ease
 the porting of existing CUDA code into the HIP environment.  This section describes the available tools 
 and provides practical suggestions on how to port CUDA code and work through common issues. 
 
diff --git a/projects/clr/hipamd/docs/markdown/hip_terms2.md b/projects/clr/hipamd/docs/markdown/hip_terms2.md
index 9603807925..be859ffb32 100644
--- a/projects/clr/hipamd/docs/markdown/hip_terms2.md
+++ b/projects/clr/hipamd/docs/markdown/hip_terms2.md
@@ -13,7 +13,7 @@ The default device can be set with hipSetDevice.
 - hcc = Heterogeneous Compute Compiler (https://bitbucket.org/multicoreware/hcc/wiki/Home).  
 
 - hipify - tool to convert CUDA(R) code to portable C++ code.
-- hipconfig - tool to report various confoguration properties of the target platform.
+- hipconfig - tool to report various configuration properties of the target platform.
 
 - nvcc = nvcc compiler, do not capitalize.
 - hcc  = heterogeneous compute compiler, do not capitalize.

From 23ddb5aec442e01fd6c4d0188f8cdc5286e11134 Mon Sep 17 00:00:00 2001
From: Martin Schleiss <schleissmartin@gmail.com>
Date: Wed, 14 Dec 2016 00:18:13 +0100
Subject: [PATCH 15/16] Fix another typo

[ROCm/clr commit: 8a4fe76de2b65d27573aaaadd076ed158b70427c]
---
 projects/clr/hipamd/docs/markdown/hip_faq.md | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/projects/clr/hipamd/docs/markdown/hip_faq.md b/projects/clr/hipamd/docs/markdown/hip_faq.md
index 70ad94ba43..7a41e38bd4 100644
--- a/projects/clr/hipamd/docs/markdown/hip_faq.md
+++ b/projects/clr/hipamd/docs/markdown/hip_faq.md
@@ -107,7 +107,7 @@ However, we can provide a rough summary of the features included in each CUDA SD
 
 ### What libraries does HIP support?
 HIP includes growing support for the 4 key math libraries using hcBlas, hcFft, hcrng, and hcsparse).
-These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaces with other HCC code.  Developers should use conditional compliation if portability to nvcc systems is desired - using calls to cu* routines on one path and hc* routines on the other.  
+These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaces with other HCC code.  Developers should use conditional compilation if portability to nvcc systems is desired - using calls to cu* routines on one path and hc* routines on the other.  
 
 - [hcblas](https://bitbucket.org/multicoreware/hcblas)
 - [hcfft](https://bitbucket.org/multicoreware/hcfft)
@@ -132,7 +132,7 @@ HIP offers several benefits over OpenCL:
 Both HIP and CUDA are dialects of C++, and thus porting between them is relatively straightforward.
 Both dialects support templates, classes, lambdas, and other C++ constructs.
 As one example, the hipify tool was originally a perl script that used simple text conversions from CUDA to HIP.
-HIP and CUDA provide similar math library calls as well.  In summary, the HIP philospohy was to make the HIP language close enough to CUDA that the porting effort is relatively simple.
+HIP and CUDA provide similar math library calls as well.  In summary, the HIP philosophy was to make the HIP language close enough to CUDA that the porting effort is relatively simple.
 This reduces the potential for error, and also makes it easy to automate the translation.  HIP's goal is to quickly get the ported program running on both platforms with little manual intervention,
 so that the programmer can focus on performance optimizations.
 

From bceb8949c7cfec279b9d12083537116d2da2b335 Mon Sep 17 00:00:00 2001
From: Brecht Carlier <erazerbrecht@gmail.com>
Date: Wed, 14 Dec 2016 21:41:21 +0100
Subject: [PATCH 16/16] Update hip_faq.md

Fixed navigation and list.


[ROCm/clr commit: cc73dd4ed53ac785d2019fe893989292815f035d]
---
 projects/clr/hipamd/docs/markdown/hip_faq.md | 5 ++---
 1 file changed, 2 insertions(+), 3 deletions(-)

diff --git a/projects/clr/hipamd/docs/markdown/hip_faq.md b/projects/clr/hipamd/docs/markdown/hip_faq.md
index 7a41e38bd4..0b79976988 100644
--- a/projects/clr/hipamd/docs/markdown/hip_faq.md
+++ b/projects/clr/hipamd/docs/markdown/hip_faq.md
@@ -10,7 +10,7 @@
 - [What specific version of CUDA does HIP support?](#what-specific-version-of-cuda-does-hip-support)
 - [What libraries does HIP support?](#what-libraries-does-hip-support)
 - [How does HIP compare with OpenCL?](#how-does-hip-compare-with-opencl)
-- [How does porting CUDA to HIP compare to porting CUDA to OpenCL?]
+- [How does porting CUDA to HIP compare to porting CUDA to OpenCL?](#how-does-porting-cuda-to-hip-compare-to-porting-cuda-to-opencl)
 - [What hardware does HIP support?](#what-hardware-does-hip-support)
 - [Does Hipify automatically convert all source code?](#does-hipify-automatically-convert-all-source-code)
 - [What is NVCC?](#what-is-nvcc)
@@ -53,7 +53,6 @@ At a high-level, the following features are not supported:
 - Graphics interoperation with OpenGL or Direct3D
 - CUDA Driver API (Under Development)
 - CUDA IPC Functions (Under Development)
-
 - CUDA array, mipmappedArray and pitched memory
 - MemcpyToSymbol functions
 - Queue priority controls
@@ -196,8 +195,8 @@ HIP will set the platform to HCC if it sees that the AMD graphics driver is inst
 Sometimes this isn't what you want - you can force HIP to recognize the platform by setting HIP_PLATFORM to hcc (or nvcc)
 ```
 export HIP_PLATFORM=hcc
-
 ```
+
 One symptom of this problem is the message "error: 'unknown error'(11) at square.hipref.cpp:56".  This can occur if you have a CUDA installation on an AMD platform, and HIP incorrectly detects the platform as nvcc.  HIP may be able to compile the application using the nvcc tool-chain, but will generate this error at runtime since the platform does not have a CUDA device. The fix is to set HIP_PLATFORM=hcc and rebuild the issue. 
 
 If you see issues related to incorrect platform detection, please file an issue with the GitHub issue tracker so we can improve HIP's platform detection logic.