diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt
index 1a6bc3636d..3348056faf 100644
--- a/projects/clr/hipamd/CMakeLists.txt
+++ b/projects/clr/hipamd/CMakeLists.txt
@@ -14,9 +14,8 @@ if(${HIPIFY_STANDLONE})
   include_directories(${LLVM_INCLUDE_DIRS})
   link_directories(${LLVM_LIBRARY_DIRS})
   add_definitions(${LLVM_DEFINITIONS})
-  add_llvm_executable( hipify src/Cuda2Hip.cpp )
+  add_llvm_executable(hipify src/Cuda2Hip.cpp )
   find_program(LIT_COMMAND lit)
-  
 else()
   set(LLVM_LINK_COMPONENTS
     Option
diff --git a/projects/clr/hipamd/src/Cuda2Hip.cpp b/projects/clr/hipamd/src/Cuda2Hip.cpp
index c2217c5bf0..1c4b812506 100644
--- a/projects/clr/hipamd/src/Cuda2Hip.cpp
+++ b/projects/clr/hipamd/src/Cuda2Hip.cpp
@@ -158,7 +158,7 @@ namespace {
       cuda2hipRename["cudaStreamDefault"] = "hipStreamDefault";
       cuda2hipRename["cudaStreamNonBlocking"] = "hipStreamNonBlocking";
 
-      // Other synchronization 
+      // Other synchronization
       cuda2hipRename["cudaDeviceSynchronize"] = "hipDeviceSynchronize";
       cuda2hipRename["cudaThreadSynchronize"] = "hipDeviceSynchronize";  // translate deprecated cudaThreadSynchronize
       cuda2hipRename["cudaDeviceReset"] = "hipDeviceReset";
@@ -166,7 +166,7 @@ namespace {
       cuda2hipRename["cudaSetDevice"] = "hipSetDevice";
       cuda2hipRename["cudaGetDevice"] = "hipGetDevice";
 
-      // Device 
+      // Device
       cuda2hipRename["cudaDeviceProp"] = "hipDeviceProp_t";
       cuda2hipRename["cudaGetDeviceProperties"] = "hipDeviceGetProperties";
 
@@ -357,11 +357,11 @@ namespace {
       }
     }
 
-    void EndOfMainFile() override 
-    { 
-   
+    void EndOfMainFile() override
+    {
+
     }
-    
+
     bool SeenEnd;
     void setSourceManager(SourceManager * sm) { _sm = sm; }
     void setPreprocessor (Preprocessor  * pp) { _pp = pp; }
@@ -376,102 +376,103 @@ namespace {
   };
 
 class Cuda2HipCallback : public MatchFinder::MatchCallback {
-  public:
+  public:
     Cuda2HipCallback(Replacements *Replace) : Replace(Replace) {}
-
+
     void run(const MatchFinder::MatchResult &Result) override {
-
-      SourceManager * SM = Result.SourceManager;
-
-	  if (const CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaCall"))
-	  {
-		  const FunctionDecl * funcDcl = call->getDirectCallee();
-		  std::string name = funcDcl->getDeclName().getAsString();
-		  if (N.cuda2hipRename.count(name)) {
-			  std::string repName = N.cuda2hipRename[name];
-        SourceLocation sl = call->getLocStart();
-        Replacement Rep(*SM, sl, name.length(), repName);
-			  Replace->insert(Rep);
-		  }
-	  }
-
-      if (const CUDAKernelCallExpr * launchKernel = Result.Nodes.getNodeAs<clang::CUDAKernelCallExpr>("cudaLaunchKernel"))
-      {
-        LangOptions DefaultLangOptions;
-        StringRef initialParamList;
-        SmallString<40> XStr;
-        raw_svector_ostream OS(XStr);
-        OS << "hipLaunchParm lp";
-        const FunctionDecl * kernelDecl = launchKernel->getDirectCallee();
-        SourceLocation l1 = kernelDecl->getNameInfo().getLocStart();
-        l1.dump(*SM);llvm::outs() << "\n";
-        SourceLocation kernelArgListStart = clang::Lexer::findLocationAfterToken(l1, clang::tok::l_paren, *SM, DefaultLangOptions, true);
-        kernelArgListStart.dump(*SM);llvm::outs() << "\n";
-        size_t replacementLength = 0;
-        if (kernelDecl->getNumParams() > 0) {
-          const ParmVarDecl * pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams()-1);
-          SourceLocation kernelArgListEnd = SourceLocation(pvdLast->getLocEnd());
-          SourceLocation stop = clang::Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions);
-          replacementLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart);
-          initialParamList = StringRef(SM->getCharacterData(kernelArgListStart), replacementLength);
-          OS << ", " << initialParamList;
-        }
-          
-        llvm::outs() << "initial paramlist: " << initialParamList << "\n";
-        llvm::outs() << "new paramlist: " << OS.str() << "\n";
-        Replacement Rep0(*(Result.SourceManager), kernelArgListStart, replacementLength, OS.str());
-        Replace->insert(Rep0);
-
-      XStr.clear();
-      OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << kernelDecl->getName() << "), ";
-
-      const CallExpr * config = launchKernel->getConfig();
-      llvm::outs() << "\nKernel config arguments:\n";
-      for (unsigned argno = 0; argno < config->getNumArgs(); argno++)
-      {
-        const Expr * arg = config->getArg(argno);
-        if (!isa<CXXDefaultArgExpr>(arg)) {
-          const ParmVarDecl * pvd = config->getDirectCallee()->getParamDecl(argno);
-          
-          SourceLocation sl(arg->getLocStart());
-          SourceLocation el(arg->getLocEnd());
-          SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions);
-          StringRef outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl));
-          llvm::outs() << "args[ " << argno << "]" << outs << " <" << pvd->getType().getAsString() << ">\n";
-          if (pvd->getType().getAsString().compare("dim3") == 0)
-            OS << " dim3(" << outs << "),";
-          else
-            OS << " " << outs << ",";
-        } else
-          OS << " 0,";
-      }
- 
-      for (unsigned argno = 0; argno < launchKernel->getNumArgs(); argno++)
-      {
-        const Expr * arg = launchKernel->getArg(argno);
-        SourceLocation sl(arg->getLocStart());
-        SourceLocation el(arg->getLocEnd());
-        SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions);
-        std::string outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl));
-        llvm::outs() << outs.c_str() << "\n";
-        OS << " " << outs << ",";
-      }
-      XStr.pop_back();
-      OS << ")";
-      size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchKernel->getLocEnd(), 0, *SM, DefaultLangOptions)) -
-        SM->getCharacterData(launchKernel->getLocStart());
-      Replacement Rep(*SM, launchKernel->getLocStart(), length, OS.str());
-      Replace->insert(Rep);
-	  }
-
-	  if (const MemberExpr * threadIdx = Result.Nodes.getNodeAs<clang::MemberExpr>("cudaBuiltin"))
-	  {
-		  if (const OpaqueValueExpr * refBase = dyn_cast<OpaqueValueExpr>(threadIdx->getBase())) {
-        if (const DeclRefExpr * declRef = dyn_cast<DeclRefExpr>(refBase->getSourceExpr())) {
-          std::string name = declRef->getDecl()->getNameAsString();
-          std::string memberName = threadIdx->getMemberDecl()->getNameAsString();
-          size_t pos = memberName.find_first_not_of("__fetch_builtin_");
-          memberName = memberName.substr(pos, memberName.length() - pos);
+
+      SourceManager * SM = Result.SourceManager;
+
+      if (const CallExpr * call = Result.Nodes.getNodeAs<clang::CallExpr>("cudaCall"))
+      {
+        const FunctionDecl * funcDcl = call->getDirectCallee();
+        std::string name = funcDcl->getDeclName().getAsString();
+        if (N.cuda2hipRename.count(name)) {
+	  std::string repName = N.cuda2hipRename[name];
+          SourceLocation sl = call->getLocStart();
+	  Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ?
+          SM->getImmediateSpellingLoc(sl) : sl, name.length(), repName);
+	  Replace->insert(Rep);
+        }
+      }
+
+      if (const CUDAKernelCallExpr * launchKernel = Result.Nodes.getNodeAs<clang::CUDAKernelCallExpr>("cudaLaunchKernel"))
+      {
+        LangOptions DefaultLangOptions;
+        StringRef initialParamList;
+        SmallString<40> XStr;
+        raw_svector_ostream OS(XStr);
+        OS << "hipLaunchParm lp";
+        const FunctionDecl * kernelDecl = launchKernel->getDirectCallee();
+        SourceLocation l1 = kernelDecl->getNameInfo().getLocStart();
+        l1.dump(*SM);llvm::outs() << "\n";
+        SourceLocation kernelArgListStart = clang::Lexer::findLocationAfterToken(l1, clang::tok::l_paren, *SM, DefaultLangOptions, true);
+        kernelArgListStart.dump(*SM);llvm::outs() << "\n";
+        size_t replacementLength = 0;
+        if (kernelDecl->getNumParams() > 0) {
+          const ParmVarDecl * pvdLast = kernelDecl->getParamDecl(kernelDecl->getNumParams()-1);
+          SourceLocation kernelArgListEnd = SourceLocation(pvdLast->getLocEnd());
+          SourceLocation stop = clang::Lexer::getLocForEndOfToken(kernelArgListEnd, 0, *SM, DefaultLangOptions);
+          replacementLength = SM->getCharacterData(stop) - SM->getCharacterData(kernelArgListStart);
+          initialParamList = StringRef(SM->getCharacterData(kernelArgListStart), replacementLength);
+          OS << ", " << initialParamList;
+        }
+
+        llvm::outs() << "initial paramlist: " << initialParamList << "\n";
+        llvm::outs() << "new paramlist: " << OS.str() << "\n";
+        Replacement Rep0(*(Result.SourceManager), kernelArgListStart, replacementLength, OS.str());
+        Replace->insert(Rep0);
+
+      XStr.clear();
+      OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << kernelDecl->getName() << "), ";
+
+      const CallExpr * config = launchKernel->getConfig();
+      llvm::outs() << "\nKernel config arguments:\n";
+      for (unsigned argno = 0; argno < config->getNumArgs(); argno++)
+      {
+        const Expr * arg = config->getArg(argno);
+        if (!isa<CXXDefaultArgExpr>(arg)) {
+          const ParmVarDecl * pvd = config->getDirectCallee()->getParamDecl(argno);
+
+          SourceLocation sl(arg->getLocStart());
+          SourceLocation el(arg->getLocEnd());
+          SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions);
+          StringRef outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl));
+          llvm::outs() << "args[ " << argno << "]" << outs << " <" << pvd->getType().getAsString() << ">\n";
+          if (pvd->getType().getAsString().compare("dim3") == 0)
+            OS << " dim3(" << outs << "),";
+          else
+            OS << " " << outs << ",";
+        } else
+          OS << " 0,";
+      }
+
+      for (unsigned argno = 0; argno < launchKernel->getNumArgs(); argno++)
+      {
+        const Expr * arg = launchKernel->getArg(argno);
+        SourceLocation sl(arg->getLocStart());
+        SourceLocation el(arg->getLocEnd());
+        SourceLocation stop = clang::Lexer::getLocForEndOfToken(el, 0, *SM, DefaultLangOptions);
+        std::string outs(SM->getCharacterData(sl), SM->getCharacterData(stop) - SM->getCharacterData(sl));
+        llvm::outs() << outs.c_str() << "\n";
+        OS << " " << outs << ",";
+      }
+      XStr.pop_back();
+      OS << ")";
+      size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchKernel->getLocEnd(), 0, *SM, DefaultLangOptions)) -
+        SM->getCharacterData(launchKernel->getLocStart());
+      Replacement Rep(*SM, launchKernel->getLocStart(), length, OS.str());
+      Replace->insert(Rep);
+	  }
+
+	  if (const MemberExpr * threadIdx = Result.Nodes.getNodeAs<clang::MemberExpr>("cudaBuiltin"))
+	  {
+		  if (const OpaqueValueExpr * refBase = dyn_cast<OpaqueValueExpr>(threadIdx->getBase())) {
+        if (const DeclRefExpr * declRef = dyn_cast<DeclRefExpr>(refBase->getSourceExpr())) {
+          std::string name = declRef->getDecl()->getNameAsString();
+          std::string memberName = threadIdx->getMemberDecl()->getNameAsString();
+          size_t pos = memberName.find_first_not_of("__fetch_builtin_");
+          memberName = memberName.substr(pos, memberName.length() - pos);
           name += "." + memberName;
      		  std::string repName = N.cuda2hipRename[name];
           SourceLocation sl = threadIdx->getLocStart();
@@ -558,23 +559,24 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback {
 
     if (const StringLiteral * stringLiteral = Result.Nodes.getNodeAs<clang::StringLiteral>("stringLiteral"))
     {
-      StringRef s = stringLiteral->getString();
-      size_t begin = 0;
-      while ((begin = s.find("cuda", begin)) != StringRef::npos) {
-        const size_t end = s.find_first_of(" ", begin + 4);
-        StringRef name = s.slice(begin, end);
-        llvm::outs() << "\nToken: <" << name << "> found in string literal " << s << "\n";
-        StringRef repName = N.cuda2hipRename[name];
-        if (!repName.empty()) {
-          llvm::outs() << "\nWill be replaced with: <" << repName << "\n";
-          SourceLocation sl = stringLiteral->getLocationOfByte(begin, *SM,
-            Result.Context->getLangOpts(), Result.Context->getTargetInfo());
-          Replacement Rep(*SM, sl, name.size(), repName);
-          Replace->insert(Rep);
-        }
-        if (end == StringRef::npos) break;
-        begin = end + 1;
-      }
+      StringRef s = stringLiteral->getString();
+      size_t begin = 0;
+      while ((begin = s.find("cuda", begin)) != StringRef::npos) {
+        const size_t end = s.find_first_of(" ", begin + 4);
+        StringRef name = s.slice(begin, end);
+        llvm::outs() << "\nToken: <" << name << "> found in string literal " << s << "\n";
+        StringRef repName = N.cuda2hipRename[name];
+        if (!repName.empty()) {
+          llvm::outs() << "\nWill be replaced with: <" << repName << "\n";
+          SourceLocation sl = stringLiteral->getLocationOfByte(begin, *SM,
+            Result.Context->getLangOpts(), Result.Context->getTargetInfo());
+          Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ?
+            SM->getImmediateSpellingLoc(sl) : sl, name.size(), repName);
+          Replace->insert(Rep);
+        }
+        if (end == StringRef::npos) break;
+        begin = end + 1;
+      }
     }
 
     if (const UnaryExprOrTypeTraitExpr * expr = Result.Nodes.getNodeAs<clang::UnaryExprOrTypeTraitExpr>("cudaStructSizeOf"))
@@ -609,7 +611,7 @@ int main(int argc, const char **argv) {
   llvm::sys::PrintStackTraceOnErrorSignal();
 
   int Result;
-  
+
   CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::Required);
   std::string dst = OutputFilename;
   std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
@@ -642,9 +644,9 @@ int main(int argc, const char **argv) {
   Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(enumDecl(matchesName("cuda.*")))).bind("cudaEnumConstantDecl"), &Callback);
   Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(cxxRecordDecl(matchesName("cuda.*")))).bind("cudaStructVar"), &Callback);
   Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(pointsTo(cxxRecordDecl(matchesName("cuda.*"))))).bind("cudaStructVarPtr"), &Callback);
-  Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(namedDecl(matchesName("cuda.*")))).bind("cudaParamDecl"), &Callback);
+  Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(namedDecl(matchesName("cuda.*")))).bind("cudaParamDecl"), &Callback);
   Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(pointsTo(namedDecl(matchesName("cuda.*"))))).bind("cudaParamDeclPtr"), &Callback);
-  Finder.addMatcher(expr(sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration(cxxRecordDecl(matchesName("cuda.*"))))))).bind("cudaStructSizeOf"), &Callback);
+  Finder.addMatcher(expr(sizeOfExpr(hasArgumentOfType(recordType(hasDeclaration(cxxRecordDecl(matchesName("cuda.*"))))))).bind("cudaStructSizeOf"), &Callback);
   Finder.addMatcher(stringLiteral().bind("stringLiteral"), &Callback);
 
   auto action = newFrontendActionFactory(&Finder, &PPCallbacks);
@@ -663,7 +665,7 @@ int main(int argc, const char **argv) {
 
 	  Tool.clearArgumentsAdjusters();
   }
-  
+
   LangOptions DefaultLangOptions;
   IntrusiveRefCntPtr<DiagnosticOptions> DiagOpts = new DiagnosticOptions();
   TextDiagnosticPrinter DiagnosticPrinter(llvm::errs(), &*DiagOpts);
diff --git a/projects/clr/hipamd/test/CMakeLists.txt b/projects/clr/hipamd/test/CMakeLists.txt
index 4f74091b0b..5a0250381d 100644
--- a/projects/clr/hipamd/test/CMakeLists.txt
+++ b/projects/clr/hipamd/test/CMakeLists.txt
@@ -6,7 +6,7 @@ if( NOT PYTHONINTERP_FOUND )
           "Please install Python or specify the PYTHON_EXECUTABLE CMake variable.")
 endif()
 
-set(BINARY_DIR       ${CMAKE_CURRENT_BINARY_DIR}  )
+set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}/..  )
 
 configure_file(
   ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.in
@@ -16,7 +16,7 @@ configure_file(
 add_lit_testsuite(check-hipify "Running HIPify regression tests"
   ${CMAKE_CURRENT_SOURCE_DIR}
   PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg
-  DEPENDS hipify
+  DEPENDS hipify FileCheck lit
   )
 
 add_custom_target(check)
diff --git a/projects/clr/hipamd/test/lit.cfg b/projects/clr/hipamd/test/lit.cfg
index 775cdbf0d4..d9606b48fc 100644
--- a/projects/clr/hipamd/test/lit.cfg
+++ b/projects/clr/hipamd/test/lit.cfg
@@ -51,4 +51,5 @@ if not tool_path:
     lit_config.note('Did not find ' + tool_name + ' in ' + llvm_tools_dir)
     tool_path = llvm_tools_dir + '/' + tool_name
 config.substitutions.append((tool_name, tool_path))
+config.substitutions.append(("hipify", obj_root+"/hipify"))
 
diff --git a/projects/clr/hipamd/test/lit.site.cfg.in b/projects/clr/hipamd/test/lit.site.cfg.in
index 6fc35440d4..4511316ac7 100644
--- a/projects/clr/hipamd/test/lit.site.cfg.in
+++ b/projects/clr/hipamd/test/lit.site.cfg.in
@@ -1,7 +1,7 @@
 import sys
 
 config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
-config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
+config.obj_root = "@BINARY_DIR@"
 
 # Support substitution of the tools and libs dirs with user parameters. This is
 # used when we can't determine the tool dir at configuration time.