From 5db57f36d7743b0e1b7368e4b3db3fa0b5ffa972 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 14 Jan 2019 13:01:57 +0530 Subject: [PATCH 1/6] Switch to faster ocml variants for trig instrinsics Change-Id: If62821e2fe1b0da91ad1b8c5580ebf1a009405e9 [ROCm/hip commit: 2b9d52b21025954f97a349fec474bf234fa41ffe] --- projects/hip/include/hip/hcc_detail/math_functions.h | 11 ++++------- projects/hip/include/hip/hcc_detail/math_fwd.h | 4 ++++ 2 files changed, 8 insertions(+), 7 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index 63e48fab29..8ac9ca068c 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -507,7 +507,7 @@ 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); } @@ -655,15 +655,12 @@ __DEVICE__ inline void __sincosf(float x, float* sptr, float* cptr) { - float tmp; - - *sptr = - __ocml_sincos_f32(x, (__attribute__((address_space(5))) float*) &tmp); - *cptr = tmp; + *sptr = __ocml_native_sin_f32(x); + *cptr = __ocml_native_cos_f32(x); } __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); } diff --git a/projects/hip/include/hip/hcc_detail/math_fwd.h b/projects/hip/include/hip/hcc_detail/math_fwd.h index e5594924ba..ab57ebb909 100644 --- a/projects/hip/include/hip/hcc_detail/math_fwd.h +++ b/projects/hip/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); @@ -227,6 +229,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__ From 937d78963bb500d550a594a65e689ad95ab7f770 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 14 Jan 2019 18:23:54 +0300 Subject: [PATCH 2/6] [HIPIFY] Introduce '-I' option as hipify-clang's one + '-I' might be set as hipify-clang option (before separator '--' or without specifying separator at all); + '-I' as a clang option might be specified as well (after options separator '--'). [ROCm/hip commit: 93aeb04cec1c6f3ae7be0c46ca74e1a57bdfae89] --- projects/hip/hipify-clang/README.md | 1 - projects/hip/hipify-clang/src/ArgParse.cpp | 7 +++++++ projects/hip/hipify-clang/src/ArgParse.h | 1 + projects/hip/hipify-clang/src/main.cpp | 6 ++++++ 4 files changed, 14 insertions(+), 1 deletion(-) diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index 812500ad34..deb951efbb 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/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 ``` diff --git a/projects/hip/hipify-clang/src/ArgParse.cpp b/projects/hip/hipify-clang/src/ArgParse.cpp index 9696cb74aa..d63a328b5b 100644 --- a/projects/hip/hipify-clang/src/ArgParse.cpp +++ b/projects/hip/hipify-clang/src/ArgParse.cpp @@ -89,4 +89,11 @@ cl::opt DashDash("-", cl::value_desc("--"), cl::cat(ToolTemplateCategory)); +cl::list I("I", + cl::desc("Add directory to include search path"), + cl::value_desc("directory"), + cl::ZeroOrMore, + cl::Prefix, + cl::cat(ToolTemplateCategory)); + cl::extrahelp CommonHelp(ct::CommonOptionsParser::HelpMessage); diff --git a/projects/hip/hipify-clang/src/ArgParse.h b/projects/hip/hipify-clang/src/ArgParse.h index 97f172d812..222062bae4 100644 --- a/projects/hip/hipify-clang/src/ArgParse.h +++ b/projects/hip/hipify-clang/src/ArgParse.h @@ -33,6 +33,7 @@ extern cl::opt OutputFilename; extern cl::opt OutputDir; extern cl::opt TemporaryDir; extern cl::opt CudaPath; +extern cl::list I; extern cl::opt Inplace; extern cl::opt SaveTemps; extern cl::opt NoBackup; diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index 6df8ae9679..e219f73023 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/hipify-clang/src/main.cpp @@ -199,6 +199,12 @@ int main(int argc, const char **argv) { #if defined(HIPIFY_CLANG_RES) Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); #endif + if (!I.empty()) { + for (std::string s : I) { + 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! From 225a0f331046f3295e8af4b9c6a73b9dcba0fb56 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 14 Jan 2019 19:34:42 +0300 Subject: [PATCH 3/6] [HIPIFY] Introduce '-D' option as hipify-clang's one + '-D' might be set as hipify-clang option (before separator '--' or without specifying separator at all); + '-D' as a clang option might be specified as well (after options separator '--'); + All defines specified as hipify-clang defines (before options separator '--') will be set after clang's ones and will redefine the same clang's ones. [ROCm/hip commit: 13ffeb47ceed8458e33a76ecbf2a31a220173574] --- projects/hip/hipify-clang/src/ArgParse.cpp | 9 ++++++++- projects/hip/hipify-clang/src/ArgParse.h | 1 + projects/hip/hipify-clang/src/main.cpp | 6 ++++++ 3 files changed, 15 insertions(+), 1 deletion(-) diff --git a/projects/hip/hipify-clang/src/ArgParse.cpp b/projects/hip/hipify-clang/src/ArgParse.cpp index d63a328b5b..e81ada1ddc 100644 --- a/projects/hip/hipify-clang/src/ArgParse.cpp +++ b/projects/hip/hipify-clang/src/ArgParse.cpp @@ -85,7 +85,7 @@ cl::opt Examine("examine", cl::cat(ToolTemplateCategory)); cl::opt 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)); @@ -96,4 +96,11 @@ cl::list I("I", cl::Prefix, cl::cat(ToolTemplateCategory)); +cl::list D("D", + cl::desc("Define to or 1 if omitted"), + cl::value_desc("macro>= OutputDir; extern cl::opt TemporaryDir; extern cl::opt CudaPath; extern cl::list I; +extern cl::list D; extern cl::opt Inplace; extern cl::opt SaveTemps; extern cl::opt NoBackup; diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index e219f73023..db6c803c27 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/hipify-clang/src/main.cpp @@ -199,6 +199,12 @@ int main(int argc, const char **argv) { #if defined(HIPIFY_CLANG_RES) Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); #endif + if (!D.empty()) { + for (std::string s : D) { + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-D", ct::ArgumentInsertPosition::END)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(s.c_str(), ct::ArgumentInsertPosition::END)); + } + } if (!I.empty()) { for (std::string s : I) { Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-I", ct::ArgumentInsertPosition::END)); From 2f5951d54c13caff20b31a82e42a9e354ebe0981 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 14 Jan 2019 21:19:34 +0300 Subject: [PATCH 4/6] [HIPIFY] rename variables as requested in review [ROCm/hip commit: e6c09bb9afaa0eabdbac060d1152049f73d18967] --- projects/hip/hipify-clang/src/ArgParse.cpp | 4 ++-- projects/hip/hipify-clang/src/ArgParse.h | 4 ++-- projects/hip/hipify-clang/src/main.cpp | 8 ++++---- 3 files changed, 8 insertions(+), 8 deletions(-) diff --git a/projects/hip/hipify-clang/src/ArgParse.cpp b/projects/hip/hipify-clang/src/ArgParse.cpp index e81ada1ddc..3cc372adab 100644 --- a/projects/hip/hipify-clang/src/ArgParse.cpp +++ b/projects/hip/hipify-clang/src/ArgParse.cpp @@ -89,14 +89,14 @@ cl::opt DashDash("-", cl::value_desc("--"), cl::cat(ToolTemplateCategory)); -cl::list I("I", +cl::list IncludeDirs("I", cl::desc("Add directory to include search path"), cl::value_desc("directory"), cl::ZeroOrMore, cl::Prefix, cl::cat(ToolTemplateCategory)); -cl::list D("D", +cl::list MacroNames("D", cl::desc("Define to or 1 if omitted"), cl::value_desc("macro>= OutputFilename; extern cl::opt OutputDir; extern cl::opt TemporaryDir; extern cl::opt CudaPath; -extern cl::list I; -extern cl::list D; +extern cl::list IncludeDirs; +extern cl::list MacroNames; extern cl::opt Inplace; extern cl::opt SaveTemps; extern cl::opt NoBackup; diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index db6c803c27..d0fcc17e66 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/hipify-clang/src/main.cpp @@ -199,14 +199,14 @@ int main(int argc, const char **argv) { #if defined(HIPIFY_CLANG_RES) Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); #endif - if (!D.empty()) { - for (std::string s : D) { + 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 (!I.empty()) { - for (std::string s : I) { + 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)); } From 09fe2d6bc5e2ed3b4d4e99c3a197685865c5dd49 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 15 Jan 2019 19:40:37 +0530 Subject: [PATCH 5/6] Map more instrincis to ocml native Change-Id: I54e55d5902931bafdabd0956e4e8c1d7b39f7173 [ROCm/hip commit: dd9bea7d25f2417b6f55159f592488614ffe7970] --- .../include/hip/hcc_detail/math_functions.h | 12 ++++++------ projects/hip/include/hip/hcc_detail/math_fwd.h | 18 ++++++++++++++++++ 2 files changed, 24 insertions(+), 6 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index 8ac9ca068c..3e9db0d1f6 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -510,10 +510,10 @@ inline 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); } diff --git a/projects/hip/include/hip/hcc_detail/math_fwd.h b/projects/hip/include/hip/hcc_detail/math_fwd.h index ab57ebb909..cf8aeb9c6c 100644 --- a/projects/hip/include/hip/hcc_detail/math_fwd.h +++ b/projects/hip/include/hip/hcc_detail/math_fwd.h @@ -94,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)) @@ -154,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)) @@ -239,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)) From 6789cec88f717808d931c11a60fcfd5a2e53a29f Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 16 Jan 2019 19:30:09 +0300 Subject: [PATCH 6/6] [HIPIFY][fix][#867] Relative paths routines fixes and improvements + Always check for existence any file or dir, first of all, because real_path and make_absolute differ: real_path additionally performs access on file/dir. + Relative path routines in source files implemented similarly to dirs. + Error handling improvements. [ROCm/hip commit: 1be38618e2ba6ed5b1f75594724c93006bba4552] --- projects/hip/hipify-clang/src/main.cpp | 89 ++++++++++++++++---------- 1 file changed, 56 insertions(+), 33 deletions(-) diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index d0fcc17e66..6df2f57073 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/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(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 (!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 (sys::fs::exists(sDir)) { + if (sys::fs::is_regular_file(sDir)) { + 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 ""; + } + } 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(std::errc::no_such_file_or_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 = llcompat::real_path(sDir, dirAbsPath, true); - if (EC) { - llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n"; - 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 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 sTmpDirAbsParh = getAbsoluteDirectory(TemporaryDir, EC); + std::string sTmpFileName, sSourceAbsPath; + 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 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));