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..3cc372adab 100644 --- a/projects/hip/hipify-clang/src/ArgParse.cpp +++ b/projects/hip/hipify-clang/src/ArgParse.cpp @@ -85,8 +85,22 @@ 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)); +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 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 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 6df8ae9679..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)); @@ -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! diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index 63e48fab29..3e9db0d1f6 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/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_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..cf8aeb9c6c 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); @@ -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))