Merge 'master' into 'amd-master'

Change-Id: I7a1d8c05593540710d27f3013633bbc545cb19db

hipify-clang/src/ArgParse.cpp

@@ -29,6 +29,11 @@ cl::opt<std::string> OutputFilename("o",
   cl::value_desc("filename"),
   cl::cat(ToolTemplateCategory));
 
+cl::opt<std::string> OutputDir("o-dir",
+  cl::desc("Output direcory"),
+  cl::value_desc("directory"),
+  cl::cat(ToolTemplateCategory));
+
 cl::opt<std::string> TemporaryDir("temp-dir",
   cl::desc("Temporary direcory"),
   cl::value_desc("directory"),

hipify-clang/src/ArgParse.h

@@ -30,6 +30,7 @@ namespace ct = clang::tooling;
 
 extern cl::OptionCategory ToolTemplateCategory;
 extern cl::opt<std::string> OutputFilename;
+extern cl::opt<std::string> OutputDir;
 extern cl::opt<std::string> TemporaryDir;
 extern cl::opt<bool> Inplace;
 extern cl::opt<bool> SaveTemps;

hipify-clang/src/main.cpp

@@ -36,13 +36,47 @@ THE SOFTWARE.
 
 #define DEBUG_TYPE "cuda2hip"
 
+std::string sHipify = "[HIPIFY] ", sConflict = "conflict: ", sError = "error: ";
 namespace ct = clang::tooling;
 
+std::string getAbsoluteDirectory(const std::string& sDir, std::error_code& EC,
+                                 const std::string& sDirType = "temporary",
+                                 bool bCreateDir = true) {
+  if (sDir.empty()) {
+    return sDir;
+  }
+  SmallString<256> dirAbsPath;
+  EC = sys::fs::real_path(sDir, dirAbsPath, true);
+  if (!EC && sys::fs::is_regular_file(dirAbsPath)) {
+    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());
+    return "";
+  }
+  if (EC && bCreateDir) {
+    EC = sys::fs::create_directory(sDir);
+    if (EC) {
+      llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
+      return "";
+    }
+    EC = sys::fs::real_path(sDir, dirAbsPath, true);
+    if (EC) {
+      llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
+      return "";
+    }
+  }
+  return dirAbsPath.c_str();
+}
+
 int main(int argc, const char **argv) {
   llcompat::PrintStackTraceOnErrorSignal();
   ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore);
   std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
-  std::string dst = OutputFilename, sHipify = "[HIPIFY] ", sConflict = "conflict: ", sError = "error: ";
+  std::string dst = OutputFilename, dstDir = OutputDir;
+  std::error_code EC;
+  std::string sOutputDirAbsPath = getAbsoluteDirectory(OutputDir, EC, "output");
+  if (EC) {
+    return 1;
+  }
   if (!dst.empty()) {
     if (fileSources.size() > 1) {
       llvm::errs() << sHipify << sConflict << "-o and multiple source files are specified.\n";
@@ -56,39 +90,29 @@ int main(int argc, const char **argv) {
       llvm::errs() << sHipify << sConflict << "both -no-output and -o options are specified.\n";
       return 1;
     }
+    if (!dstDir.empty()) {
+      dst = sOutputDirAbsPath + "/" + dst;
+    }
   }
   if (NoOutput && Inplace) {
     llvm::errs() << sHipify << sConflict << "both -no-output and -inplace options are specified.\n";
     return 1;
   }
+  if (!dstDir.empty() && Inplace) {
+    llvm::errs() << sHipify << sConflict << "both -o-dir and -inplace options are specified.\n";
+    return 1;
+  }
   if (Examine) {
     NoOutput = PrintStats = true;
   }
   int Result = 0;
-  std::error_code EC;
   SmallString<128> tmpFile;
-  SmallString<256> sourceAbsPath, tmpDirAbsPath;
-  StringRef sourceFileName, sourceDir, ext = "hip";
-  std::string sTmpDirAbsParh, sTmpFileName;
-  if (!TemporaryDir.empty()) {
-    EC = sys::fs::real_path(TemporaryDir, tmpDirAbsPath, true);
-    if (!EC && sys::fs::is_regular_file(tmpDirAbsPath)) {
-      llvm::errs() << "\n" << sHipify << sError << TemporaryDir << " is not a directory\n";
-      return 1;
-    }
-    if (EC) {
-      EC = sys::fs::create_directory(TemporaryDir);
-      if (EC) {
-        llvm::errs() << "\n" << sHipify << sError << EC.message() << ": temporary directory: " << TemporaryDir << "\n";
-        return 1;
-      }
-      EC = sys::fs::real_path(TemporaryDir, tmpDirAbsPath, true);
-      if (EC) {
-        llvm::errs() << "\n" << sHipify << sError << EC.message() << ": temporary directory: " << TemporaryDir << "\n";
-        return 1;
-      }
-    }
-    sTmpDirAbsParh = tmpDirAbsPath.c_str();
+  SmallString<256> sourceAbsPath;
+  StringRef sourceFileName, ext = "hip";
+  std::string sTmpFileName;
+  std::string sTmpDirAbsParh = getAbsoluteDirectory(TemporaryDir, EC);
+  if (EC) {
+    return 1;
   }
   // Arguments for the Statistics print routines.
   std::unique_ptr<std::ostream> csv = nullptr;
@@ -99,13 +123,15 @@ int main(int argc, const char **argv) {
   if (PrintStats) {
     statPrint = &llvm::errs();
   }
-  int FD;
   for (const auto & src : fileSources) {
     if (dst.empty()) {
       if (Inplace) {
         dst = src;
       } else {
         dst = src + "." + ext.str();
+        if (!dstDir.empty()) {
+          dst = sOutputDirAbsPath + "/" + dst;
+        }
       }
     }
     // Create a copy of the file to work on. When we're done, we'll move this onto the
@@ -118,7 +144,7 @@ int main(int argc, const char **argv) {
       continue;
     }
     if (TemporaryDir.empty()) {
-      EC = sys::fs::createTemporaryFile(src, ext, FD, tmpFile);
+      EC = sys::fs::createTemporaryFile(src, ext, tmpFile);
       if (EC) {
         llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << tmpFile << "\n";
         Result = 1;

include/hip/hcc_detail/math_functions.h

@@ -518,9 +518,11 @@ float __expf(float x) { return __ocml_exp_f32(x); }
 __DEVICE__
 inline
 float __fadd_rd(float x, float y) { return __ocml_add_rtn_f32(x, y); }
+#endif
 __DEVICE__
 inline
-float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); }
+float __fadd_rn(float x, float y) { return x + y; }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 float __fadd_ru(float x, float y) { return __ocml_add_rtp_f32(x, y); }
@@ -530,9 +532,11 @@ float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); }
 __DEVICE__
 inline
 float __fdiv_rd(float x, float y) { return __ocml_div_rtn_f32(x, y); }
+#endif
 __DEVICE__
 inline
-float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); }
+float __fdiv_rn(float x, float y) { return x / y; }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 float __fdiv_ru(float x, float y) { return __ocml_div_rtp_f32(x, y); }
@@ -550,12 +554,14 @@ float __fmaf_rd(float x, float y, float z)
 {
     return __ocml_fma_rtn_f32(x, y, z);
 }
+#endif
 __DEVICE__
 inline
 float __fmaf_rn(float x, float y, float z)
 {
-    return __ocml_fma_rte_f32(x, y, z);
+    return __ocml_fma_f32(x, y, z);
 }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 float __fmaf_ru(float x, float y, float z)
@@ -571,9 +577,11 @@ float __fmaf_rz(float x, float y, float z)
 __DEVICE__
 inline
 float __fmul_rd(float x, float y) { return __ocml_mul_rtn_f32(x, y); }
+#endif
 __DEVICE__
 inline
-float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); }
+float __fmul_rn(float x, float y) { return x * y; }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 float __fmul_ru(float x, float y)  { return __ocml_mul_rtp_f32(x, y); }
@@ -583,24 +591,30 @@ float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); }
 __DEVICE__
 inline
 float __frcp_rd(float x) { return __llvm_amdgcn_rcp_f32(x); }
+#endif
 __DEVICE__
 inline
 float __frcp_rn(float x) { return __llvm_amdgcn_rcp_f32(x); }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 float __frcp_ru(float x) { return __llvm_amdgcn_rcp_f32(x); }
 __DEVICE__
 inline
 float __frcp_rz(float x) { return __llvm_amdgcn_rcp_f32(x); }
+#endif
 __DEVICE__
 inline
 float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); }
+#endif
 __DEVICE__
 inline
-float __fsqrt_rn(float x) { return __ocml_sqrt_rte_f32(x); }
+float __fsqrt_rn(float x) { return __ocml_sqrt_f32(x); }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 float __fsqrt_ru(float x) { return __ocml_sqrt_rtp_f32(x); }
@@ -610,9 +624,11 @@ float __fsqrt_rz(float x) { return __ocml_sqrt_rtz_f32(x); }
 __DEVICE__
 inline
 float __fsub_rd(float x, float y) { return __ocml_sub_rtn_f32(x, y); }
+#endif
 __DEVICE__
 inline
-float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); }
+float __fsub_rn(float x, float y) { return x - y; }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 float __fsub_ru(float x, float y) { return __ocml_sub_rtp_f32(x, y); }
@@ -1042,9 +1058,11 @@ double yn(int n, double x)
 __DEVICE__
 inline
 double __dadd_rd(double x, double y) { return __ocml_add_rtn_f64(x, y); }
+#endif
 __DEVICE__
 inline
-double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); }
+double __dadd_rn(double x, double y) { return x + y; }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 double __dadd_ru(double x, double y) { return __ocml_add_rtp_f64(x, y); }
@@ -1054,9 +1072,11 @@ double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); }
 __DEVICE__
 inline
 double __ddiv_rd(double x, double y) { return __ocml_div_rtn_f64(x, y); }
+#endif
 __DEVICE__
 inline
-double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); }
+double __ddiv_rn(double x, double y) { return x / y; }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 double __ddiv_ru(double x, double y) { return __ocml_div_rtp_f64(x, y); }
@@ -1066,9 +1086,11 @@ double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); }
 __DEVICE__
 inline
 double __dmul_rd(double x, double y) { return __ocml_mul_rtn_f64(x, y); }
+#endif
 __DEVICE__
 inline
-double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); }
+double __dmul_rn(double x, double y) { return x * y; }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 double __dmul_ru(double x, double y) { return __ocml_mul_rtp_f64(x, y); }
@@ -1078,9 +1100,11 @@ double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); }
 __DEVICE__
 inline
 double __drcp_rd(double x) { return __llvm_amdgcn_rcp_f64(x); }
+#endif
 __DEVICE__
 inline
 double __drcp_rn(double x) { return __llvm_amdgcn_rcp_f64(x); }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 double __drcp_ru(double x) { return __llvm_amdgcn_rcp_f64(x); }
@@ -1090,9 +1114,11 @@ double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); }
 __DEVICE__
 inline
 double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); }
+#endif
 __DEVICE__
 inline
-double __dsqrt_rn(double x) { return __ocml_sqrt_rte_f64(x); }
+double __dsqrt_rn(double x) { return __ocml_sqrt_f64(x); }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 double __dsqrt_ru(double x) { return __ocml_sqrt_rtp_f64(x); }
@@ -1102,9 +1128,11 @@ double __dsqrt_rz(double x) { return __ocml_sqrt_rtz_f64(x); }
 __DEVICE__
 inline
 double __dsub_rd(double x, double y) { return __ocml_sub_rtn_f64(x, y); }
+#endif
 __DEVICE__
 inline
-double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); }
+double __dsub_rn(double x, double y) { return x - y; }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 double __dsub_ru(double x, double y) { return __ocml_sub_rtp_f64(x, y); }
@@ -1117,12 +1145,14 @@ double __fma_rd(double x, double y, double z)
 {
     return __ocml_fma_rtn_f64(x, y, z);
 }
+#endif
 __DEVICE__
 inline
 double __fma_rn(double x, double y, double z)
 {
-    return __ocml_fma_rte_f64(x, y, z);
+    return __ocml_fma_f64(x, y, z);
 }
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
 __DEVICE__
 inline
 double __fma_ru(double x, double y, double z)

tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp

@@ -36,31 +36,45 @@ THE SOFTWARE.
 __device__ void double_precision_intrinsics() {
 #if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dadd_rd(0.0, 1.0);
+#endif
     __dadd_rn(0.0, 1.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dadd_ru(0.0, 1.0);
     __dadd_rz(0.0, 1.0);
     __ddiv_rd(0.0, 1.0);
+#endif
     __ddiv_rn(0.0, 1.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __ddiv_ru(0.0, 1.0);
     __ddiv_rz(0.0, 1.0);
     __dmul_rd(1.0, 2.0);
+#endif
     __dmul_rn(1.0, 2.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dmul_ru(1.0, 2.0);
     __dmul_rz(1.0, 2.0);
     __drcp_rd(2.0);
+#endif
     __drcp_rn(2.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __drcp_ru(2.0);
     __drcp_rz(2.0);
     __dsqrt_rd(4.0);
+#endif
     __dsqrt_rn(4.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dsqrt_ru(4.0);
     __dsqrt_rz(4.0);
     __dsub_rd(2.0, 1.0);
+#endif
     __dsub_rn(2.0, 1.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __dsub_ru(2.0, 1.0);
     __dsub_rz(2.0, 1.0);
     __fma_rd(1.0, 2.0, 3.0);
+#endif
     __fma_rn(1.0, 2.0, 3.0);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fma_ru(1.0, 2.0, 3.0);
     __fma_rz(1.0, 2.0, 3.0);
 #endif

tests/src/deviceLib/hipFloatMath.cpp

@@ -38,10 +38,12 @@ __global__ void floatMath(float* In, float* Out) {
     Out[tid] = __cosf(In[tid]);
     Out[tid] = __exp10f(Out[tid]);
     Out[tid] = __expf(Out[tid]);
-#if defined OCML_BASIC_ROUNDED_OPERATIONS
     Out[tid] = __frsqrt_rn(Out[tid]);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     Out[tid] = __fsqrt_rd(Out[tid]);
+#endif
     Out[tid] = __fsqrt_rn(Out[tid]);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     Out[tid] = __fsqrt_ru(Out[tid]);
     Out[tid] = __fsqrt_rz(Out[tid]);
 #endif

tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp

@@ -41,35 +41,51 @@ __device__ void single_precision_intrinsics() {
     __expf(0.0f);
 #if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fadd_rd(0.0f, 1.0f);
+#endif
     __fadd_rn(0.0f, 1.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fadd_ru(0.0f, 1.0f);
     __fadd_rz(0.0f, 1.0f);
     __fdiv_rd(4.0f, 2.0f);
+#endif
     __fdiv_rn(4.0f, 2.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fdiv_ru(4.0f, 2.0f);
     __fdiv_rz(4.0f, 2.0f);
 #endif
     __fdividef(4.0f, 2.0f);
 #if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fmaf_rd(1.0f, 2.0f, 3.0f);
+#endif
     __fmaf_rn(1.0f, 2.0f, 3.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fmaf_ru(1.0f, 2.0f, 3.0f);
     __fmaf_rz(1.0f, 2.0f, 3.0f);
     __fmul_rd(1.0f, 2.0f);
+#endif
     __fmul_rn(1.0f, 2.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fmul_ru(1.0f, 2.0f);
     __fmul_rz(1.0f, 2.0f);
     __frcp_rd(2.0f);
+#endif
     __frcp_rn(2.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __frcp_ru(2.0f);
     __frcp_rz(2.0f);
+#endif
     __frsqrt_rn(4.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fsqrt_rd(4.0f);
+#endif
     __fsqrt_rn(4.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fsqrt_ru(4.0f);
     __fsqrt_rz(4.0f);
     __fsub_rd(2.0f, 1.0f);
+#endif
     __fsub_rn(2.0f, 1.0f);
+#if defined OCML_BASIC_ROUNDED_OPERATIONS
     __fsub_ru(2.0f, 1.0f);
     __fsub_rz(2.0f, 1.0f);
 #endif