From 58dfeff27a198843bc35ec0e8cb024b6962e0129 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 18 Dec 2018 19:31:54 +0000 Subject: [PATCH 1/3] Add *_rn functions back into HIP intrinsics Add back the round-to-nearest-even intrinsics back to HIP math intrinsics as it caused regression. --- include/hip/hcc_detail/math_functions.h | 54 ++++++++++++++----- .../hipDoublePrecisionIntrinsics.cpp | 14 +++++ tests/src/deviceLib/hipFloatMath.cpp | 4 +- .../hipSinglePrecisionIntrinsics.cpp | 16 ++++++ 4 files changed, 75 insertions(+), 13 deletions(-) diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 557257b2b0..63e48fab29 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/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) diff --git a/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp b/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp index 939bdae743..24a8320840 100644 --- a/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp +++ b/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 diff --git a/tests/src/deviceLib/hipFloatMath.cpp b/tests/src/deviceLib/hipFloatMath.cpp index c6a07e26a9..ab1deb477a 100644 --- a/tests/src/deviceLib/hipFloatMath.cpp +++ b/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 diff --git a/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp b/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp index b216b3cb54..751ea88641 100644 --- a/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp +++ b/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 From c8d720081f58d976b87686e5d2a3d4cc6178d1f1 Mon Sep 17 00:00:00 2001 From: emankov Date: Wed, 26 Dec 2018 16:30:41 +0300 Subject: [PATCH 2/3] [HIPIFY] Introduce option -o-dir Option -o-dir for output directory: + if not specified source file(s) directory is used; + creates the directory if the directory doesn't exist (only one level in a tree); + if -o and -o-dir both are specified the hipified file is placed to "-o-dir" + "-o"; + reports an error in case of a wrong directory specified, in case of necessity of creating a tree of subfolders, or in case of a filename specified. --- hipify-clang/src/ArgParse.cpp | 5 +++ hipify-clang/src/ArgParse.h | 1 + hipify-clang/src/main.cpp | 75 ++++++++++++++++++++++++----------- 3 files changed, 57 insertions(+), 24 deletions(-) diff --git a/hipify-clang/src/ArgParse.cpp b/hipify-clang/src/ArgParse.cpp index 87aa8b6755..e5915aedb9 100644 --- a/hipify-clang/src/ArgParse.cpp +++ b/hipify-clang/src/ArgParse.cpp @@ -29,6 +29,11 @@ cl::opt OutputFilename("o", cl::value_desc("filename"), cl::cat(ToolTemplateCategory)); +cl::opt OutputDir("o-dir", + cl::desc("Output direcory"), + cl::value_desc("directory"), + cl::cat(ToolTemplateCategory)); + cl::opt TemporaryDir("temp-dir", cl::desc("Temporary direcory"), cl::value_desc("directory"), diff --git a/hipify-clang/src/ArgParse.h b/hipify-clang/src/ArgParse.h index efffe9e5de..70a6a9b28d 100644 --- a/hipify-clang/src/ArgParse.h +++ b/hipify-clang/src/ArgParse.h @@ -30,6 +30,7 @@ namespace ct = clang::tooling; extern cl::OptionCategory ToolTemplateCategory; extern cl::opt OutputFilename; +extern cl::opt OutputDir; extern cl::opt TemporaryDir; extern cl::opt Inplace; extern cl::opt SaveTemps; diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index 5417dda783..7af894e16c 100644 --- a/hipify-clang/src/main.cpp +++ b/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(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 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 csv = nullptr; @@ -106,6 +130,9 @@ int main(int argc, const char **argv) { 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 From 912c4710ed6cf82ebcde95bc0af9912f68d381e5 Mon Sep 17 00:00:00 2001 From: emankov Date: Thu, 27 Dec 2018 17:00:18 +0300 Subject: [PATCH 3/3] [HIPIFY] Use sys::fs::createTemporaryFile without file creation [Reasons] + File is opened further by clang + There is a regression with LLVM 5.x, when system error occurs: temporary file couldn't be renamed - permission denied. --- hipify-clang/src/main.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index 7af894e16c..06981a9d70 100644 --- a/hipify-clang/src/main.cpp +++ b/hipify-clang/src/main.cpp @@ -123,7 +123,6 @@ int main(int argc, const char **argv) { if (PrintStats) { statPrint = &llvm::errs(); } - int FD; for (const auto & src : fileSources) { if (dst.empty()) { if (Inplace) { @@ -145,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;