Merge 'master' into 'amd-master'

Change-Id: I03be43978ccec5e0248abfd1e6e33a7bfa309ba8


[ROCm/clr commit: 2e5dcfa432]

projects/clr/hipamd/hipify-clang/README.md

@@ -321,7 +321,6 @@ For example:
 ./hipify-clang \
   square.cu \
   --cuda-path=/usr/local/cuda-8.0 \
-  -- \
   -I /usr/local/cuda-8.0/samples/common/inc
 ```
 

projects/clr/hipamd/hipify-clang/src/ArgParse.cpp

@@ -85,8 +85,22 @@ cl::opt<bool> Examine("examine",
   cl::cat(ToolTemplateCategory));
 
 cl::opt<bool> DashDash("-",
-  cl::desc("Separator between hipify-clang and clang options;\ndon't specify if there are no clang options."),
+  cl::desc("Separator between hipify-clang and clang options;\ndon't specify if there are no clang options"),
   cl::value_desc("--"),
   cl::cat(ToolTemplateCategory));
 
+cl::list<std::string> IncludeDirs("I",
+  cl::desc("Add directory to include search path"),
+  cl::value_desc("directory"),
+  cl::ZeroOrMore,
+  cl::Prefix,
+  cl::cat(ToolTemplateCategory));
+
+cl::list<std::string> MacroNames("D",
+  cl::desc("Define <macro> to <value> or 1 if <value> omitted"),
+  cl::value_desc("macro>=<value"),
+  cl::ZeroOrMore,
+  cl::Prefix,
+  cl::cat(ToolTemplateCategory));
+
 cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage);

projects/clr/hipamd/hipify-clang/src/ArgParse.h

@@ -33,6 +33,8 @@ extern cl::opt<std::string> OutputFilename;
 extern cl::opt<std::string> OutputDir;
 extern cl::opt<std::string> TemporaryDir;
 extern cl::opt<std::string> CudaPath;
+extern cl::list<std::string> IncludeDirs;
+extern cl::list<std::string> MacroNames;
 extern cl::opt<bool> Inplace;
 extern cl::opt<bool> SaveTemps;
 extern cl::opt<bool> NoBackup;

projects/clr/hipamd/hipify-clang/src/main.cpp

@@ -39,30 +39,56 @@ THE SOFTWARE.
 std::string sHipify = "[HIPIFY] ", sConflict = "conflict: ", sError = "error: ";
 namespace ct = clang::tooling;
 
-std::string getAbsoluteDirectory(const std::string& sDir, std::error_code& EC,
+std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) {
+  if (sFile.empty()) {
+    return sFile;
+  }
+  if (!sys::fs::exists(sFile)) {
+    llvm::errs() << "\n" << sHipify << sError << "source file: " << sFile << " doesn't exist\n";
+    EC = std::error_code(static_cast<int>(std::errc::no_such_file_or_directory), std::generic_category());
+    return "";
+  }
+  SmallString<256> fileAbsPath;
+  EC = llcompat::real_path(sFile, fileAbsPath, true);
+  if (EC) {
+    llvm::errs() << "\n" << sHipify << sError << EC.message() << ": source file: " << sFile << "\n";
+    return "";
+  }
+  EC = std::error_code();
+  return fileAbsPath.c_str();
+}
+
+std::string getAbsoluteDirectoryPath(const std::string& sDir, std::error_code& EC,
                                  const std::string& sDirType = "temporary",
                                  bool bCreateDir = true) {
   if (sDir.empty()) {
     return sDir;
   }
+  EC = std::error_code();
   SmallString<256> dirAbsPath;
-  EC = llcompat::real_path(sDir, dirAbsPath, true);
+  if (sys::fs::exists(sDir)) {
-  if (!EC && sys::fs::is_regular_file(dirAbsPath)) {
+    if (sys::fs::is_regular_file(sDir)) {
-    llvm::errs() << "\n" << sHipify << sError << sDir << " is not a directory\n";
+      llvm::errs() << "\n" << sHipify << sError << sDir << " is not a directory\n";
-    EC = std::error_code(static_cast<int>(std::errc::not_a_directory), std::generic_category());
+      EC = std::error_code(static_cast<int>(std::errc::not_a_directory), std::generic_category());
-    return "";
+      return "";
+    }
+  } else {
+    if (bCreateDir) {
+      EC = sys::fs::create_directory(sDir);
+      if (EC) {
+        llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
+        return "";
+      }
+    } else {
+      llvm::errs() << "\n" << sHipify << sError << sDirType << " directory: " << sDir << " doesn't exist\n";
+      EC = std::error_code(static_cast<int>(std::errc::no_such_file_or_directory), std::generic_category());
+      return "";
+    }
   }
-  if (EC && bCreateDir) {
+  EC = llcompat::real_path(sDir, dirAbsPath, true);
-    EC = sys::fs::create_directory(sDir);
+  if (EC) {
-    if (EC) {
+    llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
-      llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
+    return "";
-      return "";
-    }
-    EC = llcompat::real_path(sDir, dirAbsPath, true);
-    if (EC) {
-      llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
-      return "";
-    }
   }
   return dirAbsPath.c_str();
 }
@@ -80,7 +106,7 @@ int main(int argc, const char **argv) {
   std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
   std::string dst = OutputFilename, dstDir = OutputDir;
   std::error_code EC;
-  std::string sOutputDirAbsPath = getAbsoluteDirectory(OutputDir, EC, "output");
+  std::string sOutputDirAbsPath = getAbsoluteDirectoryPath(OutputDir, EC, "output");
   if (EC) {
     return 1;
   }
@@ -114,10 +140,9 @@ int main(int argc, const char **argv) {
   }
   int Result = 0;
   SmallString<128> tmpFile;
-  SmallString<256> sourceAbsPath;
   StringRef sourceFileName, ext = "hip";
-  std::string sTmpFileName;
+  std::string sTmpFileName, sSourceAbsPath;
-  std::string sTmpDirAbsParh = getAbsoluteDirectory(TemporaryDir, EC);
+  std::string sTmpDirAbsParh = getAbsoluteDirectoryPath(TemporaryDir, EC);
   if (EC) {
     return 1;
   }
@@ -131,26 +156,24 @@ int main(int argc, const char **argv) {
     statPrint = &llvm::errs();
   }
   for (const auto & src : fileSources) {
+    // Create a copy of the file to work on. When we're done, we'll move this onto the
+    // output (which may mean overwriting the input, if we're in-place).
+    // Should we fail for some reason, we'll just leak this file and not corrupt the input.
+    sSourceAbsPath = getAbsoluteFilePath(src, EC);
+    if (EC) {
+      continue;
+    }
+    sourceFileName = sys::path::filename(sSourceAbsPath);
     if (dst.empty()) {
       if (Inplace) {
         dst = src;
       } else {
         dst = src + "." + ext.str();
         if (!dstDir.empty()) {
-          dst = sOutputDirAbsPath + "/" + dst;
+          dst = sOutputDirAbsPath + "/" + sourceFileName.str() + "." + ext.str();
         }
       }
     }
-    // Create a copy of the file to work on. When we're done, we'll move this onto the
-    // output (which may mean overwriting the input, if we're in-place).
-    // Should we fail for some reason, we'll just leak this file and not corrupt the input.
-    EC = llcompat::real_path(src, sourceAbsPath, true);
-    sourceFileName = sys::path::filename(sourceAbsPath);
-    if (EC) {
-      llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << src << "\n";
-      Result = 1;
-      continue;
-    }
     if (TemporaryDir.empty()) {
       EC = sys::fs::createTemporaryFile(sourceFileName, ext, tmpFile);
       if (EC) {
@@ -176,7 +199,7 @@ int main(int argc, const char **argv) {
     ct::RefactoringTool Tool(OptionsParser.getCompilations(), std::string(tmpFile.c_str()));
     ct::Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile.c_str());
     ReplacementsFrontendActionFactory<HipifyAction> actionFactory(&replacementsToUse);
-    std::string sInclude = "-I" + sys::path::parent_path(sourceAbsPath).str();
+    std::string sInclude = "-I" + sys::path::parent_path(sSourceAbsPath).str();
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN));
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN));
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN));
@@ -199,6 +222,18 @@ int main(int argc, const char **argv) {
 #if defined(HIPIFY_CLANG_RES)
     Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES));
 #endif
+    if (!MacroNames.empty()) {
+      for (std::string s : MacroNames) {
+        Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::END));
+        Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::END));
+      }
+    }
+    if (!IncludeDirs.empty()) {
+      for (std::string s : IncludeDirs) {
+        Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-I", ct::ArgumentInsertPosition::END));
+        Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::END));
+      }
+    }
     Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster());
     Statistics& currentStat = Statistics::current();
     // Hipify _all_ the things!

projects/clr/hipamd/include/hip/hcc_detail/math_functions.h

@@ -507,13 +507,13 @@ float ynf(int n, float x)
 // BEGIN INTRINSICS
 __DEVICE__
 inline
-float __cosf(float x) { return __ocml_cos_f32(x); }
+float __cosf(float x) { return __ocml_native_cos_f32(x); }
 __DEVICE__
 inline
-float __exp10f(float x) { return __ocml_exp10_f32(x); }
+float __exp10f(float x) { return __ocml_native_exp10_f32(x); }
 __DEVICE__
 inline
-float __expf(float x) { return __ocml_exp_f32(x); }
+float __expf(float x) { return __ocml_native_exp_f32(x); }
 #if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
@@ -613,7 +613,7 @@ float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); }
 #endif
 __DEVICE__
 inline
-float __fsqrt_rn(float x) { return __ocml_sqrt_f32(x); }
+float __fsqrt_rn(float x) { return __ocml_native_sqrt_f32(x); }
 #if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
@@ -638,13 +638,13 @@ float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); }
 #endif
 __DEVICE__
 inline
-float __log10f(float x) { return __ocml_log10_f32(x); }
+float __log10f(float x) { return __ocml_native_log10_f32(x); }
 __DEVICE__
 inline
-float __log2f(float x) { return __ocml_log2_f32(x); }
+float __log2f(float x) { return __ocml_native_log2_f32(x); }
 __DEVICE__
 inline
-float __logf(float x) { return __ocml_log_f32(x); }
+float __logf(float x) { return __ocml_native_log_f32(x); }
 __DEVICE__
 inline
 float __powf(float x, float y) { return __ocml_pow_f32(x, y); }
@@ -655,15 +655,12 @@ __DEVICE__
 inline
 void __sincosf(float x, float* sptr, float* cptr)
 {
-    float tmp;
+    *sptr = __ocml_native_sin_f32(x);
-
+    *cptr = __ocml_native_cos_f32(x);
-    *sptr =
-        __ocml_sincos_f32(x, (__attribute__((address_space(5))) float*) &tmp);
-    *cptr = tmp;
 }
 __DEVICE__
 inline
-float __sinf(float x) { return __ocml_sin_f32(x); }
+float __sinf(float x) { return __ocml_native_sin_f32(x); }
 __DEVICE__
 inline
 float __tanf(float x) { return __ocml_tan_f32(x); }

projects/clr/hipamd/include/hip/hcc_detail/math_fwd.h

@@ -63,6 +63,8 @@ float __ocml_copysign_f32(float, float);
 __device__
 float __ocml_cos_f32(float);
 __device__
+float __ocml_native_cos_f32(float);
+__device__
 __attribute__((pure))
 __device__
 float __ocml_cosh_f32(float);
@@ -92,12 +94,18 @@ __attribute__((pure))
 float __ocml_exp10_f32(float);
 __device__
 __attribute__((pure))
+float __ocml_native_exp10_f32(float);
+__device__
+__attribute__((pure))
 float __ocml_exp2_f32(float);
 __device__
 __attribute__((pure))
 float __ocml_exp_f32(float);
 __device__
 __attribute__((pure))
+float __ocml_native_exp_f32(float);
+__device__
+__attribute__((pure))
 float __ocml_expm1_f32(float);
 __device__
 __attribute__((const))
@@ -152,17 +160,26 @@ __attribute__((pure))
 float __ocml_log10_f32(float);
 __device__
 __attribute__((pure))
+float __ocml_native_log10_f32(float);
+__device__
+__attribute__((pure))
 float __ocml_log1p_f32(float);
 __device__
 __attribute__((pure))
 float __ocml_log2_f32(float);
 __device__
+__attribute__((pure))
+float __ocml_native_log2_f32(float);
+__device__
 __attribute__((const))
 float __ocml_logb_f32(float);
 __device__
 __attribute__((pure))
 float __ocml_log_f32(float);
 __device__
+__attribute__((pure))
+float __ocml_native_log_f32(float);
+__device__
 float __ocml_modf_f32(float, __attribute__((address_space(5))) float*);
 __device__
 __attribute__((const))
@@ -227,6 +244,8 @@ float __ocml_sincospi_f32(float, __attribute__((address_space(5))) float*);
 __device__
 float __ocml_sin_f32(float);
 __device__
+float __ocml_native_sin_f32(float);
+__device__
 __attribute__((pure))
 float __ocml_sinh_f32(float);
 __device__
@@ -235,6 +254,9 @@ __device__
 __attribute__((const))
 float __ocml_sqrt_f32(float);
 __device__
+__attribute__((const))
+float __ocml_native_sqrt_f32(float);
+__device__
 float __ocml_tan_f32(float);
 __device__
 __attribute__((pure))